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