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));
...