title: CUDA Reduction minimum value and index
body: <p>I implemented a minimum reduce using <code>CUDA 8</code> by following <a href="https://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/" rel="nofollow noreferrer">this</a> great explanation and modifying it</p>
  
  <pre><code>__inline__ __device__ int warpReduceMin(int val) 
  {
      for (int offset = warpSize / 2; offset &gt; 0; offset /= 2)
      {
          int tmpVal = __shfl_down(val, offset);
          if (tmpVal &lt; val)
          {
              val = tmpVal;
          }
      }
      return val;
  }
  
  __inline__ __device__ int blockReduceMin(int val) 
  {
  
      static __shared__ int shared[32]; // Shared mem for 32 partial mins
      int lane = threadIdx.x % warpSize;
      int wid = threadIdx.x / warpSize;
  
      val = warpReduceMin(val);     // Each warp performs partial reduction
  
      if (lane == 0)
      {
          shared[wid] = val; // Write reduced value to shared memory
      }
  
      __syncthreads();              // Wait for all partial reductions
  
      //read from shared memory only if that warp existed
      val = (threadIdx.x &lt; blockDim.x / warpSize) ? shared[lane] : INT_MAX;
  
      if (wid == 0)
      {
          val = warpReduceMin(val); //Final reduce within first warp
      }
  
      return val;
  }
  
  __global__ void deviceReduceBlockAtomicKernel(int *in, int* out, int N) {
      int minVal = INT_MAX;
      for (int i = blockIdx.x * blockDim.x + threadIdx.x;
          i &lt; N;
          i += blockDim.x * gridDim.x) 
      {
          minVal = min(minVal, in[i]);
      }
      minVal = blockReduceMin(minVal);
      if (threadIdx.x == 0)
      {
          atomicMin(out, minVal);
      }
  }
  </code></pre>
  
  <p>and it works great and I'm getting the minimum value. However, I don't care about the minimum value, only about its index in the original input array.</p>
  
  <p>I tried modifying my code a bit</p>
  
  <pre><code>__inline__ __device__ int warpReduceMin(int val, int* idx) // Adding output idx
  {
      for (int offset = warpSize / 2; offset &gt; 0; offset /= 2)
      {
          int tmpVal = __shfl_down(val, offset);
          if (tmpVal &lt; val)
          {
              *idx = blockIdx.x * blockDim.x + threadIdx.x + offset; // I guess I'm missing something here
              val = tmpVal;
          }
      }
      return val;
  }
  
  ...
  blockReduceMin stayed the same only adding idx to function calls
  ...
  
  __global__ void deviceReduceBlockAtomicKernel(int *in, int* out, int N) {
      int minVal = INT_MAX;
      int minIdx = 0; // Added this
      for (int i = blockIdx.x * blockDim.x + threadIdx.x;
          i &lt; N;
          i += blockDim.x * gridDim.x) 
      {
          if (in[i] &lt; minVal)
          {
              minVal = in[i];
              minIdx = i; // Added this
          }
      }
      minVal = blockReduceMin(minVal, &amp;minIdx);
      if (threadIdx.x == 0)
      {
          int old = atomicMin(out, minVal);
          if (old != minVal) // value was updated
          {
              atomicExch(out + 1, minIdx);
          }
      }
  }
  </code></pre>
  
  <p>But it doesn't work. I feel that I'm missing something important and that this is not the way to go about it, but my search turned up no results.</p>