本篇主要介绍CUDA编程中的Memory coalescing概念。
翻译整理自 Programming Massively Parallel Processors 4th
memory coalescing就是内存合并，通常用于高效地在全局内存、共享内存、寄存器之间传输数据。
CUDA kernel 性能的一个最重要因素是访问全局内存中数据的速度，有限的带宽很可能成为瓶颈。毕竟CUDA kernel一下子要处理很多数据，数据从哪来，得先有数据才能算了。所以如何高效传输数据当然就比较重要了。
这里聊到的memory coalescing通常与tile技术结合使用，从而高效利用带宽。
全局内存 global memory 的特性
Global memory 是用 DRAM 实现的。数据位存储在小电容中的 DRAM 单元中，其中电荷的存在或缺失区分了 1 和 0 值。从 DRAM 单元读取数据需要小电容使用其微小的电荷来驱动连接到传感器的高电容线路，并触发传感器的检测机制，以确定电容中是否有足够的电荷作为 “1”。这个过程在目前 DRAM 芯片中需要几十纳秒。
但GPU算的更快，所以一般使用并行性来增加其数据访问速率，也就是所谓的 memory access throughput。设备每次访问DRAMs中的一个位置时，实际上包含目标请求位置在内的一系列连续位置都会被访问到。
说白了就是处理器很快，但内存读取慢，所以希望有一种方式可以增大读取速度
为什么 DRAM 如此慢？
下图显示了一个 DRAM 单元及其内容访问路径。decoder（解码器）是一个电子电路，它使用晶体管来驱动连接到成千上万单元的输出门的线路。充电或放电到期望水平可能需要很长时间。更难的是让单元驱动连接到感应放大器的垂直线，并允许感应放大器检测其内容。这是基于电荷共享。门会释放存储在单元中的微小电荷。如果单元内容是“1”，微小电荷必须将长位线的大电容的电势提升到足够高的水平，以便触发感应放大器的检测机制。一个很好的类比是，有人在长走廊的一端拿着一小杯咖啡，而走廊另一端的人需要通过沿走廊传播的香气来确定咖啡的味道。
为了提升处理速度，理论上可以在每个存储单元中使用体积更大、性能更强的电容。但实际上，DRAM 的发展趋势恰恰相反。为了在每块芯片上存储更多数据，每个单元中的电容体积不断缩小，其性能也随之降低。
如何利用这个特性
每次访问 DRAM 位置时，都会访问包括请求位置在内的一系列连续位置。每个 DRAM 芯片中都配备了许多传感器，它们同时并行工作。每个传感器都能感知这些连续位置中一个比特的内容。一旦由传感器检测到，所有这些连续位置的数据都可以高速传输到处理器。这些被访问和传递的连续位置被称为 DRAM bursts。如果应用程序集中使用这些突发中的数据，DRAM 可以以远高于访问真正随机位置序列的情况下的速率提供数据。
在任何给定时间点，一个 warp 中的线程执行相同的指令。当一个 warp 中的所有线程执行加载指令时，硬件会检测它们是否访问连续的全局内存位置。换句话说，最有利的访问模式是当一个 warp 中的所有线程访问连续的全局内存位置时。在这种情况下，硬件将所有这些访问合并或合并为对连续 DRAM 位置的一次集中访问。例如，对于一个 warp 的给定加载指令，如果线程 0 访问全局内存位置 X，线程 1 访问位置 X + 1，线程 2 访问位置 X + 2，等等，所有这些访问将被合并或组合为单一请求，以便在访问 DRAM 时访问连续位置。这种合并访问允许 DRAM 以 bursts 形式传递数据。
举个栗子
C 和 CUDA 中的多维数组元素是根据行优先约定放置在线性寻址的内存空间中的。请记住，行优先这个术语表示在内存中连续存放数组的每一行。在行优先存储中，二维数组的第一行元素先连续存储，然后是第二行，依此类推。在 C 和 CUDA 中，这意味着数组元素是按照它们在行中的顺序存储的。 如下图展示：
上图中第一行的四个元素首先按照行中出现的顺序放置。然后放置行 1 的元素，接着是行 2 的元素，然后是行 3 的元素。尽管 M0,0 和 M1,0 在二维矩阵中看起来是连续的，但在线性地址的内存中它们相隔四个位置。
为了有效利用合并硬件，理解这种内存布局是非常重要的。当访问多维数组时，如果访问模式与内存中数据的布局一致，那么这些访问就可能被硬件合并为较少的、更高效的内存访问操作。
可以合并访问的情况
如果一个 warp 的线程依次访问同一行中的连续元素，那么这些访问就可能被合并成一个单一的 DRAM bursts 访问，从而提高内存访问效率。反之，如果访问模式与内存布局不一致，比如跨行访问元素，那么合并硬件就无法有效地工作，从而降低了内存访问的效率。
如果刚上述提到的二维数组是矩阵乘法中使用的第二个输入矩阵。在这种情况下，分配给连续输出元素的 warp 中的连续线程将遍历这个输入矩阵的连续列。
看下图的左上部分的代码，右上部分显示了访问模式的逻辑视图：连续线程遍历连续列。通过检查代码可以发现，对 M 的访问可以被合并。数组 M 的索引是 k×Width+col。变量 k 和 Width 在 warp 中的所有线程中都有相同的值。变量 col 被定义为 blockIdx.x×blockDim.x+threadIdx.x，这意味着连续线程（具有连续 threadIdx.x 值）将具有连续的 col 值，并因此访问 M 的连续元素。
图中的下半部分显示了访问模式的物理视图。在迭代 0 中，连续的线程将访问位于内存中相邻的行 0 中的连续元素，如图 6.2 中的“迭代 0 的加载”所示。在迭代 1 中，连续的线程将访问同样位于内存中相邻的行 1 中的连续元素，如图 6.2 中的“迭代 1 的加载”所示。这个过程对所有行都继续进行。我们可以看到，线程在此过程中形成的内存访问模式是一个有利于合并的模式。
不可以合并访问的情况
现在假设矩阵是以列优先顺序而不是行优先顺序存储的。这可能有各种原因。例如，我们可能在乘以以行优先顺序存储的矩阵的转置。在线性代数中，我们经常需要使用矩阵的原始形式和转置形式。最好避免创建和存储两种形式。
图 6.3 描述了当矩阵以列优先顺序存储时，连续的线程如何遍历连续的列。图 6.3 的左上部分显示了代码，右上部分显示了内存访问的逻辑视图。程序仍然尝试让每个线程访问矩阵 M 的一列。通过检查代码可以看出，对 M 的访问并不利于合并。数组 M 的索引是 col×Width+k。与之前一样，col 被定义为 blockIdx.x×blockDim.x+threadIdx.x，这意味着连续的线程（具有连续 threadIdx.x 值）将具有连续的 col 值。然而，在 M 的索引中，col 乘以 Width，这意味着连续的线程将访问 M 中相隔 Width 的元素。因此，这些访问并不利于合并。
在图 6.3 的下半部分，我们可以看到，内存访问的物理视图与图 6.2 中的相当不同。在迭代 0 中，连续的线程将逻辑上访问行 0 中的连续元素，但由于列优先布局，这次它们在内存中并不相邻。这些加载在图 6.3 中显示为“迭代 0 的加载”。类似地，在迭代 1 中，连续的线程将访问行 1 中的连续元素，它们也不在内存中相邻。对于一个真实的矩阵，每个维度通常有数百甚至数千个元素。每次迭代中相邻线程访问的 M 元素可能相隔数百甚至数千个元素。硬件将确定这些元素彼此相隔较远，无法合并。
如何优化
当计算本身不自然适合合并时，有各种策略可以优化代码以实现内存合并。一种策略是重新安排线程映射到数据的方式；另一种策略是重新安排数据本身的布局。
还有种策略是以合并的方式在全局内存和共享内存之间传输数据，并在共享内存中执行不利的访问模式，这提供了更快的访问latency。
图 6.4 举例说明了如何应用角转换（corner turning）。在这个示例中，A 是一个以行优先布局存储在全局内存中的输入矩阵，B 是一个以列优先布局存储在全局内存中的输入矩阵。它们相乘以产生一个以行优先布局存储在全局内存中的输出矩阵 C。该示例说明了四个线程如何负责输出tile顶部的四个连续元素并加载输入tile元素。
对矩阵 A 中输入块的访问类似于第 5 章“内存架构和数据局部性”中的情况。四个线程加载输入块顶部的四个元素。每个线程加载的输入元素的局部行和列索引与线程输出元素在输出块中的索引相同。这些访问是合并的，因为连续线程按行优先布局在内存中相邻地访问 A 中同一行的连续元素。
另一方面，对矩阵 B 中输入块的访问需要与第 5 章“内存架构和数据局部性”中的情况不同。图 6.4(A) 显示了如果我们使用与第 5 章相同的布局，访问模式会是什么样子。尽管四个线程在逻辑上加载输入块顶部的四个连续元素，但连续线程加载的元素在内存中彼此相距很远，这是因为 B 元素的列优先布局。换句话说，负责输出块中同一行连续元素的连续线程在内存中加载的位置是不连续的，这导致了未合并的内存访问。
这个问题可以通过指派四个连续线程来加载输入块左侧边缘（同一列）的四个连续元素来解决，如图 6.4(B) 所示。直观上，我们在每个线程计算加载 B 输入块的线性化索引时交换了 threadIdx.x 和 threadIdx.y 的角色。由于 B 是列优先布局，同一列中的连续元素在内存中是相邻的。因此，连续线程加载的输入元素在内存中是相邻的，这确保了内存访问是合并的。
如果把 B 元素的一小块（tile）放入共享内存中，无论是列优先布局还是行优先布局，因为shared memory基于SRAM且访问延迟较低，不需要这种合并操作
内存合并的主要优点是它通过将多个内存访问合并为单一访问来减少全局内存流量。当访问同时发生并且访问相邻的内存位置时，可以合并访问。
换种方式理解
交通拥堵不仅仅出现在计算中。我们大多数人都经历过高速公路系统中的交通拥堵，高速公路交通拥堵的根本原因是有太多的车辆试图在为远远较少的车辆设计的道路上行驶。当拥堵发生时，每辆车的行驶时间大大增加。上下班通勤时间在交通拥堵时很容易增加一倍或三倍。
减少交通拥堵的大多数解决方案涉及减少道路上的汽车数量。假设通勤者数量是恒定的，人们需要共享乘车以减少道路上的汽车数量。共享乘车的常见方式是拼车，其中一群通勤者的成员轮流驾驶一辆车将团体送到工作地点。政府通常需要制定政策来鼓励拼车。在一些国家，政府简单地不允许某些类型的车辆每天在道路上行驶。例如，奇数车牌的汽车可能不允许在星期一、星期三或星期五上路。这鼓励不同天数车辆被允许上路的人们组成拼车小组。在其他国家，政府可能会为减少道路上汽车数量的行为提供激励。例如，在某些国家，拥挤的高速公路的某些车道被指定为拼车车道；只有超过两到三人的汽车才被允许使用这些车道。也有国家政府使汽油非常昂贵，以至于人们为了省钱而组成拼车小组。所有这些鼓励拼车的措施都是为了克服拼车需要额外努力的事实，正如我们在图 6.6 中展示的那样。
拼车要求希望拼车的工作者在共同的通勤时间表上做出妥协和一致。图 6.6 的上半部分展示了一个适合拼车的良好时间表模式。时间从左到右推移。工作者 A 和工作者 B 有相似的睡眠、工作和晚餐时间表。这使得这两位工作者能够轻松地一起开车上班和回家。他们相似的时间表使他们更容易就共同的出发时间和返回时间达成一致。但在图 6.6 的下半部分所示的时间表中情况并非如此。在这种情况下，工作者 A 和工作者 B 有非常不同的时间表。工作者 A 通宵狂欢至日出，白天睡觉，晚上去上班。工作者 B 夜间睡觉，早上去上班，并在下午 6:00 回家吃晚餐。这些时间表差异如此之大，以至于这两位工作者不可能协调出共同的时间来开车上下班。
内存合并与拼车安排非常相似。我们可以将数据视为通勤者，将 DRAM 访问请求视为车辆。当 DRAM 请求的速率超过 DRAM 系统提供的访问带宽时，交通拥堵情况加剧，算术单元变得空闲。如果多个线程从同一 DRAM 位置访问数据，它们可以形成一个“拼车团队”并将它们的访问合并成一个 DRAM 请求。然而，这要求线程具有类似的执行时间表，以便它们的数据访问可以合并为一个。同一 warp 中的线程是完美的候选者，因为它们都同时执行加载指令，这是由SIMD 执行的特性。
参考
- Programming Massively Parallel Processors 4th
- CUDA 编程之 Memory Coalescing - 知乎
- Irregular memory access patterns and the cache - CUDA Programming and Performance - NVIDIA Developer Forums