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

共享內存（shared memory）是位於SM上的on-chip（片上）一塊內存，每個SM都有，就是內存比較小，早期的GPU只有16K（16384），現在生產的GPU一般都是48K（49152）。

共享內存由於是片上內存，因而帶寬高，延遲小（較全局內存而言），合理使用共享內存對程序效率具有很大提升。

下面是使用共享內存對一個數組進行求和，使用全局內存進行歸約求和可以瀏覽https://www.cnblogs.com/xiaoxiaoyibu/p/11397205.html

#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;

    //----使用shared memory--------------------------------------------------------------

    __shared__ double s_Para[N];    　　　　//定義長度為N的共享內存數組
    if (tid < N)                    　　　　//循環整個數組，每個線程負責將一個元素從全局內存載入共享內存
        s_Para[tid] = Para[tid];
    __syncthreads();                　　　　//(紅色下波浪線提示由於VS不識別，不影響運行)同步，等待所有線程把自己負責的元素載入到共享內存再執行下面代碼

    for (int index = 1; index < blockDim.x; index *= 2)
    {
　　　　　__syncthreads();　　　　　　　　　　 //同步，以防止歸約過程中某個線程運行速度過快導致計算錯誤（后面線程計算使用錯誤的前面線程結果值）

　　　　if (tid % (2 * index) == 0) 
　　　　{ 
　　　　　　s_Para[tid] += s_Para[tid + index]; 
　　　　} 
　　} 

　　if (tid == 0)　　　　　　　　　　　　　　//整個數組相加完成后，將共享內存數組0號元素的值賦給全局內存數組0號元素，最后返回CPU端 
　　　　Para[tid] = s_Para[tid];  
}

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