在这篇 NVIDIA 博客文章中，为什么通过共享内存进行复制的速度更快?

Question

几年前，NVIDIA 的 Mark Harris 发布了以下内容：

CUDA C/C++ 中的高效矩阵转置

其中他描述了如何使用共享内存而不是朴素方法更快地执行矩阵转置。出于方法论目的，他还实现了基于共享内存磁贴的简单矩阵复制版本。

有点令人惊讶的是，通过共享内存磁贴进行复制的速度比"朴素"拷贝(使用 2D 网格(更快：朴素副本为 136 GB/秒，基于共享内存磁贴的拷贝为 152.3 GB/秒。那是在开普勒微架构卡上，特斯拉K20c。

我的问题：为什么这有意义？也就是说，当所做的只是合并的读取和写入时，为什么有效带宽没有降低？具体来说，它是否与未使用__restrict(因此可能未使用__ldg()(有关？

注意：这个问题与换位无关。这篇文章是关于换位的，它的教训很好。它没有讨论涉及简单、非转置复制的奇怪现象。

Answer 1

这不太可能是 GDDR5 读/写，因为它应该完全由 L2 缓存缓冲并被高占用率掩盖。合并的读取/写入(或缺乏(都没有，即使开普勒很容易被这些减慢速度。

我们在这里看到的只是读取和写入之间的更长的管道，它掩盖了读取操作上留下的任何延迟。

<小时 />

for (int j = 0; j < TILE_DIM; j+= BLOCK_ROWS)
odata[(y+j)*width + x] = idata[(y+j)*width + x];

如果没有__restrict，编译器必须假定循环迭代之间的数据依赖性，因此每次迭代都必须隐式同步前一个迭代。这甚至不是不使用__ldg()的效果(如果没有数据重用的可能性，则通过纹理单元不会有什么区别(，而是全局内存读取的直接停滞。

<小时 />

for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x] = idata[(y+j)*width + x];
__syncthreads();
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x];

另一方面，除了同步前的最后几行外，这不必停滞。假设编译器已经展开了这些简单的循环，它变得很明显。

在这种特定情况下，那里的__syncthreads();甚至适得其反，没有充分的理由在开始写出之前等待最后一行完成读取。