title: From given vector find max value and its index by reduction method in CUDA
body: <p>I am new to CUDA, for the vector to find max value and its index I use CUDA</p>
  
  <p>here its my code:</p>
  
  <pre><code>#include &lt; cuda.h &gt;
  #include &lt; stdio.h &gt;
  #include &lt; time.h &gt;
  #include &lt;iostream&gt;
  using namespace std;
  #define tbp 256
  #define nblocks 1
  __global__ void kernel_max(int *a, int *d, int *index,int *idx)
  {
      __shared__ int sdata[tbp]; //"static" shared memory
  
      int tid = threadIdx.x;
      int i = blockIdx.x * blockDim.x + threadIdx.x;
  
      sdata[tid] = a[i];
      index[tid] = i;
      __syncthreads();
      for(int s=tbp/2 ; s &gt;= 1 ; s=s/2)
      {
          if(tid &lt; s)
          {
              if(sdata[tid] &lt; sdata[tid + s])
              {
                  sdata[tid] = sdata[tid + s];
                  index[tid] = index[tid+s];
               __syncthreads();
              }
              __syncthreads();
          }
          __syncthreads();
      }
      __syncthreads();
      if(tid == 0 ) 
      {
          d[blockIdx.x] = sdata[0];
          idx[blockIdx.x] = index[0];
      }
      __syncthreads();
  }
  
  int main()
  {
      int i;
      const int N=tbp*nblocks;
      srand(time(NULL));
  
      int *a;
      a = (int*)malloc(N * sizeof(int));
      int *d;
      d = (int*)malloc(nblocks * sizeof(int));
      int *index;
      index = (int*)malloc(N * sizeof(int));
      int *idx;
      idx = (int*)malloc(nblocks * sizeof(int));
  
      int *dev_a, *dev_d, *dev_index,*dev_idx;
  
      cudaMalloc((void **) &amp;dev_a, N*sizeof(int));
      cudaMalloc((void **) &amp;dev_d, nblocks*sizeof(int));
      cudaMalloc((void **) &amp;dev_index, N*sizeof(int));
      cudaMalloc((void **) &amp;dev_idx, nblocks*sizeof(int));
  
      int mmm=0;
      int ddd=0;
      for( i = 0 ; i &lt; N ; i++)
      {
          a[i] = rand()% 100 + 5;
          index[i]=i;
          //printf("%d\n",a[i]);
          if(mmm&lt;a[i]) 
          {
              mmm=a[i];
              ddd=i;  
          }
      }
      printf("");
      printf("");
      printf("");
      printf("");
      cudaMemcpy(dev_a , a, N*sizeof(int),cudaMemcpyHostToDevice);
      cudaMemcpy(dev_index , index, N*sizeof(int),cudaMemcpyHostToDevice);
      kernel_max &lt;&lt;&lt; nblocks,tbp &gt;&gt;&gt;(dev_a,dev_d,dev_index,dev_idx);
  
      cudaMemcpy(d, dev_d, nblocks*sizeof(int),cudaMemcpyDeviceToHost);
      cudaMemcpy(index, dev_index, N*sizeof(int),cudaMemcpyDeviceToHost);
      cudaMemcpy(idx, dev_idx, nblocks*sizeof(int),cudaMemcpyDeviceToHost);
  
      printf("cpu max= %d, gpu_max = %d ,cpu index: %d, gpu index: %d",mmm,d[0],ddd,idx[0]);
      printf("\n");
  
      if(ddd!=idx[0])
      {
          cout&lt;&lt;"index mismatch!damn!!"&lt;&lt;endl;
      }
      else
      {
          cout&lt;&lt;"congratulations!!"&lt;&lt;endl;
      }
      /*
      for(i=0;i&lt;N;i++)
          cout&lt;&lt;*(index+i)&lt;&lt;endl;
      */
      cudaFree(dev_a);
      cudaFree(dev_d);
      cudaFree(dev_index);
      cudaFree(dev_idx);
  
      free(a);
      free(d);
      free(index);
      free(idx);
  
      return 0;
  }
  </code></pre>
  
  <p>The problem is that for the tbp &lt; 128 it can get correct result both in value and index
  when increase to 256,512,1024, the result will sometimes go wrong.
  Can anyone given a explanation for this situation?Thanks.</p>