CUDA_常量內存

常量內存簡介

常量內存其實只是全局內存的一種虛擬地址形式，並無特殊保留的常量內存塊。常量內存有兩個特性：

• 高速緩存,擁有緩存加速；
• 支持將單個值廣播到線程束中的每一個線程

常量內存的大小限定爲64K，每一個SM擁有8KB的常數存儲器緩存，在編譯期時聲明一塊常量內存，須要用到__constant__關鍵字，例如

__constant__ float my_array[1024] = { 0.0F, 1.0F, 1.3F, ...};

不一樣於c/c++中的const常量，cuda中常量內存在聲明後是能夠修改的。若是要在運行時改變常量內存中內容，只須要在調用GPU內核以前簡單地調用cudaCopyToSymbol函數。若是在編譯階段或主機端運行階段都沒有定義常量內存，那麼常量內存區將未定義。

主機與設備常量內存

使用預約義宏__CUDA_ARCH__來支持主機與設備的常量內存複製，這會方便CPU和GPU對內存的讀取。

__constant__ double dc_vals[2] = {0.0, 1000.0};
       const double hc_vals[2] = {0.0, 1000.0};
      
__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
	return dc_vals[i];
#else
	return hc_vals[i];
#endif
}

訪問常量內存

cuda運行時api

cuda運行時應用程序可使用函數cudaMemecpyToSymbol()和cudaMemcpyFromSymbol()分別複製數據到常量內存和從常量內存複製數據。常量內存的指針可使用cudaGetSymbolAddress()函數查詢。

驅動程序api

驅動程序api應用程序可使用函數cuModuleGlobal()查詢常量內存的設備指針。因爲驅動程序api不包括cuda運行時的語言集成特性。驅動程序api不包括像cudaMemcpyToSymbol()這樣的特殊內存複製函數。因此必須使用cuModuleGetGlobal()查詢地址，以後使用cuMemcpyHtoD()或cuMemcpyDtoH().

常量存儲器使用案例

__constant__ char p_HelloCUDA[11];
__constant__ int t_HelloCUDA[11] = {0,1,2,3,4,5,6,7,8,9,10};
__constant__ int num = 11;

__global__ static void HelloCUDA(char * result)
{
    int i = 0;
    for(i=0; i<num; i++)
    {
        result[i] = p_HelloCUDA[i] + t_HelloCUDA[i];
    }
}

int main(int argc, char* argv[])
{
    if(!InitCUDA())
        return 0;
    
    char helloCUDA[] = "hdjik CUDA!";
    char *device_result = 0;
    char host_result[12] = {0};

    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, argv);
    return 0;
}

注意到定義常量存儲器時，須要將其定義在全部函數以外，做用範圍爲整個文件，而且對主機端和設備端函數均可見。同時，上述代碼中說明了使用常量內存的兩種方法。

• 在定義時直接初始化常數寄存器，而後在kernel裏面直接使用。

__constant__ int t_HelloCUDA[11] = {0,1,2,3,4,5,6,7,8,9,10};
__constant__ int num = 11;

• 定義一個constant數組，(先聲明)，而後使用函數進行賦值。

__constant__ char p_HelloCUDA[11]; // 聲明
...
// 使用cudaMemcpyToSymbol進行賦值
CUDA_SAFE_CALL(cudaMemcpyToSymbol(p_HelloCUDA, helloCUDA, sizeof(char) * 11));
...