【CUDA】常量内存

一、认识常量内存

二、使用常量内存实现1D模板

三、与只读缓存比较

一、认识常量内存

常量内存对内核代码而言是只读的，但它对主机而言即是可读的又是可写的。常量内存位于设备的DRAM上（和全局内存一样）。有一个专用的片上缓存，从每个SM的常量缓存中读取的延迟，比直接从常量内存中读取要低得多。每个SM常量内存缓存大小的限制为64KB

constant Memory 的获取方式不同于其它的 GPU 内存，对于 constant Memory 来说，最佳获取方式是 warp 中的 32 个thread获取 constant Memory 中的同一个地址。若 warp 中的线程访问不同的地址，那么访问就需要串行

一个常量内存读取的成本与 warp 中 thread 读取唯一地址的数量呈线性关系。在全局作用域中必须用以下修饰符声明常量内存：__constant__。常量内存变量的生存期与应用程序的生存期相同，其对于网格内的所有线程都是可访问的，并且主机可以通过运行时函数访问

当使用 CUDA 独立编译能力时，常量内存变量跨多个源文件可见。因为设备只能读取常量内存，所以常量内存中的值必须使用以下运行时函数进行初始化：

cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset, cudaMemcpyKind kind);

cudaMemcpyToSymbol 函数将 src 指向的数据复制到设备上有 symbol 指定的常量内存中。枚举变量 kind 指定了传输方向，默认情况下，kind是 cudaMemcoyHostToDevice

二、使用常量内存实现1D模板

在一维中，在位置 x 周围九点模板会给这些位置上的值应用一些函数：{x-4h，x-3h，x-2h，x-h，x，x+h，x+2h，x+3h，x+4h}

一个九点模板的例子是实变量函数 f 在点 x 上一阶导数的第八阶中心差分公式。裂解这个公式的应用并不重要，只要简单的了解到其会将上述的九点作为输入并产生单一输出。接下来将这个公式作为一个示例模板：

那么要放到 constant Memory 中的便是其中的 c0、c1、c2 ... ...

const int RADIUS = 4;  // 半径, 一个数左右两边各4个数, 组成9个数 
const int LENGTH = 6; // 共有多少有效数字// 在常量内存中声明coef数组
__constant__ float coef[RADIUS + 1];void set_coef_constant()
{const float host_coef[] = {0, 1, 2, 3, 4};cudaMemcpyToSymbol(coef, host_coef, ((RADIUS + 1) * sizeof(float)));
}__global__ 
void stencil_1d(float* input, float* output)
{// 因为每个线程需要 9 个点来计算一个点，所以使用共享内存来优化缓存数据，从而减少对全局内存的冗余访问__shared__ float data[LENGTH + 2 * RADIUS];// 访问全局内存的线程索引int globalThreadIndex = blockIdx.x * blockDim.x + threadIdx.x; // 0 到 5// 访问共享内存的线程索引int sharedThreadIndex = threadIdx.x + RADIUS; // 4 到 9// gloabl -> shareddata[sharedThreadIndex] = input[globalThreadIndex + RADIUS];// 前四个线程负责 读取左右边界if(threadIdx.x < RADIUS) {data[sharedThreadIndex - RADIUS] = input[sharedThreadIndex - RADIUS];data[LENGTH + sharedThreadIndex] = input[LENGTH + sharedThreadIndex];}__syncthreads();float tmp = 0;// 提示 CUDA 编译器，表明这个循环将自动展开#pragma unrollfor(int i = 1; i < RADIUS; ++i)tmp += coef[i] * (data[sharedThreadIndex - RADIUS] - data[sharedThreadIndex + RADIUS]);output[globalThreadIndex] = tmp;
}void test(float* input, float* output)
{stencil_1d<<<1, 6>>>(input, output);
}

三、与只读缓存比较

Kepler 系列的GPU允许使用 texture pipeline 作为一个 global Memory 只读缓存。因为这是一个独立的使用单独带宽的只读缓存，所以对带宽限制的kernel性能有很大的提升

每个 Kepler SM 都有48KB的只读缓存。一般来说，只读缓存在分散读取方面比一级缓存更好，当线程束的线程都读取相同地址时，不应使用只读缓存。只读缓存的粒度为32个字节

当通过只读缓存访问全局内存时，需要向编译器指出在内核的持续时间里数据是只读的。有两种方法可以实现这一点：

使用内部函数 __ldg
全局内存的限定指针。内部函数 __ldg 用于替代标准指针解引用，并且强制加载通过只读数据缓存，限定指针为 const __restrict__，以表明应该通过只读缓存被访问

__global__ void kernel(float* output, float* input) {...output[idx] += __ldg(&input[idx]);...
}
void kernel(float* output, const float* __restrict__ input) {...output[idx] += input[idx];
}

一般使用 __ldg 是更好的选择。通过 constant 缓存存储的数据必须相对较小而且必须获取同一个地址以便获取最佳性能，相反，只读缓存则可以存放较大的数据，且不必地址一致

constant 缓存和只读缓存对于device来说，都是只读的。二者都有大小限制，前者每个 SM 只能有64KB，后者则是48KB。对于读同一个地址，constant 缓存表现好，只读缓存则对地址较分散的情况表现好