   x |= x >> 16;
   return ++x;
 }
-
 
 unsigned int n_reduce_blocks(unsigned int n) {
   return min(max_blocks, (n + (2*reduce_cta_size - 1)) / (2*reduce_cta_size));

    return a + b;
  }
};


template <typename T, typename Op, unsigned int block_size>
__global__ void reduce(const T *a, unsigned int n, T *b) {

    // Reduce in shared memory
    sdata[tid] = reduced;
    __syncthreads();

    #pragma unroll
    for (unsigned int i=512; i >= 2; i >>= 1) {
      if (block_size >= i) {
        if (tid < (i << 1)) {
          sdata[tid] = reduced = op(reduced, sdata[tid + i]);
        }
        // No need to sync threads in the same warp
        if (tid >= 32)
          __syncthreads();
      }
    }

