重新排序工作维度可能会带来巨大的性能提升，但为什么呢？答案

【问题标题】：Reordering of work dimensions may cause a huge performance boost, but why?重新排序工作维度可能会带来巨大的性能提升，但为什么呢？
【发布时间】：2015-09-12 04:13:04
【问题描述】：

我在 GPU 上使用 OpenCL 进行立体图像处理，在我将 C++ 实现移植到 OpenCL 后，我开始尝试优化。一个非常简单的实验是交换尺寸。

考虑一个简单的内核，它针对二维工作空间（例如 640x480）上的每个像素执行。在我的情况下，这是一个人口普查转换。

交换自：

int globalU = get_global_id(0);
int globalV = get_global_id(1);

也

int globalU = get_global_id(1);
int globalV = get_global_id(0);

在以相同方式调整 NDRange 时，性能提升了约 500%。 3d 空间中的其他实验实现了从 72 毫秒到 2 毫秒的执行时间，仅对维度进行了重新排序。

谁能解释一下，这是怎么回事？它只是内存管道和缓存使用的影响吗？

编辑：图像具有标准的mamory 布局。这就是为什么我想知道效果。我期望最好的速度，当迭代像图像存储在内存中一样时，情况并非如此。

在阅读了一些 AMD APP SDK 文档后，我发现了一些关于内存通道的有趣细节。这可能是一个原因。

【问题讨论】：

我不知道 opencl，但通常二维数组可以将它们的元素放在内存中，这样行或列都在连续的内存中（但不能同时在内存中）。如果您从同一行（或列）访问大部分时间元素，这将解释为什么交换维度会产生如此巨大的影响。
除了@tobi303 所说的，查找“coalesced reads”——它们有很大的不同！

标签： c++ performance opencl gpu

【解决方案1】：

当您访问内存中的元素时，首先将其加载到 CPU 的缓存中。 CPU 不会加载单个元素（例如 1 个字节），而是将单个行（例如 64 个相邻字节）加载到缓存中。这是因为您通常很可能会访问后续元素，因此 CPU 不需要再次访问 RAM。

这有很大的不同，因为要访问高速缓存，电信号甚至不应该离开 CPU 芯片，而如果 CPU 需要从 RAM 加载数据，则信号需要传输到单独的芯片，并且可能不止一个信号from CPU 是必需的，因为它通常需要指定 RAM 中的行和列来访问它的一部分（阅读What every programmer should know about memory 了解更多信息）。实际上，访问缓存的时间可能只需要 0.5 ns，而访问 RAM 需要 100 ns。

所以计算机算法应该考虑到这一点。如果您遍历矩阵中的所有元素，则应该遍历它们，以便您可以几乎同时访问彼此靠近的元素。因此，如果您的矩阵在内存中具有以下布局：

m_0_0, m_0_1, m_0_2, ... m_1_0, m_1_1（第一列、第二列等）

您应该按顺序访问元素：m_0_0、m_0_1、m_0_2（按列）

如果您使用不同的访问顺序（在这种情况下按行），当您访问第一列中的第一个元素时，CPU 将在缓存中加载第一列的一部分，然后当您访问第二列中的第一个元素时，第二列的一部分等等。当您将遍历第一行并访问第一列中的第二个元素时，第一列的缓存值将不再存在于缓存中，因为它的大小有限（并且非常小）。因此，这样的算法将有效地消除缓存。

【讨论】：