GPU 內存的分級（gpu memory hierarchy）
小普 中科院化學所在讀博士研究生
研究課題，計算機模擬並行軟件的開發與應用
Email: yaopu2019@126.com （歡迎和我討論問題）
摘要(Abstact)
GPU 的存儲是多樣化的, 其速度和數量並不相同，了解GPU存儲對於程序的性能調優有着重要的意義。本文介紹如下幾個問題:
1.內存類型有什么？2)查詢自己設備的內存大小 3）內存訪問速度4）不同級別的存儲關系5）使用注意事項。各種存儲結構的優缺點。
正文
GPU結構圖
①寄存器內存（Register memory）
優點：訪問速度的冠軍！
缺點：數量有限
使用：在__global__函數 ，或者___device__ 函數內，定義的普通變量，就是寄存器變量。
例子：
1 //kernel.cu 2 3 __global__ void register_test() 4 5 { 6 7 int a = 1.0; 8 9 double b = 2.0; 10 11 } 12 13 14 15 //main.cu 16 17 int main() 18 19 { 20 21 int nBlock = 100; 22 23 register_test <<<nBlock,128>>>(); 24 25 return 0; 26 27 } 28 29 30 31
②共享內存（Shared memory）
優點：
1緩存速度快 比全局內存 快2兩個數量級
2 線程塊內，所有線程可以讀寫。
3 生命周期與線程塊同步
缺點：大小有限制
使用：關鍵詞 __shared__ 如 __shared__ double A[128];
適用條件：
使用場合，如規約求和 : a = sum A[i]
如果不是頻繁修改的變量，比如矢量加法。
是編程優化中的重要手段！
C[i] = A[i] + B[i] 則沒有必要將A,B進行緩存到shared memory 中。
1 /kernel.cu 2 3 __global__ void shared_test() 4 5 { 6 7 __shared__ double A[128]; 8 9 int a = 1.0; 10 11 double b = 2.0; 12 13 int tid = threadIdx.x; 14 15 A[tid] = a; 16 17 }
另外一種開辟shared memory 的方式
kernel 函數內，聲明方式
extern __shared__ unsigned int s_out[];
執行 kernel_func<<n_block,block_size,shared_mem_size>>>();
③全局內存 (Global Memory)
優點：
1空間最大（GB級別）
2.可以通過cudaMemcpy 等與Host端，進行交互。
3.生命周期比Kernel函數長
4.所有線程都能訪問
缺點：訪存最慢
1 //kernel.cu 2 3 __global__ void shared_test(int *B) 4 5 { 6 7 double b = 2.0; 8 9 int tid = threadIdx.x; 10 11 int id = blockDim.x*128 + threadIdx.x; 12 13 int a = B[id] ; 14 15 }
④紋理內存
優點，比普通的global memory 快
缺點：使用起來，需要四個步驟，麻煩一點
適用場景：比較大的只需要讀取array，采用紋理方式訪問，會實現加速
使用的四個步驟（這里以1維float數組為例子），初學者，自己手敲一遍代碼！！！
第一步，聲明紋理空間,全局變量：
texture<float, 1, cudaReadModeElementType> tex1D_load;
第二步,綁定紋理
第三步，使用
第四步，解綁定
具體看代碼，（最好自己敲一遍！）
1 #include <iostream> 2 3 #include <time.h> 4 5 #include <assert.h> 6 7 #include <cuda_runtime.h> 8 9 #include "helper_cuda.h" 10 11 #include <iostream> 12 13 #include <ctime> 14 15 #include <stdio.h> 16 17 18 19 using namespace std; 20 21 22 23 texture<float, 1, cudaReadModeElementType> tex1D_load; 24 25 //第一步，聲明紋理空間,全局變量 26 27 28 29 __global__ void kernel(float *d_out, int size) 30 31 { 32 33 //tex1D_load 為全局變量，不在參數表中 34 35 int index; 36 37 index = blockIdx.x * blockDim.x + threadIdx.x; 38 39 if (index < size) 40 41 { 42 43 d_out[index] = tex1Dfetch(tex1D_load, index); //第三步，抓取紋理內存的值 44 45 //從紋理中抓取值 46 47 printf("%f\n", d_out[index]); 48 49 } 50 51 } 52 53 54 55 int main() 56 57 { 58 59 int size = 120; 60 61 size_t Size = size * sizeof(float); 62 63 float *harray; 64 65 float *d_in; 66 67 float *d_out; 68 69 70 71 harray = new float[size]; 72 73 checkCudaErrors(cudaMalloc((void **)&d_out, Size)); 74 75 checkCudaErrors(cudaMalloc((void **)&d_in, Size)); 76 77 78 79 //initial host memory 80 81 82 83 for (int m = 0; m < 4; m++) 84 85 { 86 87 printf("m = %d\n", m); 88 89 for (int loop = 0; loop < size; loop++) 90 91 { 92 93 harray[loop] = loop + m * 1000; 94 95 } 96 97 //拷貝到d_in中 98 99 checkCudaErrors(cudaMemcpy(d_in, harray, Size, cudaMemcpyHostToDevice)); 100 101 102 103 //第二步,綁定紋理 104 105 checkCudaErrors(cudaBindTexture(0, tex1D_load, d_in, Size)); 106 107 //0表示沒有偏移 108 109 110 111 int nBlocks = (Size - 1) / 128 + 1; 112 113 kernel<<<nBlocks, 128>>>(d_out, size); //第三步 114 115 cudaUnbindTexture(tex1D_load); //第四，解紋理 116 117 getLastCudaError("Kernel execution failed"); 118 119 checkCudaErrors(cudaDeviceSynchronize()); 120 121 } 122 123 delete[] harray; 124 125 cudaUnbindTexture(&tex1D_load); 126 127 checkCudaErrors(cudaFree(d_in)); 128 129 checkCudaErrors(cudaFree(d_out)); 130 131 return 0; 132 133 }
總結如下表
要點：
1 在同一個warp內，多線線程訪問一個bank的不同地址，造成confict,影響shared memory 的速度。
2 解決bank confict的方法： padding。
3 const memory 用於存儲固定常量，比如固定的參數等。
結束語
小普 中科院化學所在讀博士研究生
研究課題，計算機模擬並行軟件的開發與應用
Email: yaopu2019@126.com （歡迎和我討論問題，私信和郵件都OK!）
讓程序使得更多人受益！
參考文獻
1) CUDA專家手冊 GPU編程權威指南 [M] 2014
2) CUDA Toolkit Documentation v10.1.168 https://docs.nvidia.com/cuda/