CUDA 常量記憶體

常量記憶體的大小限定為64k，每個sm擁有8kb的常數儲存器快取，在編譯期時宣告一塊常量記憶體，需要用到__constant__關鍵字，例如

__constant__ float my_array[1024] = ;

不同於c/c++中的const常量，cuda中常量內存在宣告後是可以修改的。如果要在執行時改變常量記憶體中內容，只需要在呼叫gpu核心之前簡單地呼叫cudacopytosymbol函式。如果在編譯階段或主機端執行階段都沒有定義常量記憶體，那麼常量記憶體區將未定義。

使用預定義巨集__cuda_arch__來支援主機與裝置的常量記憶體複製，這會方便cpu和gpu對記憶體的讀取。

__constant__ double dc_vals[2] = ;
const double hc_vals[2] = ;
__device__ __host__ double f(size_t i)

cuda執行時應用程式可以使用函式cudamemecpytosymbol()和cudamemcpyfromsymbol()分別複製資料到常量記憶體和從常量記憶體複製資料。常量記憶體的指標可以使用cudagetsymboladdress()函式查詢。

驅動程式api應用程式可以使用函式cumoduleglobal()查詢常量記憶體的裝置指標。由於驅動程式api不包括cuda執行時的語言整合特性。驅動程式api不包括像cudamemcpytosymbol()這樣的特殊記憶體複製函式。所以必須使用cumodulegetglobal()查詢位址，之後使用cumemcpyhtod()或cumemcpydtoh().

__constant__ char p_hellocuda[11]
;__constant__ int t_hellocuda[11]
=;__constant__ int num =11;
__global__ static
void
hellocuda
(char
* result)
}int
main
(int argc,
char
* ar**)
;cuda_safe_call
(cudamalloc
(void**
)&device_result,
sizeof
(char)*
11);cuda_safe_call
(cudamemcpytosymbol
(p_hellocuda, hellocuda,
sizeof
(char)*
11));
hellocuda<<
<1,
1,0>>
>
(device_result)
;cut_check_error
("kernel excution failed\n");
cuda_safe_call
(cudamemcpy
(&host_result, device_result,
sizeof
(char)*
11, cudamemcpydevicetohost));
;printf
("%s\n"
, host_result)
;cuda_safe_call
(cudafree
(device_result));
cut_exit
(argc, ar**)
;return0;
}

注意到定義常量儲存器時，需要將其定義在所有函式之外，作用範圍為整個檔案，並且對主機端和裝置端函式都可見。同時，上述**中說明了使用常量記憶體的兩種方法。

__constant__ int t_hellocuda[11] = ;
__constant__ int num = 11;

__constant__ char p_hellocuda[11]; // 宣告
...// 使用cudamemcpytosymbol進行賦值
cuda_safe_call(cudamemcpytosymbol(p_hellocuda, hellocuda, sizeof(char) * 11));
...

CUDA 常量記憶體

CUDA學習記憶體處理之常量記憶體（4）

CUDA記憶體使用

CUDA記憶體簡介

CUDA 常量記憶體

CUDA學習 記憶體處理之常量記憶體（4）

CUDA記憶體使用

CUDA記憶體簡介

相關推薦

CUDA學習記憶體處理之常量記憶體（4）