其中闆載顯存主要包括全局内存（global memory）、本地内存（local memory）、常量内存（constant memory）和紋理内存（texture memory），而片上内存主要包括寄存器（register）和共享内存（shared memory）。它們的主要特點如下表所列：

片上内存的讀寫效率通常比闆載顯存更快，而寄存器又是所有顯存類型中最快的存儲器。本文我們将分别介紹這些存儲器的應用場景及用法。

01

寄存器與本地内存

從代碼實現上看，寄存器變量與本地内存變量的定義方式是一樣的，它們都是定義于cuda核函數中的變量，很像C/C 函數中的局部變量定義，比如以下核函數代碼中的變量A、index、tmp都屬于寄存器變量或本地變量：

__global__ void cuda_kernel(float *a, float b, float *c, 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) { float A[200]; int index = y*col x; float tmp = a[index]*b[index] b[index]*b[index]; c[index] = tmp*tmp; } }

那麼定義于CUDA核函數中的變量，什麼時候是寄存器變量，什麼時候是本地變量呢？

通常以下三種情況下定義的變量為本地變量，其餘情況則是寄存器變量：

由以上可知，一般不會在核函數中定義太多或者太大的變量，不然系統自動将超出寄存器限制的變量分配到本地内存，影響程序運行效率。

02

共享内存

共享内存的主要特點在于“共享”，也即同一個線程塊中的所有線程都可以對這一塊存儲進行讀寫操作，所以“共享”是針對同一個線程塊中所有線程而言的。一旦共享内存被定義并指定大小，系統将給所有線程塊都分配相同大小的共享内存，比如定義一個大小為8 bytes的unsigned char型共享内存，那麼所有線程塊都會被分配一個8 bytes的unsigned char型共享内存。

前文我們已經詳細介紹共享内存的特點與應用，此處不再重複：

CUDA加速——共享内存介紹及其應用

03

常量内存

常量内存在device端（GPU端）隻讀，在host端（CPU端）可讀可寫，通常情況下使用__constant__修飾的變量，其數據存儲于常量内存，并且該變量為全局變量，對同一個.cu文件中且定義于其後面的所有核函數都可見。比如以下代碼中，變量A的值存儲于常量内存，kernel1和kernel2都可以使用A，但kernel0因定義在A前面而不能使用A。

__global__ void kernel0() { } __constant__ float A[128]; __global__ void kernel1() { } __global__ void kernel2() { }

常量内存在device端隻讀，因此隻能在host端對其初始化和修改，通過調用cudaMemcpyToSymbol函數實現，比如以下代碼：

__constant__ float A[10]; void init_constant(void) { float B[10] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; //數組A定義在device端的常量内存，數組B定義在host端，使用數組B的内容初始化數組A cudaMemcpyToSymbol(A, B, 10*sizeof(float)); }

什麼情況下使用常量内存呢？

首先我們來講一下warp和half-warp的概念。

從軟件的角度來看，一個block包含的最大線程數通常為512或1024，多個block包含的所有線程都是并行執行的。然而從硬件的角度來看并不是這樣，硬件上把每個block中的所有線程分成每32個一組的線程束，一個線程束稱為一個warp。同一個warp中的線程才是真正意義上的并行執行，并且它們使用各自的數據執行相同的處理指令。不同warp的執行由系統調度，所以不同的warp不一定并行執行。

顧名思義，half-warp就是半個warp，也即同一個warp中的16個線程。

接下來我們介紹使用常量内存的好處，主要有兩個：

由以上可知，在每個half-warp中16個線程需要訪問相同内存地址的情況下，使用常量内存能夠大大提升效率，但是如果每個half-warp中16個線程需要訪問不同的内存地址，這種情況下則不适合使用常量内存。

04

全局内存

全局内存是GPU上容量最大的存儲器，可達到10 GB，所以CUDA編程時通常把較大的數據存儲在全局内存，因此全局内存也是所有GPU存儲器中最常用的存儲器。下面我們分别介紹全局内存的申請、拷貝、使用和釋放。

通過調用cudaMalloc函數可以方便地申請全局内存，不過需要注意該函數第三個參數單位是byte，因此如果數據類型不是char/unsigned char類型，那麼第三個參數需要将數據長度再乘以數據類型所占字節數（sizeof）：

//定義指針 unsigned char *A; int *B; float *C; //定義數據長度 const int data_len = 128; //申請全局内存 cudaMalloc((void**)&A, data_len * sizeof(unsigned char)); cudaMalloc((void**)&B, data_len * sizeof(int)); cudaMalloc((void**)&C, data_len * sizeof(float));

通常調用cudaMemcpy函數将數據從host端内存拷貝到device端全局内存，或者從device端全局内存拷貝到host端内存。同樣，需要注意該函數第三個參數單位也是byte，因此第三個參數需要将數據長度再乘以數據類型所占字節數：

float A[8] = {0, 1, 2, 3, 4, 5, 6, 7}; float *C; cudaMalloc((void**)&C, 8 * sizeof(float)); //将數據從host端内存拷貝到device端全局内存 cudaMemcpy(C, A, 8 * sizeof(float), cudaMemcpyHostToDevice); //将數據從device端全局内存拷貝到host端内存 cudaMemcpy(A, C, 8 * sizeof(float), cudaMemcpyDeviceToHost);

注意以上代碼中決定拷貝方向的是第四個參數，如果是從host拷貝到device，那麼該參數是cudaMemcpyHostToDevice，如果是device到host，那麼該參數是cudaMemcpyDeviceToHost。

一個典型的CUDA并行任務流程是這樣的：

(1) 把數據從host端拷貝到device端（通常是拷貝到device端的全局内存）。

(2) 在device端開啟多線程并行處理數據。

(3) 待開啟的所有線程處理數據完畢，将最後處理結果從device端再拷貝回host端。

所以全局内存通常在CUDA核函數中使用，通常使用線程id号來索引全局内存中的數據，使每個線程與存儲數據的全局内存地址一一對應。下面我們舉一個簡單例子來說明全局内存的使用。

假如有兩個相同尺寸的矩陣A和矩陣B，且A、B都是float型數據矩陣，現在要使用CUDA并行計算A、B相同坐标點數據的平方差。

首先，是CPU實現代碼，循環遍曆所有點計算平方差即可：

void CPU_cal(Mat A, Mat B) { Mat C(A.size(), CV_32FC1); for(int i = 0; i < A.rows; i ) //行遍曆 { float *pA = A.ptr<float>(i); float *pB = B.ptr<float>(i); float *pC = C.ptr<float>(i); for(int j = 0; j < A.cols; j ) //列遍曆 { pC[j] = (pA[j] - pB[j])*(pA[j] - pB[j]); } } }

接着是GPU實現代碼，開啟多線程并行計算每個點的平方差：

/* CUDA核函數 */ __global__ void GPU_cal_kernel(float *A_cuda, float *B_cuda, float *C_cuda, int row, int col) { //線程的x方向id int x = threadIdx.x blockDim.x * blockIdx.x; //col //線程的y方向id int y = threadIdx.y blockDim.y * blockIdx.y; //row if(x < col && y < row) { //将線程的二維id轉換為全局内存的一維地址索引，并保存到寄存器變量index int index = y * col x; //根據一一對應關系，線程(x, y)對應全局内存地址index = y * col x //也即線程(x, y)負責處理全局内存地址index保存的數據 //使用index來索引A_cuda、B_cuda，就相當于從全局内存的index地址讀取數據 //将從A_cuda、B_cuda讀取的index地址數據相減，并把差值保存到寄存器變量diff float diff = A_cuda[index] - B_cuda[index]; //寄存器變量diff保存了差值，因此計算平方的時候可直接使用diff中保存的值 //得到平方值之後，再将結果保存到全局内存C_cuda的index位置 //使用index索引C_cuda，并對其賦值，相當于對全局内存C_cuda的index地址進行寫操作 C_cuda[index] = diff * diff; } } /* 調用以上核函數 */ void GPU_cal(Mat A, Mat B) { float *A_cuda, *B_cuda, *C_cuda; //計算數據長度，注意如果是byte長度還需再乘以sizeof(float) const int data_len = A.rows * A.cols; //申請全局内存 cudaMalloc((void**)&A_cuda, data_len * sizeof(float)); cudaMalloc((void**)&B_cuda, data_len * sizeof(float)); cudaMalloc((void**)&C_cuda, data_len * sizeof(float)); //将數據從host内存拷貝到device全局内存 cudaMemcpy(A_cuda, (float *)A.data, data_len * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(B_cuda, (float *)B.data, data_len * sizeof(float), cudaMemcpyHostToDevice); //定義線程塊、線程塊中的線程都為二維索引 dim3 cuda_Block(16, 16); //每個線程塊有16*16個線程 int M = (A.cols cuda_Block.x - 1) / cuda_Block.x; int N = (A.rows cuda_Block.y - 1) / cuda_Block.y; dim3 cuda_Grid(M, N); //線程網格總共有M*N個線程塊 //調用核函數并行處理 GPU_cal_kernel<<<cuda_Grid, cuda_Block>>>(A_cuda, B_cuda, C_cuda, A.rows, A.cols); Mat C(A.size(), CV_32FC1); //将并行計算結果從device全局内存拷貝到host内存 cudaMemcpy((float *)C.data, C_cuda, data_len * sizeof(float), cudaMemcpyDeviceToHost); //釋放申請的全局内存 cudaFree(A_cuda); cudaFree(B_cuda); cudaFree(C_cuda); }

需注意，對于申請的全局内存，如果不再使用，必須調用cudaFree函數将其釋放。而且以上代碼中線程塊、線程的索引都是二維的，其實它們的索引還可以是一維或三維，以後我們再詳細講怎麼通過線程塊、線程的一維或二維或三維索引來确定其對應的全局内存地址。

此外，針對全局内存的訪問有合并、對齊的說法，隻有在合并、對齊的情況下才能高效地訪問全局内存，在下篇文章我們再詳細探讨這個話題。

05

紋理内存

紋理内存是GPU中的一種隻讀存儲器，其使用方式為将某一段全局内存綁定到紋理内存，這段全局内存通常的表現形式為一維CUDA數組/全局内存、二維或三維CUDA數組，然後通過讀取紋理内存（也稱為紋理拾取）來獲取全局内存的數據。相比全局内存的訪問要求對齊、合并，紋理内存對非對齊訪問和随機訪問具有良好的加速效果。

一維紋理可以綁定到CUDA數組，也可以直接綁定到全局内存。下面舉一個簡單的例子來介紹一維紋理内存的使用。

首先是一維紋理的定義，紋理内存通常定義為全局變量：

//float表示數據類型 //cudaTextureType1D、cudaTextureType2D、cudaTextureType3D分别表示一維、二維、三維 //cudaReadModeElementType表示隻讀模式 //tex_1D為定義的紋理内存變量 texture<float, cudaTextureType1D, cudaReadModeElementType> tex_1D;

其次，是紋理内存的綁定，這裡我們直接把全局内存綁定到紋理：

const int data_len = 128; const int data_size = data_len * sizeof(float); //初始化host端數組 float *data_host = (float *)malloc(data_size); for(int i = 0; i < data_len; i ) { data_host[i] = i; } //申請device端全局内存 float *data_device; cudaMalloc((void**)&data_device, data_size); //将數據從host端拷貝到device端全局内存 cudaMemcpy(data_device, data_host, data_size, cudaMemcpyHostToDevice); //參數一表示以bytes為單位的偏移量，也即綁定到紋理的全局内存的起始偏移地址 //參數二為紋理内存變量 //參數三為全局内存地址變量 cudaBindTexture(0, tex_1D, data_device);

接着是在核函數中紋理拾取，通過調用tex1Dfetch函數實現：

//功能：将紋理内存的數據拷貝到全局内存A_cuda __global__ void cuda_kernel(float *A_cuda, int data_len) { //線程id int x = threadIdx.x blockDim.x * blockIdx.x; if(x < data_len) { //使用線程id來索引全局内存A_cuda和紋理内存tex_1D A_cuda[x] = tex1Dfetch(tex_1D, x); } }

最後是紋理内存的解綁，函數執行完畢之後需要對紋理内存進行解綁：

cudaUnbindTexture(tex_1D);

完整代碼：

texture<float, cudaTextureType1D, cudaReadModeElementType> tex_1D; //功能：将紋理内存的數據拷貝到全局内存A_cuda __global__ void cuda_kernel(float *A_cuda, int data_len) { //線程id int x = threadIdx.x blockDim.x * blockIdx.x; if(x < data_len) { //使用線程id來索引全局内存A_cuda和紋理内存tex_1D A_cuda[x] = tex1Dfetch(tex_1D, x); } } void cuda_copy_data(float *data_host_dst) { const int data_len = 128; const int data_size = data_len * sizeof(float); //初始化host端數組 float *data_host = (float *)malloc(data_size); for(int i = 0; i < data_len; i ) { data_host[i] = i; } //申請device端全局内存 float *data_device, *data_dst; cudaMalloc((void**)&data_device, data_size); cudaMalloc((void**)&data_dst, data_size); //将數據從host端拷貝到device端全局内存 cudaMemcpy(data_device, data_host, data_size, cudaMemcpyHostToDevice); //參數一表示以bytes為單位的偏移量，也即綁定到紋理的全局内存的起始偏移地址 //參數二為紋理内存變量 //參數三為全局内存地址變量 cudaBindTexture(0, tex_1D, data_device); dim3 tex_Block(16); //每個block有16個線程 //總共有((data_len 15) / 16)個block dim3 tex_Grid((data_len tex_Block.x - 1) / tex_Block.x); //調用核函數 cuda_kernel<<<tex_Grid, tex_Block>>>(data_dst, data_len); //将數據從device端拷貝到host端 cudaMemcpy(data_host_dst, data_dst, data_size, cudaMemcpyDeviceToHost); //紋理解綁 cudaUnbindTexture(tex_1D); //釋放全局内存 cudaFree(data_device); cudaFree(data_dst); //釋放host内存 free(data_host); }

通常将存儲一張二維圖像的全局内存綁定到二維紋理，在核函數中可高效地随機訪問二維紋理數據。

以全局變量的方式定義二維紋理：

texture<float, cudaTextureType2D, cudaReadModeElementType> tex_src;

定義CUDA數組，并申請CUDA數組内存：

//聲明數據類型為float cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); //定義CUDA數組 cudaArray *cuArray_src; //分配大小為col_c*row_c的CUDA數組 int col_c = 512; int row_c = 512; cudaMallocArray(&cuArray_src, &channelDesc, col_c, row_c);

設置紋理内存參數，并将CUDA數組綁定到紋理内存：

//尋址方式 //cudaAddressModeWrap--循環尋址，如果超出最大地址則轉成從最小地址開始 //cudaAddressModeClamp--鉗位尋址，如果超出最大地址則訪問最大地址 tex_src1.addressMode[0] = cudaAddressModeWrap; tex_src1.addressMode[1] = cudaAddressModeWrap; //是否對紋理坐标歸一化 tex_src1.normalized = false; //紋理的濾波模式： //cudaFilterModePoint--最鄰近插值 //cudaFilterModeLinear--雙線性插值 tex_src1.filterMode = cudaFilterModePoint; //紋理綁定，将CUDA數組綁定到紋理tex_src cudaBindTextureToArray(&tex_src, cuArray_src, &channelDesc);

将數據從host内存拷貝到CUDA數組：

Mat M_float = Mat::zeros(row_c, col_c, CV_32FC1); cudaMemcpyToArray(cuArray_src, 0, 0, (float *)M_float.data, row_c*col_c*sizeof(float), cudaMemcpyHostToDevice);

在CUDA核函數中調用tex2D函數進行紋理拾取：

//tex_src--要拾取的紋理内存 //x--紋理内存的x坐标 //y--紋理内存的y坐标 //功能：将紋理内存中(x,y)坐标處的數據加載到寄存器變量d float d = tex2D(tex_src, x, y);

紋理解綁并釋放CUDA數組：

cudaUnbindTexture(tex_src); cudaFreeArray(cuArray_src);

有時候需要使用CUDA處理多幀的圖像，把多幀時間序列的圖像傳入到GPU中，此時就可以把保存圖像的全局内存綁定到二維紋理内存（将每幀圖像展開拼接為一行），核函數通過紋理拾取來訪問輸入的圖像數據。不過二維紋理内存的寬是有限制的：

cudaMallocArray函數的第三個參數img_size為寬，也即每幀圖像的總數據個數。然而二維紋理對寬是有限制的，如果寬超過64K就會出錯。所以如果每幀圖像的大小超過了64K，則不能使用二維紋理内存，這時候可以使用三維紋理内存。

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); cudaMallocArray((cudaArray**)&arr_mat_x, &channelDesc, img_size, IIR_N_X);

三維紋理的使用與一維、二維紋理區别較大，下面将詳細說明三維紋理内存的使用。

以全局變量的方式定義三維紋理：

texture<float, cudaTextureType3D, cudaReadModeElementType> tex_mat; //定義為3D類型的紋理内存

定義三維的CUDA數組：

cudaArray *arr_mat; //定義CUDA數組 cudaExtent extent; //定義圖像的尺寸和幀數結構體 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); //定義數據類型為float //創建extent時，在舊版本中寬度w以字節為單位，即必須乘上sizeof(DTYPE)，新版本已經不用乘以sizeof(DTYPE)，否則反而會出錯!文檔和函數說明都沒有改過來，這裡是坑！ extent.width = col; //每幀圖像的列數，這裡不需要再乘以sizeof(float) extent.height = row; //每幀圖像的行數 extent.depth = picnum; //圖像的總幀數 //創建picnum幀row*col的存儲空間 cudaMalloc3DArray((cudaArray**)&arr_mat, &channelDesc, extent);

将數據從host内存拷貝到CUDA數組中，其中pic為Mat類型的vector數組，其包含了picnum幀圖像：

vector<Mat> pic;

cudaMemcpy3DParms HostToDev = {0}; //定義數據傳輸的結構體 HostToDev.dstArray = arr_mat; //指定數據傳輸的目标地址為cuda數組 HostToDev.extent = make_cudaExtent(col, row, 1); //創建extent時，在舊版本中寬度w以字節為單位，即必須乘上sizeof(DTYPE)，新版本已經不用乘以sizeof(DTYPE)，否則反而會出錯!文檔和函數說明都沒有改過來，這裡是坑！ HostToDev.kind = cudaMemcpyHostToDevice; //定義傳輸方向為CPU到GPU顯存 HostToDev.srcPos = make_cudaPos(0, 0, 0); //定義數據傳輸的源地址的偏移量(w, h, img_index) for(int i = 0; i < picnum; i ) //拷貝多幀圖像到cuda數組 { //指定數據傳輸的源地址，注意這裡的第二個參數需要乘以數據類型所占的字節數 HostToDev.srcPtr = make_cudaPitchedPtr((void *)pic[i].data, col*sizeof(float), col, row); HostToDev.dstPos = make_cudaPos(0, 0, i); //指定目标地址的偏移量，分别為x，y，z地址 cudaMemcpy3D(&HostToDev); //根據以上設置的參數實行拷貝 }

設置三維紋理參數，并将CUDA數組綁定到三維紋理：

tex_mat.normalized = 0; //索引地址不歸一化 //filterMode：濾波模式。僅對綁定 CUDA 數組的紋理有效。當使用浮點型的坐标尋址紋理時，将根據設定返回不同類型的值。設定可以有：cudaFilterModePoint 和 cudaFilterModeLinear。分别表示最近鄰插值和線性插值 tex_mat.filterMode = cudaFilterModePoint; tex_mat.addressMode[0] = cudaAddressModeClamp; //尋址模式，即如何處理越界的紋理坐标。可設置：cudaAddressModeClamp 和 cudaAddressModeWrap。Clamp 即鉗位模式，Wrap 為循環模式。循環模式隻支持歸一化的紋理坐标 tex_mat.addressMode[1] = cudaAddressModeClamp; tex_mat.addressMode[2] = cudaAddressModeClamp; tex_mat.channelDesc = channelDesc; //描述紋理返回值類型，同cuda數組部分的内容 cudaBindTextureToArray(tex_mat, (cudaArray *)arr_mat, channelDesc); //綁定紋理内存

在核函數中調用tex3D函數執行紋理拾取：

//後面三個參數分别是x，y，z坐标 tex3D(tex_mat, x, y, z);

最後是紋理解綁和釋放CUDA數組（這個與二維紋理一樣）：

cudaUnbindTexture(tex_mat); cudaFreeArray(arr_mat);

紋理内存具有硬件插值功能，包括最鄰近插值和雙線性插值這兩種插值方式。如果紋理拾取時輸入的訪問坐标地址是浮點數，紋理内存将自動根據設置插值方式對浮點坐标進行插值，然後返回插值結果。這個插值過程不需要開發者來實現，是硬件自動完成的，開發者隻需要設置好插值方式為最鄰近插值或者雙線性插值即可，因此可以節省很多計算時間。

這個我們前文已經介紹，詳細請參考：

CUDA紋理内存--硬件插值功能的應用

更多精彩资讯请关注tft每日頭條，我们将持续为您更新最新资讯!