4

几年前，NVIDIA 的 Mark Harris 发布了这个：

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

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

有点令人惊讶的是，通过共享内存块进行复制的速度比“原始”复制（使用 2D 网格）要快：原始复制为 136 GB/秒，基于共享内存块的复制为 152.3 GB/秒。那是在 Kepler 微架构卡 Tesla K20c 上。

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

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

4

1 回答 1

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();特定情况下甚至会适得其反，在开始写出之前没有充分的理由等待最后一行完成读取。

于 2020-08-03T08:19:02.030 回答