CUDA學習（六）之使用共享內存（shared memory）進行歸約求和（M個包含N個線程的線程塊）

在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結果一致）：