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