问题:使用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"
#include <iostream>
using 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<N; i++)
    {
        Para[i] = (i + 1) * 0.1;                          //数组赋值
        ParaSum += Para[i];                                //CPU端数组累加
    }
    cout << " CPU result = " << ParaSum << endl;          //显示CPU端结果
    double d_ParaSum;
    d_ParallelTest << < 1, N >> > (Para);                  //调用核函数(一个包含N个线程的线程块)
    cudaDeviceSynchronize();                              //同步
    d_ParaSum = Para[0];                                  //从累加过后数组的0号元素得出结果
    cout << " GPU result = " << d_ParaSum << endl;         //显示GPU端结果
}
int main() {
//并行归约
    ParallelTest();                        //调用归约函数
    system("pause");
    return 0;
}
结果如下所示(CPU和GPU计算结果一致):

原文:https://www.cnblogs.com/xiaoxiaoyibu/p/11397205.html