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. 通过减少（即消除）“多余”的负载和存储

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

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上的内存绑定操作（如此类操作）而言，这种节省可能非常重要。通过减少所需的装载和存储数量，通常可以与减少装载/存储操作数量成比例地改善整体性能。