Test with line numbers

  1.  
  2.  
  3.  
  4.  
  5.  
  6.  
  7.  
  8.  
  9.  
  10.  
  11.  
  12.  
  13.  
  14.  
  15.  
  16.  
  17.  
  18.  
  19.  
  20.  
  21.  
  22.  
  23.  
  24.  
  25.  
  26.  
  27.  
  28.  
  29.  
  30.  
  31.  
  32.  
template <typename T, typename BinaryOperator>
__global__ void dev_vector_op(T* a, T* b, T* out, unsigned int len) {
    unsigned int const idx = blockDim.x * blockIdx.x + threadIdx.x;
    BinaryOperator op;
    out[idx] = op(a[idx], b[idx]);
}

template <typename T, typename BinOp>
__global__ void dev_reduce_block(T* gdata, T* result, unsigned int const len) {
    extern __shared__ T sdata[];
    unsigned int const offset = threadIdx.x;
    BinOp const op;

    sdata[2 * offset] = gdata[2 * offset];
    sdata[2 * offset + 1] = gdata[2 * offset + 1];

#define IDX(x) ((x) + 2 * offset) * d - 1

    for (unsigned int d = 1; d <= len / 2; d *= 2) {
        __syncthreads();
        unsigned int const ai = IDX(1);
        if (ai >= len) continue;
        unsigned int const bi = IDX(2);
        sdata[bi] = op(sdata[ai], sdata[bi]);
    }

#undef IDX
    
    __syncthreads();
    if (threadIdx.x == 0)
        *result = sdata[len - 1];
}