CUDA学习（六）之使用共享内存（shared memory）进行归约求和（M个包含N个线程的线程块）...

快来打我* 2023-08-17 17:17 8阅读 0赞

在[https://www.cnblogs.com/xiaoxiaoyibu/p/11402607.html][https_www.cnblogs.com_xiaoxiaoyibu_p_11402607.html]中介绍了使用一个包含N个线程的线程块和共享内存进行数组归约求和，

基本思路：

定义M个包含N个线程的线程块时（NThreadX = ((NX + ThreadX - 1) / ThreadX)），全局线程索引需使用tid = blockIdx.x \* blockDim.x + threadIdx.x，而在每个线程块中局部线程索引是i = threadIdx.x，

每个线程块只计算一部分求和，求和结果保存在该线程块中的共享内存数组0号元素中，线程结束后将该值赋给对应全局数组（blockIdx.x \* blockDim.x）元素中，最后在CPU端使用循环将每个线程块所求和相加，即得到最后结果。

代码如下：

#pragma once
    #include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include "device_functions.h"
    
    #include <iostream>
    
    using namespace std;
    const int NX = 10240;            //数组长度
    const int ThreadX = 256;        //线程块大小
    //使用shared memory和多个线程块
    __global__ void d_SharedMemoryTest(double *para)
    {
        int i = threadIdx.x;                                    //该线程块中线程索引
        int tid = blockIdx.x * blockDim.x + threadIdx.x;        //M个包含N个线程的线程块中相对应全局内存数组的索引（全局线程）
    
        __shared__ double s_Para[ThreadX];                        //定义固定长度（线程块长度）的共享内存数组
        if (tid < NX)                                            //判断全局线程小于整个数组长度NX，防止数组越界
            s_Para[i] = para[tid];                                //将对应全局内存数组中一段元素的值赋给共享内存数组
        __syncthreads();                                        　//(红色下波浪线提示由于VS不识别，不影响运行)同步，等待所有线程把自己负责的元素载入到共享内存再执行下面代码
    
        
        for (int index = 1; index < blockDim.x; index *= 2)        //归约求和
        {
            __syncthreads();
            if (i % (2 * index) == 0)
            {
                s_Para[i] += s_Para[i + index];
            }
        }
        
        if (i == 0)                                                //求和完成，总和保存在共享内存数组的0号元素中
            para[blockIdx.x * blockDim.x + i] = s_Para[i];        //在每个线程块中，将共享内存数组的0号元素赋给全局内存数组的对应元素，即线程块索引*线程块维度+i（blockIdx.x * blockDim.x + i）
    
    }
    
    
    //使用shared memory和多个线程块
    void s_ParallelTest()
    {
        double *Para;
        cudaMallocManaged((void **)&Para, sizeof(double) * NX);        //统一内存寻址，CPU和GPU都可以使用
    
        double ParaSum = 0;
        for (int i = 0; i<NX; i++)
        {
            Para[i] = (i + 1) * 0.01;                        //数组赋值
            ParaSum += Para[i];                                //CPU端数组累加
        }
    
        cout << " CPU result = " << ParaSum << endl;        //显示CPU端结果
        double d_ParaSum;
    
        int NThreadX = ((NX + ThreadX - 1) / ThreadX);
        cout << " 线程块大小 ：" << ThreadX << "                线程块数量 ：" << NThreadX << endl;
    
        d_SharedMemoryTest << < NThreadX, ThreadX >> > (Para);                //调用核函数（M个包含N个线程的线程块）
    
        cudaDeviceSynchronize();                            //同步
    
        for (int i=0; i<NThreadX; i++)
        {
            d_ParaSum += Para[i*ThreadX];                    //将每个线程块相加求的和（保存在对应全局内存数组中）相加求和
        }
                                        
        cout << " GPU result = " << d_ParaSum << endl;        //显示GPU端结果
    
    }
    
    int main() {
        
        s_ParallelTest();
    
    
        system("pause");
        return 0;
    }

结果如下（CPU和GPU结果一致）：

![1265435-20190823222747526-51766284.png][]

转载于:https://www.cnblogs.com/xiaoxiaoyibu/p/11403003.html

[https_www.cnblogs.com_xiaoxiaoyibu_p_11402607.html]: https://www.cnblogs.com/xiaoxiaoyibu/p/11402607.html
[1265435-20190823222747526-51766284.png]: /images/20230808/675ce4a3712e4894a077f5038fee5737.png