CUDA 顯示卡 GPU memory
CUDA儲存器型別：
每個執行緒擁有自己的register and loacal memory;
每個執行緒塊擁有一塊shared memory;
所有執行緒都可以訪問global memory;
還有，可以被所有執行緒訪問的只讀儲存器：constant memory and texture memory
1、 暫存器Register
暫存器是GPU上的快取記憶體器，其基本單元是暫存器檔案，每個暫存器檔案大小為32bit.
Kernel中的區域性(簡單型別)變數第一選擇是被分配到Register中。
特點：每個執行緒私有，速度快。
2、 區域性儲存器 local memory
當register耗盡時，資料將被儲存到local memory。如果每個執行緒中使用了過多的暫存器，或聲明瞭大型結構體或陣列，或編譯器無法確定陣列大小，執行緒的私有資料就會被分配到local memory中。
特點：每個執行緒私有；沒有快取，慢。
注：在宣告區域性變數時，儘量使變數可以分配到register。如：
unsigned int mt[3];
改為： unsigned int mt0, mt1, mt2;
3、 共享儲存器 shared memory
可以被同一block中的所有執行緒讀寫
特點：block中的執行緒共有；訪問共享儲存器幾乎與register一樣快.
-
//u(i)= u(i)^2 + u(i-1) -
//Static -
__global__ example(float* u) { -
int i=threadIdx.x; -
__shared__ int tmp[4]; -
tmp[i]=u[i]; -
u[i]=tmp[i]*tmp[i]+tmp[3-i]; -
} -
int main() { -
float hostU[4] = {1, 2, 3, 4}; -
float* devU; -
size_t size = sizeof(float)*4; -
cudaMalloc(&devU, size); -
cudaMemcpy(devU, hostU, size, -
cudaMemcpyHostToDevice); -
example<<<1,4>>>(devU, devV); -
cudaMemcpy(hostU, devU, size, -
cudaMemcpyDeviceToHost); -
cudaFree(devU); -
return 0; -
} -
//Dynamic -
extern __shared__ int tmp[]; -
__global__ example(float* u) { -
int i=threadIdx.x; -
tmp[i]=u[i]; -
u[i]=tmp[i]*tmp[i]+tmp[3-i]; -
} -
int main() { -
float hostU[4] = {1, 2, 3, 4}; -
float* devU; -
size_t size = sizeof(float)*4; -
cudaMalloc(&devU, size); -
cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice); -
example<<<1,4,size>>>(devU, devV); -
cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost); -
cudaFree(devU); -
return 0; -
}
4、 全域性儲存器 global memory
特點：所有執行緒都可以訪問；沒有快取
-
//Dynamic -
__global__ add4f(float* u, float* v) { -
int i=threadIdx.x; -
u[i]+=v[i]; -
} -
int main() { -
float hostU[4] = {1, 2, 3, 4}; -
float hostV[4] = {1, 2, 3, 4}; -
float* devU, devV; -
size_t size = sizeof(float)*4; -
cudaMalloc(&devU, size); -
cudaMalloc(&devV, size); -
cudaMemcpy(devU, hostU, size, -
cudaMemcpyHostToDevice); -
cudaMemcpy(devV, hostV, size, -
cudaMemcpyHostToDevice); -
add4f<<<1,4>>>(devU, devV); -
cudaMemcpy(hostU, devU, size, -
cudaMemcpyDeviceToHost); -
cudaFree(devV); -
cudaFree(devU); -
return 0; -
} -
//static -
__device__ float devU[4]; -
__device__ float devV[4]; -
__global__ addUV() { -
int i=threadIdx.x; -
devU[i]+=devV[i]; -
} -
int main() { -
float hostU[4] = {1, 2, 3, 4}; -
float hostV[4] = {1, 2, 3, 4}; -
size_t size = sizeof(float)*4; -
cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice); -
cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice); -
addUV<<<1,4>>>(); -
cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost); -
return 0; -
}
5、 常數儲存器constant memory
用於儲存訪問頻繁的只讀引數
特點：只讀；有快取；空間小(64KB)
注：定義常數儲存器時，需要將其定義在所有函式之外，作用於整個檔案
1 __constant__ int devVar; 2 cudaMemcpyToSymbol(devVar, hostVar, sizeof(int), 0, cudaMemcpyHostToDevice) 3 cudaMemcpyFromSymbol(hostVar, devVar, sizeof(int), 0, cudaMemcpyDeviceToHost)
6、 紋理儲存器 texture memory
是一種只讀儲存器，其中的資料以一維、二維或者三維陣列的形式儲存在視訊記憶體中。在通用計算中，其適合實現影象處理和查詢，對大量資料的隨機訪問和非對齊訪問也有良好的加速效果。
特點：具有紋理快取，只讀。
TNE END
--------------------- 本文來自 Augusdi 的CSDN 部落格 ，全文地址請點選：https://blog.csdn.net/augusdi/article/details/12186939?utm_source=copy
=============================================================================================
CUDA儲存器模型：
GPU片內：register，shared memory；
板載視訊記憶體：local memory,constant memory, texture memory, texture memory,global memory;
host 記憶體： host memory, pinned memory.
register: 訪問延遲極低；
基本單元：register file （32bit/each）
計算能力1.0/1.1版本硬體：8192/SM；
計算能力1.2/1.3版本硬體： 16384/SM;
每個執行緒佔有的register有限，程式設計時不要為其分配過多私有變數；
local memory：暫存器被使用完畢，資料將被儲存在區域性儲存器中；
大型結構體或者陣列；
無法確定大小的陣列；
執行緒的輸入和中間變數；
定義執行緒私有陣列的同時進行初始化的陣列被分配在暫存器中；
shared memory：訪問速度與暫存器相似；
實現執行緒間通訊的延遲最小；
儲存公用的計數器或者block的公用結果；
硬體1.0~1.3中，16KByte/SM,被組織為16個bank；
宣告關鍵字 _shared_ int sdata_static[16];
global memory:存在於視訊記憶體中，也稱為線性記憶體（視訊記憶體可以被定義為線性儲存器或者CUDA陣列）；
cudaMalloc（）函式分配，cudaFree（）函式釋放，cudaMemcpy（）進行主機端與裝置端的資料傳輸；
初始化共享儲存器需要呼叫cudaMemset（）；
二維三維陣列：cudaMallocPitch（）和cudaMalloc3D（）分配線性儲存空間，可以確保分配滿足對齊要求；
cudaMemcpy2D（），cudaMemcpy3D（）與裝置端儲存器進行拷貝；
host記憶體：分為pageable memory 和 pinned memory
pageable memory： 通過作業系統API（malloc（），new（））分配的儲存器空間；、
pinned memory：始終存在於實體記憶體中，不會被分配到低速的虛擬記憶體中，能夠通過DMA加速與裝置端進行通訊；
cudaHostAlloc(), cudaFreeHost()來分配和釋放pinned memory；
使用pinned memory優點：主機端-裝置端的資料傳輸頻寬高；
某些裝置上可以通過zero-copy功能對映到裝置地址空間，從GPU直接訪問，省掉主存與視訊記憶體間進行資料拷貝的工作；
pinned memory 不可以分配過多：導致作業系統用於分頁的實體記憶體變， 導致系統整體效能下降；通常由哪個cpu執行緒分配，就只有這個執行緒才有訪問許可權；
cuda2.3版本中，pinned memory功能擴充：
portable memory：讓控制不同GPU的主機端執行緒操作同一塊portable memory，實現cpu執行緒間通訊；使用cudaHostAlloc（）分配頁鎖定記憶體時，加上cudaHostAllocPortable標誌；
write-combined Memory：提高從cpu向GPU單向傳輸資料的速度；不使用cpu的L1，L2 cache對一塊pinned memory中的資料進行緩衝，將cache資源留給其他程式使用；在pci-e匯流排傳輸期間不會被來自cpu的監視打斷；在呼叫cudaHostAlloc（）時加上cudaHostAllocWriteCombined標誌；cpu從這種儲存器上讀取的速度很低；
mapped memory：兩個地址：主機端地址（記憶體地址），裝置端地址（視訊記憶體地址）。 可以在kernnel程式中直接訪問mapped memory中的資料，不必在記憶體和視訊記憶體之間進行資料拷貝，即zero-copy功能；在主機端可以由cudaHostAlloc（）函式獲得，在裝置端指標可以通過cudaHostGetDevicePointer（）獲得；通過cudaGetDeviceProperties（）函式返回的canMapHostMemory屬性知道裝置是否支援mapped memory；在呼叫cudaHostAlloc（）時加上cudaHostMapped標誌，將pinned memory對映到裝置地址空間；必須使用同步來保證cpu和GPu對同一塊儲存器操作的順序一致性；視訊記憶體中的一部分可以既是portable memory又是mapped memory；在執行CUDA操作前，先呼叫cudaSetDeviceFlags（）（加cudaDeviceMapHost標誌）進行頁鎖定記憶體對映。
constant memory：只讀地址空間；位於視訊記憶體，有快取加速；64Kb；用於儲存需要頻繁訪問的只讀引數 ；只讀；使用_constant_ 關鍵字，定義在所有函式之外；兩種常數儲存器的使用方法：直接在定義時初始化常數儲存器；定義一個constant陣列，然後使用函式進行賦值；
texture memory：只讀；不是一塊專門的儲存器，而是牽涉到視訊記憶體、兩級紋理快取、紋理拾取單元的紋理流水線；資料常以一維、二維或者三維陣列的形式儲存在視訊記憶體中；快取加速；可以宣告大小比常數儲存器大得多；適合實現影象樹立和查詢表；對大量資料的隨機訪問或非對齊訪問有良好的加速效果；在kernel中訪問紋理儲存器的操作成為紋理拾取（texture fetching）；紋理拾取使用的座標與資料在視訊記憶體中的位置可以不同，通過紋理參照系約定二者的對映方式；將視訊記憶體中的資料與紋理參照系關聯的操作，稱為將資料與紋理繫結（texture binding）；視訊記憶體中可以繫結到紋理的資料有：普通線性儲存器和cuda陣列；存在快取機制；可以設定濾波模式，定址模式等；
--------------------- 本文來自 caoeryingzi 的CSDN 部落格 ，全文地址請點選：https://blog.csdn.net/caoeryingzi/article/details/21323475?utm_source=copy