核心融合如何提高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内核和一个乘法内核，并且假设每个内核在一个向量上工作，每个线程执行以下操作：

1. 从全局内存中加载向量A的元素
2. 将常量添加到我的向量元素中，或由常量乘以。
3. 将我的元素保存回向量A(在全局内存中)

此操作要求每个线程一个读，每个线程写一个。如果我们两个都是背靠背的话，操作的顺序应该是：

添加内核：

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

乘核：

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

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

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

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

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

这种节省对于GPU上的内存绑定操作(如这些)非常重要。通过减少所需的负载和存储数量，总体性能得到改善，通常与减少负载/存储操作的数量成正比。