CUDA C编程（十四）合并的全局内存访问

最新推荐文章于 2025-09-26 19:55:59 发布

原创

最新推荐文章于 2025-09-26 19:55:59 发布 · 2k 阅读

6 ·

CC 4.0 BY-SA版权

文章标签：

#cuda

本文探讨了如何通过共享内存优化矩阵转置过程，避免全局内存交叉访问，包括使用二维共享内存缓存、填充和展开技巧，以及调整线程块配置以提高并行性能。作者给出了基准和改进版本的核函数，展示了在不同GPU上的性能测试结果。

使用共享内存也能帮助避免对未合并的全局内存的访问，矩阵转置就是一个典型的例子：读操作被自然合并，但写操作是按照交叉访问的。其中交叉访问是全局内存中最糟糕的访问模式，因为它浪费总线带宽，在共享内存的帮助下，可以先在共享内存中进行转置操作，然后再对全局内存进行合并写操作。接下来将介绍使用多个线程块对基于交叉的全局内存访问重新排序到合并访问。
基准转置内核
作为基准，下面的核函数是一个仅使用全局内存的矩阵转置的朴素实现：

__global__ void naiveGmem(float *out, float *in, const int nx, const int ny)
{
   
   
   unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
   unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
   
   if(ix < nx && iy < ny)
   {
   
   
      out[ix * ny + iy] = in[iy * nx + ix];
   }
}

使用共享内存的矩阵转置
为了避免交叉全局访问，可以使用二维共享内存来缓存原始矩形的数据。从二维共享内存中读取的一列可以被转移到转置矩阵行中，它被存储在全局内存中。虽然朴素实现将导致共享内存存储体冲突，但是这个结果将比未合并的全局访问好得多。下图显示了在矩阵转置中是如何使用共享内存的。
在这里插入图片描述
下面的核函数实现了使用共享内存的矩阵转置。它可以被看作前面所讨论的setRowReadCol函数的扩展，这两个核函数的差别在于setRowReadCol使用一个线程块处理输入矩阵的单块转置，而transposeSmem扩展了转置操作，使用了多个线程块和多个数据块：

//存储操作是无冲突的，但是加载操作显示有16路冲突
__global__ void setRowReadCol(int *out)
{
   
   
   __shared__ int tile[</

最低0.47元/天解锁文章