title: Finding min and max by the cuda kernel reduction
body: <p>Here is my kernel code</p>
  
  <pre><code>typedef unsigned char Npp8u;
  ...
      // Kernel Implementation
  __device__ unsigned int min_device;
  __device__ unsigned int max_device;
  
  
  __global__ void findMax_Min(Npp8u * data, int numEl){
      int index = blockDim.x*blockIdx.x + threadIdx.x;
      int shared_index = threadIdx.x;
  
      __shared__ Npp8u data_shared_min[BLOCKDIM];
      __shared__ Npp8u data_shared_max[BLOCKDIM];
  
      // check index condition
      if(index &lt; numEl){
          data_shared_min[shared_index] = data[index]; //pass values from global to shared memory
          __syncthreads();
          data_shared_max[shared_index] = data[index]; //pass values from global to shared memory
  
  
          for (unsigned int stride = BLOCKDIM/2; stride &gt; 0; stride &gt;&gt;= 1) {
              if(threadIdx.x &lt;  stride){
                  if(data_shared_max[threadIdx.x] &lt;  data_shared_max[threadIdx.x+stride]) data_shared_max[shared_index] = data_shared_max[shared_index+stride];
                  if(data_shared_min[threadIdx.x]&gt;  data_shared_min[threadIdx.x+stride]) data_shared_min[shared_index] = data_shared_min[shared_index+stride];
              }
              __syncthreads();
          }
          if(threadIdx.x == 0  ){
              atomicMin(&amp;(min_device), (unsigned int)data_shared_min[threadIdx.x ]);
              //min_device =10;
              __syncthreads();
              atomicMax(&amp;(max_device), (unsigned int)data_shared_max[threadIdx.x ]);
          }
      }else{
          data_shared_min[shared_index] = 9999;
      }
  }
  </code></pre>
  
  <p>I have an image that is 512x512 and I want to find the min and max pixel values. <code>data</code> is the 1-D version of the image. This code works for max but not for min value. As I checked from matlab max value is 202 and min value is 10 but it finds 0 for the min value. Here is my kernel codes and memcpy calls</p>
  
  <pre><code>int main(){
      // Host parameter declarations.
      Npp8u * imageHost;
      int   nWidth, nHeight, nMaxGray;
  
      // Load image to the host.
      std::cout &lt;&lt; "Load PGM file." &lt;&lt; std::endl;
      imageHost = LoadPGM("lena_before.pgm", nWidth, nHeight, nMaxGray);
  
      // Device parameter declarations.
      Npp8u    * imageDevice;
      unsigned int   max, min;
      size_t size = sizeof(Npp8u)*nWidth*nHeight;
  
      cudaMalloc((Npp8u**)&amp;imageDevice, size);
  
      cudaMemcpy(imageDevice, imageHost, size, cudaMemcpyHostToDevice);
  
      int numPixels = nWidth*nHeight;
  
      dim3 numThreads(BLOCKDIM);
      dim3 numBlocks(numPixels/BLOCKDIM + (numPixels%BLOCKDIM == 0 ? 0 : 1));
  
      findMax_Min&lt;&lt;&lt;numBlocks, numThreads&gt;&gt;&gt;(imageDevice,numPixels);
      cudaMemcpyFromSymbol(&amp;max,max_device, sizeof(max_device), 0, cudaMemcpyDeviceToHost);
      cudaMemcpyFromSymbol(&amp;min,min_device, sizeof(min_device), 0, cudaMemcpyDeviceToHost);
  
  
      printf("Min value for image : %i\n", min);
      printf("Max value for image : %i\n", max);
  ...
  </code></pre>
  
  <p>Another interesting thing is changing the order of <code>cudaMemcpy</code> just after the kernel call also causes malfunctioning and values both are read as zero. I do not see the problem. Is there anyone sees the obstructed part?</p>