為了更好的在kernel函數中提高資料通路效率，我們需要了解CUDA各個存儲器模型的特點，合理的配置設定存儲空間。
-------------------華麗的分割線-------------------
一、存儲器比較
每個線程擁有自己的register and loacal memory;(可讀可寫)
每個線程塊擁有一塊shared memory;(可讀可寫)
所有線程都可以通路global memory;(可讀可寫)
還有，可以被所有線程通路的隻讀存儲器：constant memory and texture memory （隻讀）
通過上述圖，我們了解為什麼shared和register擁有同樣的通路速度，shared memory 和 register 位于 GPU 片内。全局存儲三個：global memory,constant memory,texture memory.
存儲器關系圖，如圖所示：
二、各類存儲器細講
1、 寄存器Register
寄存器是GPU上的高速緩存器，其基本單元是寄存器檔案，每個寄存器檔案大小為32bit. Kernel中的局部(簡單類型)變量第一選擇是被配置設定到Register中。
特點：每個線程私有，速度快。
2、 局部存儲器 local memory
當register耗盡時，資料将被存儲到local memory。如果每個線程中使用了過多的寄存器，或聲明了大型結構體或數組，或編譯器無法确定數組大小，線程的私有資料就會被配置設定到local memory中。但是local memory 的資料是被儲存在顯存中的，速度很慢。
特點：每個線程私有；沒有緩存，慢。
注：在聲明局部變量時，盡量使變量可以配置設定到register。如：
unsigned int mt[3]; 改為： unsigned int mt0, mt1, mt2;
3、共享記憶體 共享存儲器也是GPU片内的高速存儲器，通路sharedmemory幾乎和通路register一樣快，是實作線程間通信的延遲最小的方法。shared memory是一塊可以被同一block中的所有thread通路的可讀寫存儲器。每個塊配置設定48KB 。對于GPU上啟動的每個線程塊，CUDA C編譯器都将建立該共享變量的一個副本。線程塊中的每個線程都共享這塊記憶體，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個線程塊中的多個線程能夠在計算上通信和協作。 形式：關鍵字 __shared__添加到變量聲明中。如__shared__ float cache[10]。
4、全局記憶體 global memory
通俗意義上的裝置記憶體。所有線程都可以通路；沒有緩存 顯存中的全局存儲器也稱為線性記憶體，線性記憶體通常使用cudaMalloc 配置設定記憶體
5、常量記憶體 constant memory
常量記憶體用于儲存在核函數執行期間不會發生變化的資料。變量的通路限制為隻讀。NVIDIA硬體提供了 64KB的常量記憶體。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜态地配置設定空間。為了提升性能。常量記憶體采取了不同于标準全局記憶體的處理方式。在某些情況下，用常量記憶體替換全局記憶體能有效地減少記憶體帶寬。當我們需要拷貝資料到常量記憶體中應該使用 cudaMemcpyToSymbol()，而 cudaMemcpy()會複制到全局記憶體。
形式：關鍵字 __constant__添加到變量聲明中。如__constant__ float s[10]. 特點：隻讀；有緩存；空間小(64KB)
使用常量記憶體性能提升的原因：
I：對常量記憶體的單次讀操作可以廣播到其他的“鄰近”線程。這将節約15次讀取操作。（為什麼是15，因為“鄰近”指半個線程束，一個線程束包含32個線程的集合。）
II： 常量記憶體的資料将緩存起來，是以對相同位址的連續讀操作将不會産生額外的記憶體通信量。
注：定義常數存儲器時，需要将其定義在所有函數之外，作用于整個檔案
1 __constant__ int devVar;
2 cudaMemcpyToSymbol(devVar, hostVar, sizeof(int), 0, cudaMemcpyHostToDevice)
3 cudaMemcpyFromSymbol(hostVar, devVar, sizeof(int), 0, cudaMemcpyDeviceToHost) 例子：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define SIZE 124
//***declare the constant value***
//兩種方法指派：1-定義直接初始化 2-定義後調用函數初始化cudaMemcpyToSymbol
__constant__ int constArray[SIZE];
__constant__ int cosntArray_Result[SIZE];
__constant__ int constNumber = 10;
//********************************
__global__ void addKernel()
{
int idx = threadIdx.x + blockDim.x * blockIdx.x;
cosntArray_Result[idx] = constNumber + constArray[idx];
}
int main()
{
int inArray[SIZE] = {0};
size_t size = SIZE * sizeof(int);
int *result = NULL;
result = (int*)malloc(size);
//初始化資料
for (int i = 0; i < SIZE; i++) {
inArray[i] = i+1;
}
//copy host data to constant memory 如果不初始化常量記憶體，将預設數組各元素為0
cudaMemcpyToSymbol(constArray, inArray, size);
//call kernel fun
dim3 threadPerBlock(16);
dim3 blockNum((SIZE + threadPerBlock.x - 1) / threadPerBlock.x);
//計算輸入資料和輸出資料均采用常量記憶體
addKernel << <blockNum , threadPerBlock >> > ();
//copy constant data to host data
cudaMemcpyFromSymbol(result, cosntArray_Result, size);
//show data
for (int i = 0; i < SIZE; i++) {
printf("%5d", result[i]);
}
return 0;
}
結果：
6、紋理記憶體 和常量記憶體一樣，紋理記憶體是另一種類型的隻讀記憶體，在特定的通路模式中，紋理記憶體同樣能夠提升性能并減少記憶體流量。。紋理緩存是專門為那些在記憶體通路模式中存在大量空間局部性(Spatial Locality)的圖形應用程式而設計的。意味着一個線程讀取的位置可能與鄰近線程讀取的位置“非常接近”。如下圖：
紋理變量（引用）必須聲明為檔案作用域内的全局變量。
形式：分為一維紋理記憶體：texture<float> texconst 和 二維紋理記憶體 texture<float> texconst 。