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