CUDA記憶體型別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 main() ;
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 main() ;
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 main() ;
float hostv[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 main() ;
float hostv[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

CUDA記憶體型別memory

CUDA記憶體型別memory

CUDA記憶體型別memory

3 記憶體 Memory 效能監控

相關推薦