CUDA sample学习笔记（一）reduction

1. seperated threads and interleaved addressing

template <class T>
__global__ void reduce0(T *g_idata, T *g_odata, unsigned int n) {
  // Handle to thread block group
  cg::thread_block cta = cg::this_thread_block();
  T *sdata = SharedMemory<T>();

  // load shared mem
  unsigned int tid = threadIdx.x;
  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

  sdata[tid] = (i < n) ? g_idata[i] : 0;

  cg::sync(cta);

  // do reduction in shared mem
  for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    // modulo arithmetic is slow!
    if ((tid % (2 * s)) == 0) {
      sdata[tid] += sdata[tid + s];
    }

    cg::sync(cta);
  }

  // write result for this block to global mem
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

分析：根据注释，这个算法用取模来决定哪些thread进行计算，取模运算符在GPU中效率很低，其次选择的thread都分布在不同的warp中，意味着每个循环每个warp都要执行计算，而一个block中的warps在硬件层面并不是完全并行的，需求同样多的thread情况下尽可能使用来自同一个warp的thread才能实现最大性能。

 

2. contiguous threads but with shared memory bank conflict

template <class T>
__global__ void reduce1(T *g_idata, T *g_odata, unsigned int n) {
  // Handle to thread block group
  cg::thread_block cta = cg::this_thread_block();
  T *sdata = SharedMemory<T>();

  // load shared mem
  unsigned int tid = threadIdx.x;
  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

  sdata[tid] = (i < n) ? g_idata[i] : 0;

  cg::sync(cta);

  // do reduction in shared mem
  for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    int index = 2 * s * tid;

    if (index < blockDim.x) {
      sdata[index] += sdata[index + s];
    }

    cg::sync(cta);
  }

  // write result for this block to global mem
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

分析：这种方法改变了thread号与shared memory数据下标的对应关系，例如s=1时执行计算的是threadIdx.x=0-127(假设一个block中256个thread)；s=2时，进行计算的thread为0-63；s=4时为0-31...这样随着循环的进行GPU使用了尽可能少的warp。但是这种方法会带来bank confli