与CUDA
C权威编程指南的代码不同,在相邻配对2和交错配对的时候,grid.x可以除以二,因为第一周期就有一半的资源没有用上,确实优化了效率。
简介
问题:对有N个元素的整数数组求和。
当数据量很大的时候,我们可以分组求和,即
1、讲输入向量划分到更小的数据块中
2、所有的数据块并行求部分和
3、对所有的部分和进行求和得到结果
接下来介绍各种并行规约的方法。
相邻配对-1
每一个block负责计算出一个部分和。
首先有两个全局数组,g_idata和g_odata,分别存放需要求和的元素和每个block计算得出的部分和。
规约是就地完成的,也就是说每一步,全局内存g_idata中的值都会被部分和所替代。
计算结构如下图:
_syncthreads保证,进入下一次迭代之前,每个线程的部分和都已经被保存在全局内存当中。
核函数逻辑为,图中相邻两个元素的间隔为跨度,第一行跨度为1,第二行跨度为2…,每一次循环规约结束以后,跨度变成原来的两倍。
第一次循环后,偶数元素被(他自己和右侧奇数元素)的部分和替代。
第二次循环后,每四个元素产生的部分和在相应的部分。
…
核函数代码如下:
// Neighbored Pair Implementation with divergence
__global__ void reduceNeighbored (int *g_idata, int *g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if (idx >= n) return;
// in-place reduction in global memory
for (int stride = 1; stride < blockDim.x; stride *= 2)
{
if ((tid % (2 * stride)) == 0)
{
idata[tid] += idata[tid + stride];
}
// synchronize within threadblock
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = idata[0];
}
相邻配对-2
相邻配对1中判断是不是需要规约的条件为
if ((tid % (2 * stride)) == 0)
{
idata[tid]