title: Implementing Max Reduce in Cuda
body: <p>I've been learning Cuda and I am still getting to grips with parallelism. The problem I am having at the moment is implementing a max reduce on an array of values. This is my kernel</p>
  
  <pre><code>__global__ void max_reduce(const float* const d_array,
                       float* d_max,
                       const size_t elements)
  {
      extern __shared__ float shared[];
  
      int tid = threadIdx.x;
      int gid = (blockDim.x * blockIdx.x) + tid;
  
      if (gid &lt; elements)
          shared[tid] = d_array[gid];
      __syncthreads();
  
      for (unsigned int s=blockDim.x/2; s&gt;0; s&gt;&gt;=1) 
      {
          if (tid &lt; s &amp;&amp; gid &lt; elements)
              shared[tid] = max(shared[tid], shared[tid + s]);
          __syncthreads();
      }
  
      if (gid == 0)
          *d_max = shared[tid];
  }
  </code></pre>
  
  <p>I have implemented a min reduce using the same method (replacing the max function with the min) which works fine.</p>
  
  <p>To test the kernel, I found the min and max values using a serial for loop. The min and max values always come out the same in the kernel but only the min reduce matches up.</p>
  
  <p>Is there something obvious I'm missing/doing wrong? </p>