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

最新推荐文章于 2024-05-04 14:06:25 发布

伴君

最新推荐文章于 2024-05-04 14:06:25 发布

阅读量1.9k

点赞数 2

分类专栏： # CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/AAAA202012/article/details/122292905

版权

CUDA 专栏收录该内容

39 篇文章 172 订阅

订阅专栏

常量内存是一种专用的内存，它只用于只读数据和统一访问线程束的数据。常量内存对内核代码而言是只读的，但它对主机而言即是可读的又是可写的。常量内存位于设备的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.x < RADIUS)
   {
      smem[sidx - RADIUS] = in[idx - RADIUS];
      smem[sidx + BDIM] = in[idx + BDIM];
   }
   
   __synthreads();

   float tmp = 0.0f;
   //提示CUDA编译器，表明这个循环将自动展开
   #pragma unroll
   for(int i = 1; i < RADIUS; i++)
   {
      tmp += coef[i] * (smem[sidx + i] - smem[sidx - i]);
   }
   
   out[idx] = tmp;
}

//在常量内存中声明coef数组，代码如下所示
__constant__ float coef[RADIUS + 1];
//然后使用cudaMemcpyToSymbol的CUDA API调用从主机端初始化的常量内存：
void setup_coef_constant(void)
{
   const float h_coef[] = {a0, a1, a2, a3, a4};
   cudaMemcpyToSymbol(coef, h_coef,(RADIUS + 1) * sizeof(float));
}

与只缓内存的比较
Kepler GPU添加了一个功能，即使用GPU纹理流水线作为只读缓存，用于存储全局内存中的数据，因为这是一个独立地只读数据，它带有从标准全局内存读取的独立内存带宽，所以使用此功能可以为带宽限制内核提供性能优势。每个Kepler SM都有48KB的只读缓存。一般来说，只读缓存在分散读取方面比一级缓存更好，当线程束的线程都读取相同地址时，不应使用只读缓存。只读缓存的粒度为32个字节。
当通过只读缓存访问全局内存时，需要向编译器指出在内核的持续时间里数据是只读的。有两种方法可以实现这一点：1.使用内部函数__ldg;2.全局内存的限定指针。内部函数__ldg用于替代标准指针解引用，并且强制加载通过只读数据缓存，也可以限定指针为const__restrict__,以表明它们应该通过只读缓存被访问，如下所示：

__global__ void kernel(float* output, float* input)
{
   output[idx] = __dlg(&input[idx]);
}

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

在只读缓存机制需要更多显示控制的情况下，或者在代码非常复杂以至于编译器无法检测到只读缓存的使用是否是安全的情况下，内部函数__ldg是一个更好的选择。
只读缓存是独立的，而且区别于常量缓存，通过常量缓存加载的数据必须是相对较小的，而且访问必须一致以获得良好的性能(一个线程束内的所有线程在任何给定的时间内应该都访问相同的位置)，而通过只读缓存加载的数据可以是比较大的，而且能够在一个非统一的模式下进行访问。下面的内核是根据以前的模板内核修改而来，它使用只读缓存来存储之前在常量内存中的系数。比较一下这两个内核，会发现它们唯一的区别就是函数声明。

__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.x < RADIUS)
   {
      smem[sidx - RADIUS] = in[idx - RADIUS];
      smem[sidx + BDIM] = in[idx + BDIM];
   }
   
   __synthreads();

   float tmp = 0.0f;
   //提示CUDA编译器，表明这个循环将自动展开
   #pragma unroll
   for(int i = 1; i < RADIUS; i++)
   {
      tmp += coef[i] * (smem[sidx + i] - smem[sidx - i]);
   }
   
   out[idx] = tmp;
}
//因为该系数最初是存储在全局内存中并且读取缓存中的，调用内核之前必须分配和初始化全局内存以便在设备上存储系数，代码如下
const float h_coef[] =  {a0, a1, a2, a3, a4};
cudaMalloc((float**)&d_coef,(RADIUS + 1)* sizeof(float));
cudaMemcpy(d_coef,h_coef,(RADIUS + 1)* sizoef(float),cudaMemcpyHostToDevice);