问题:使用CUDA进行数组元素归约求和,归约求和的思想是每次循环取半。
详细过程如下:
假设有一个包含8个元素的数组,索引下标从0到7,现通过3次循环相加得到这8个元素的和,使用一个间隔变量,该间隔变量随循环次数改变(累乘)。
第一次循环,间隔变量stride等于1,将0与1号元素、2与3号元素、4与5号元素、6与7号元素相加并将结果分别保存在0、2、4、6号元素中(图中红色框所示)。
第二次循环,间隔变量stride等于2,将0与2号元素、4与6号元素相加并将结果分别保存在0、4号元素中(图中红色框所示)。
第三次循环,间隔变量stride等于4,将0与4号元素相加并将结果保存在0号元素中(图中红色框所示)。
三次循环过后,整个数组元素相加之和就保存在数组0号元素中。
代码如下:
#pragma once#include "cuda_runtime.h" #include "device_launch_parameters.h" #include "device_functions.h" #includeusing namespace std; const int N = 128; //数组长度 __global__ void d_ParallelTest(double *Para) { int tid = threadIdx.x; //----随循环次数的增加,stride逐次翻倍(乘以2)----------------------------------------------------- for (int stride = 1; stride < blockDim.x; stride *= 2) { if (tid % (2 * stride) == 0) { Para[tid] += Para[tid + stride]; //对应上图中红色框的元素 } __syncthreads(); } } void ParallelTest() { double *Para; cudaMallocManaged((void **)&Para, sizeof(double) * N); //统一内存寻址,CPU和GPU都可以使用的数组 double ParaSum = 0; for (int i = 0; i
结果如下所示(CPU和GPU计算结果一致):