CUDA的存儲器可以大致分為兩類：
板載顯存（On-board memory）
片上內存（On-chip memory）
其中板載顯存主要包括全局內存（global memory）、本地內存（local memory）、常量內存（constant memory）、紋理內存（texture memory）等，片上內存主要包括寄存器（register）和共享內存（shared memory）。不同類型的內存有各自不同的特點，不過片上內存通常比板載顯存要快，而寄存器又是所有存儲種類中最快的。本文我們着重介紹共享內存的基礎知識以及應用例子。
01
—
查看自己顯卡上的共享內存信息
CUDA提供了cudaGetDeviceCount和cudaGetDeviceProperties這兩個函數，分別用於獲取CUDA設備數、獲取CUDA設備屬性，通過調用這兩個函數，可以方便獲取共享內存信息和其它CUDA設備信息：
//顯示CUDA設備信息
void show_GPU_info(void)
{
int deviceCount;
//獲取CUDA設備總數
cudaGetDeviceCount(&deviceCount);
//分別獲取每個CUDA設備的信息
for(int i=0;i<deviceCount;i++)
{
//定義存儲信息的結構體
cudaDeviceProp devProp;
//將第i個CUDA設備的信息寫入結構體中
cudaGetDeviceProperties(&devProp, i);
std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
std::cout << "設備全局內存總量：" << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
std::cout << "SM的數量：" << devProp.multiProcessorCount << std::endl;
std::cout << "每個線程塊的共享內存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
std::cout << "每個線程塊的最大線程數：" << devProp.maxThreadsPerBlock << std::endl;
std::cout << "設備上一個線程塊（Block）中可用的32位寄存器數量： " << devProp.regsPerBlock << std::endl;
std::cout << "每個EM的最大線程數：" << devProp.maxThreadsPerMultiProcessor << std::endl;
std::cout << "每個EM的最大線程束數：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
std::cout << "設備上多處理器的數量：" << devProp.multiProcessorCount << std::endl;
std::cout << "======================================================" << std::endl;
}
}運行以上函數，得到共享內存信息以及其它設備信息，如下圖所示，本人使用的顯卡上，針對於每一個線程塊，其可以使用的最大共享內存為48 KB。
02
—
共享內存的特性
共享內存的主要特點在於“共享”，也即同一個線程塊中的所有線程都可以對這一塊存儲進行讀寫操作，所以“共享”是針對同一個線程塊中所有線程而言的。一旦共享內存被定義並指定大小，系統將給所有線程塊都分配相同大小的共享內存，比如定義一個大小為8 bytes的unsigned char型共享內存，那么所有線程塊都會被分配一個8 bytes的unsigned char型共享內存。如下圖所示：
共享內存在CUDA核函數中定義，通常有兩種方式：靜態方式、動態方式。
靜態方式定義。這種方式定義的特點是定義的同時指定大小：
__global__ shared_memory_kernel(uchar *inputs, int row, int col)
{
int x = threadIdx.x + blockDim.x * blockIdx.x; //col
int y = threadIdx.y + blockDim.y * blockIdx.y; //row
if (x < col && y < row)
{
__shared__ uchar s[8]; //定義的同時指定大小為8 bytes，因此每個線程塊都被分配8 bytes的共享內存
.
.
.
}
}動態方式定義。此方式特點為定義的時候不指定大小，在調用核函數的時候將共享內存大小以輸入參數的形式傳入。
__global__ shared_memory_kernel(uchar *inputs, int row, int col)
{
int x = threadIdx.x + blockDim.x * blockIdx.x; //col
int y = threadIdx.y + blockDim.y * blockIdx.y; //row
if (x < col && y < row)
{
extern __shared__ uchar s[]; //定義的時候不指定大小
.
.
.
}
}
void shared_memory_test(void)
{
.
.
.
//傳入的第1個參數block_num為線程塊總數
//第2個參數thread_num為每個線程塊包含的線程數
//第3個參數8為共享內存大小，所以動態共享內存大小通過第3個參數傳入
shared_memory_kernel<<<block_num, thread_num, 8>>>(inputs, row, col);
.
.
.
}需要注意：動態定義共享內存時，調用核函數傳入的數值必須以字節byte為單位，所以如果定義的共享內存不是byte類型，數值必須乘以類型占用的字節數。比如要動態定義長度為8的float類型共享內存，那么傳入核函數的數值為8*sizeof(float)。
shared_memory_kernel<<<block_num, thread_num, 8 * sizeof(float)>>>(inputs, row, col);03
—
共享內存的應用例子
前文我們講的數組元素規約求和算法，使用CUDA全局內存來存儲數據：
我們知道全局內存屬於板載顯存，而共享內存屬於片上內存，因此共享內存的讀寫速度比全局內存快得多。在前文代碼的核函數中有個for循環需要多次讀寫全局內存，全局內存本身就很慢，而且如果不是連續訪問會更慢，因此本文我們嘗試使用共享內存來代替全局內存實現前文講的規約求和算法。
由於前文的規約算法是在不同線程塊分別進行的，而共享內存又具有線程塊內共享的特性，故共享內存正好適合此應用場景。
前文的規約結構
本文使用共享內存的規約結構
下面我們比較使用共享內存的核函數與前文使用全局內存的核函數：
//使用全局內存
__global__ void cal_sum_ker0(float *Para, float *blocksum_cuda)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < N)
{
for (int index = 1; index < blockDim.x; index = (index*2))
{
if (threadIdx.x % (index*2) == 0)
{
Para[tid] += Para[tid + index];
}
__syncthreads(); //同步，以防止歸約過程中某個線程運行速度過快導致計算錯誤
}
if(threadIdx.x == 0) //整個數組相加完成后，將共享內存數組0號元素的值賦給全局內存數組0號元素
blocksum_cuda[blockIdx.x] = Para[tid];
}
}
//使用共享內存
//blockIdx.x為線程塊的ID號
//blockDim.x每個線程塊中包含的線程總個數
//threadIdx.x為每個線程塊中的線程ID號
__global__ void cal_sum_ker(float *Para, float *blocksum_cuda)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < N)
{
//動態方式定義float型共享內存
extern __shared__ float s_Para[];
//線程塊中的每個線程負責把其對應的數據從全局內存加載到共享內存
s_Para[threadIdx.x] = Para[tid];
__syncthreads(); //塊內線程同步，等待線程塊內所有線程加載數據完畢
for (int index = 1; index < blockDim.x; index = (index*2))
{
if (threadIdx.x % (index*2) == 0)
{
//在for循環中使用共享內存實現規約，避免頻繁讀寫全局內存
s_Para[threadIdx.x] += s_Para[threadIdx.x + index];
}
__syncthreads(); //塊內線程同步，以防止歸約過程中某個線程運行速度過快導致計算錯誤
}
if(threadIdx.x == 0) //將共享內存數組0號元素的值賦給全局內存數組
blocksum_cuda[blockIdx.x] = s_Para[threadIdx.x];
}
}接着在測試代碼中分別調用上方兩個核函數。調用時指定共享內存的長度為每個線程塊包含的線程數：
void cal_sum_test()
{
Timer_Us2 timer;
//定義CPU端數組
float *test_d = (float *)malloc(N * sizeof(float));
for (long long i = 0; i < N; i++)
{
test_d[i] = 0.5;
}
dim3 sumblock(512);//設置每個線程塊有512個線程
dim3 sumgrid(((N%sumblock.x) ? (N/sumblock.x + 1) : (N/sumblock.x)));
float *test_d_cuda;
float *blocksum_cuda;
float *blocksum_host = (float *)malloc(sizeof(float) * sumgrid.x);
cudaMalloc((void **)&test_d_cuda, sizeof(float) * N);
cudaMalloc((void **)&blocksum_cuda, sizeof(float) * sumgrid.x);
timer.start_timer();
//將數據從CPU端拷貝到GPU端
cudaMemcpy(test_d_cuda, test_d, sizeof(float) * N, cudaMemcpyHostToDevice);
//調用使用全局內存規約的核函數
cal_sum_ker0 << < sumgrid, sumblock>> > (test_d_cuda, blocksum_cuda);
//將所有線程塊的規約結果從GPU端拷貝到CPU端
cudaMemcpy(blocksum_host, blocksum_cuda, sizeof(float) * sumgrid.x, cudaMemcpyDeviceToHost);
//在CPU端對所有線程塊的規約求和結果做串行求和
double sum = 0.0;
for(int i = 0; i < sumgrid.x; i++)
{
sum += blocksum_host[i];
}
timer.stop_timer("GPU time (global memory):");
cout << " GPU result (global memory) = " << sum << endl; //顯示GPU端結果
//
timer.start_timer();
cudaMemcpy(test_d_cuda, test_d, sizeof(float) * N, cudaMemcpyHostToDevice);
//調用使用共享內存規約的核函數，sumblock.x為每個線程塊包含的線程數，sumblock.x * sizeof(float)就是傳入的共享內存字節數
cal_sum_ker << < sumgrid, sumblock, sumblock.x * sizeof(float) >> > (test_d_cuda, blocksum_cuda);
cudaMemcpy(blocksum_host, blocksum_cuda, sizeof(float) * sumgrid.x, cudaMemcpyDeviceToHost);
sum = 0.0;
for(int i = 0; i < sumgrid.x; i++)
{
sum += blocksum_host[i];
}
timer.stop_timer("GPU time (shared memory):");
cout << " GPU result (shared memory) = " << sum << endl; //顯示GPU端結果
cudaFree(test_d_cuda);
cudaFree(blocksum_cuda);
free(blocksum_host);
free(test_d);
}運行結果如下，可以看到使用共享內存之后，耗時減少了，這是因為共享內存的讀寫效率比全局內存高。
歡迎掃碼關注本微信公眾號，接下來會不定時更新更加精彩的內容，敬請期待~