CUDA编程 | 5.5 常量内存
注：此博客是对谭升的博客的一些学习感悟，详细内容请移步谭升的博客进行学习
谭升的博客
目录
5.5 常量内存
使用常量内存实现一维模板
与只读缓存的比较
总结
5.5 常量内存
常量内存是专用内存，它用于只读数据和线程束统一访问某一个数据。
常量内存对内核代码而言是只读的，但是主机是可以修改（写）只读内存的，当然也可以读。
常量内存并不是在片上的，而是在DRAM上，而其有在片上对应的缓存，其片上缓存就和一级缓存和共享内存一样， 有较低的延迟，但是容量比较小，合理使用可以提高内和效率，每个SM常量缓存大小限制为64KB。
所有的片上内存，我们是不能通过主机赋值的，我们只能对DRAM上内存进行赋值。
常量内存变量的生存周期与应用程序生存周期相同，所有网格对声明的常量内存都是可以访问的.
初始化常量内存使用以下函数完成
cudaError_t cudaMemcpyToSymbol(const void *symbol, const void * src, size_t count, size_t offset, cudaMemcpyKind kind)
使用常量内存实现一维模板
所谓模板操作就是把一组数据中中间连续的若干个产生一个输出
紫色的输入数据通过某些运算产生一个输出——绿色的块
f′(x)≈c3(f(x+4h)−f(x−4h))+c2(f(x+3h)−f(x−3h))+c1(f(x+2h)−f(x−2h))+c0(f(x+h)−f(x−h))
原文中有这样一段话：
无论是观察公式还是图示，我们会发现，如果模板大小是9，那么我们输出的前4个数据是没办法计算的因为要使用第-1，-2，-3，-4位置的数据，最后4个数据也是不能计算的，因为他要使用 n+1,n+2,n+3,n+4n+1,n+2,n+3,n+4 的数据，这些数据也是没有的，为了保证计算过程中访问不会越界，我们把输入数据两端进行扩充，也就是把上面虽然没有，但是要用的数据填充到输入数据中，当我们我们要处理的是中间的某段数据的时候，那么填充位的数据就来自前面的块对应的输入数据，或者后面线程块对应的输入数据
这段话我起初没有看明白，经过反复思考和GPT的帮助，认为应该是下面的意思：（仅代表个人观点）
对于输入的数据的每一个数据，我们都需要进行上述的操作，即使用周围的八个点来共同预估出改点的导数。因此，对每个数据导数的计算都需要周围八个点的辅助，那显然，输入数据列的前四个数据和后四个数据并没有完整的八个数据来帮助自己进行导数的计算。而输入数据列中，第一个和最后一个数据是缺少帮助数据个数最多的数据，分别都缺四个。因此，我们需要在输入数据列的前后各补四个数据，来保证输入数据列的每个数据周围都有八个数据来帮助自己计算导数。
与只读缓存的比较
只读缓存拥有从全局内存读取数据的专用带宽,所以，如果内核函数是带宽限制型的，那么这个帮助是非常大的
只读缓存对于分散访问的更好，当所有线程读取同一地址的时候常量缓存最好，只读缓存这时候效果并不好，只读缓存粒度为32
使用只读缓存需要更多的声明和控制，在代码非常复杂的情况下以至于编译器都没办法保证只读缓存使用是否安全的情况下，建议使用 __ldg()函数，更容易控制。
使用__ldg()的方法：
__global__ void kernel(float* output, float* input) {
...
output[idx] += __ldg(&input[idx]);
...
}使用限定指针的方法：
void kernel(float* output, const float* __restrict__ input) {
...
output[idx] += input[idx];
}只读缓存独立存在，区别于常量缓存，常量缓存喜欢小数据，而只读内存加载的数据比较大，可以在非统一模式下访问。
总结
常量内存和只读缓存：
- 对于核函数都是只读的
- SM上的资源有限，常量缓存64KB，只读缓存48KB
- 常量缓存对于统一读取（读同一个地址）执行更好
- 只读缓存适合分散读取