CUDA C编程（十五）常量内存

最新推荐文章于 2025-07-13 10:02:09 发布

原创

最新推荐文章于 2025-07-13 10:02:09 发布 · 2.2k 阅读

8 ·

CC 4.0 BY-SA版权

文章标签：

#cuda

常量内存是一种专用的内存，它只用于只读数据和统一访问线程束的数据。常量内存对内核代码而言是只读的，但它对主机而言即是可读的又是可写的。常量内存位于设备的DRAM上（和全局内存一样），并且有一个专用的片上缓存。和一级缓存和共享内存一样，从每个SM的常量缓存中读取的延迟，比直接从常量内存中读取要低得多。每个SM常量内存缓存大小的限制为64KB。
常量内存有一个不同的最优访问模式，在常量内存中，如果线程束中的所有线程都访问相同的位置，那么这个访问模式就是最优的。如果线程束中的线程访问不同的地址，那么访问就需要串行。因此，一个常量内存读取的成本与线程束中线曾读取唯一地址的数量呈线性关系。在全局作用域中必须用以下修饰符声明常量内存：__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。

使用常量内存实现一维模板
在数值分析中，模板计算在几何点集合上应用函数，并用输出更新单一点的值。模板是求解许多偏微分方程算法的基础。在一维中，在位置x周围九点模板会给这些位置上的值应用一些函数：{x-4h,x-3h,x-2h,x-h,x,x+h,x+2h,x+3h,x+4h}，下图展示了一个九点模板：
在这里插入图片描述
一个九点模板的例子是实变量函数f在点x上一阶导数的第八阶中心差分公式。裂解这个公式的应用并不重要，只要简单的了解到它会将上述的九点作为输入并产生单一输出。接下来将这个公式作为一个示例模板：
在这里插入图片描述
在一维数组中对该公式的应用是对一个数据进行并行操作，该操作能很好地映射到CUDA。它可以为每个线程分配位置x，并计算出f’(x)。核函数如下：

__global__ void stencil_ld(float *in, float *out)
{
   
   
   //由于每个线程需要9个点来计算一个点，所以要使用共享内存来优化缓存数据，从而减少对全局内存的冗余访问
   //RADIUS定义了点x两侧点的数量，这些点被用于计算x点的值，这里RADIUS被定义为4
   __shared__ float smem[BDIM + 2 * RADIUS];
   
   //计算访问全局内存的索引
   int idx = blockIdx.x + blockDim.x + threadIdx.x;
   //计算访问共享内存的每个线程的索引
   int sidx = threadIdx.x + RADIUS;
   smem[sidx] = in[idx];

   //从全局内存中读取数据到共享内存中时，前四个线程负责从左侧到右侧的光环中读取数据到共享内存中
   if(threadIdx