CUDA 内核融合如何提高 GPU 上的内存绑定应用程序的性能？

Question

我一直在研究比 GPU 上用于基本计算的设备可用内存更大的流数据集。主要限制之一是 PCIe 总线通常限制在 8GB/s 左右，内核融合可以帮助重用可重用的数据，并且它可以利用 GPU 内的共享内存和局部性。我发现的大多数研究论文都很难理解，并且大多数都在复杂的应用程序中实现了融合，例如https://ieeexplore.ieee.org/document/6270615。我读过很多论文，但都未能解释一些将两个内核融合在一起的简单步骤。

我的问题是融合实际上是如何工作的？。将普通内核更改为融合内核的步骤是什么？此外，是否需要多个内核才能融合它，因为融合只是一个花哨的术语，用于消除一些内存绑定问题，并利用局部性和共享内存。

我需要了解如何将内核融合用于基本 CUDA 程序，例如矩阵乘法或加法和减法内核。一个非常简单的例子（代码不正确，但应该给出一个想法）如：

int *device_A;
int *device_B;
int *device_C;

cudaMalloc(device_A,sizeof(int)*N);

cudaMemcpyAsync(device_A,host_A, N*sizeof(int),HostToDevice,stream);

KernelAdd<<<block,thread,stream>>>(device_A,device_B); //put result in C
KernelSubtract<<<block,thread,stream>>>(device_C);

cudaMemcpyAsync(host_C,device_C, N*sizeof(int),DeviceToHost,stream); //send final result through the PCIe to the CPU

Answer 1

内核融合背后的基本思想是将 2 个或更多内核转换为 1 个内核。操作是合并的。最初可能并不明显有什么好处。但它可以提供两种相关的好处：

1. 通过重用内核可能在寄存器或共享内存中填充的数据
2. 通过减少（即消除）“冗余”负载和存储

让我们使用像您这样的示例，其中我们有一个 Add kernel 和一个 multiply kernel，并假设每个内核都在一个向量上工作，并且每个线程执行以下操作：

1. 从全局内存中加载向量 A 的元素
2. 向我的向量元素添加一个常数或乘以一个常数
3. 将我的元素存储回向量 A（在全局内存中）

此操作要求每个线程读取一次，每个线程写入一次。如果我们两个都背靠背进行，操作顺序如下：

添加内核：

1. 从全局内存中加载向量 A 的元素
2. 向我的向量元素添加一个值
3. 将我的元素存储回向量 A（在全局内存中）

乘核：

1. 从全局内存中加载向量 A 的元素
2. 将我的向量元素乘以一个值
3. 将我的元素存储回向量 A（在全局内存中）

我们可以看到第一个内核中的第 3 步和第二个内核中的第 1 步所做的事情对于实现最终结果并不是真正必要的，但由于这些（独立）内核的设计，它们是必要的。一个内核无法将结果传递给另一个内核，除非通过全局内存。

但是如果我们将两个内核组合在一起，我们可以这样写一个内核：

1. 从全局内存中加载向量 A 的元素
2. 向我的向量元素添加一个值
3. 将我的向量元素乘以一个值
4. 将我的元素存储回向量 A（在全局内存中）

这个融合内核执行这两种操作，产生相同的结果，但不是 2 个全局内存加载操作和 2 个全局内存存储操作，它只需要每个操作 1 个。

这种节省对于 GPU 上的内存绑定操作（例如这些）来说非常重要。通过减少所需的加载和存储次数，可以提高整体性能，通常与加载/存储操作次数的减少成正比。