5.3 常量内存

常量内存是为多个线程只读类型的广播操作而优化的内存。正如这个名字所传达给我们的一样，编译器使用常量内存存储不能被轻易计算得出的常量或直接被编译为机器码的常量。虽然常量内存驻留在设备内存上，但是机器使用独特的指令来访问。其间，GPU会使用一个特殊的“常量缓存”来访问它。

常量的编译器拥有一个64KB的可用内存，编译器可自由地使用它。开发者同样有另一个64KB的可用内存空间，可以用关键字 constant 声明使用。这一限制被应用在每一个模块（驱动程序API 应用程序）或每一个文件（CUDA运行时应用程序）。

很多人会天真地想，关键字__constant_与C/C++中的关键字const应该是相似的，即在声明后不可更改。但是__constant__关键字声明的内存内容是可以被更改的。通过内存复制或通过查询__constant__内存指针并让内核写入常量内存均可以达到更改常量内存的目的。CUDA内核不应该写入正在被访问的常量内存，因为，内核执行时常量缓存并不是始终与其他内存结构保持一致。

5.3.1 主机与设备常量内存

使用预定义宏__CUDA_arch__来支持主机与设备的常量内存复制，这会方便CPU和GPU对内存的读取。Mark Harris曾描述过这个使用宏的理念[1]。

__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
}

5.3.2 访问常量内存

除了使用C/C++操作符隐式访问常量内存，开发者可以从常量内存复制数据或复制数据到常量内存，甚至可以查询常量内存的指针。

1. CUDA运行时

CUDA运行时应用程序可以使用函数CUDAMemcpyToSymbol()和CUDAMemcpyFrom Symbol()分别复制数据到常量内存和从常量内存复制数据。常量内存的指针可以使用CUDAGetSymbolAddress()函数查询。

cudaError_t cudaGetSymbolAddress(void **devPtr, char *symbol);

这一指针可以用在内核写入常量内存的操作，但开发者必须注意不要在其他内核读取常量内存时写入数据。

2. 驱动程序API

驱动程序API应用程序可以使用函数cuModuleGlobal()查询常量内存的设备指针。由于驱动程序API不包括CUDA运行时中的语言集成特性，驱动程序API不包括像CUDAMemcpyToSymbol()这样的特殊内存复制函数。应用程序必须使用函数cuModuleGetGlobal()查询地址，之后调用cuMemcpyHtoD()或cuMemcpyDtoH()。

内核使用的常量内存数量可以使用函数cuFuncGetAttribute(CU FUNC_ATTRIBUTE_constant_SIZE_BYTES)获取。

[1] http://bit.ly/OpMdN5。