在动态环境中使用CUDA图

通过将计算密集型部分放到 GPU 上，可以大大加快许多工作负载的速度。 在 CUDA 术语中，这称为启动内核。 当这些内核很多且持续时间很短时，启动开销有时会成为问题。

CUDA Graphs 提供了一种减少这种开销的方法。 图之所以有效，是因为它们将任意数量的异步 CUDA API 调用（包括内核启动）组合成一个只需要一次启动的操作。 它们在创建时确实会产生一些开销，因此它们最大的好处来自于多次重用它们。

在工具包版本 10 中引入时，CUDA 图可能已经更新以反映其实例化的一些细微变化。 此后，此类更新操作的覆盖范围和效率显着提高。 在这篇文章中，我描述了一些通过使用 CUDA 图来提高实际应用程序性能的场景，其中一些包括图更新功能。

Context

考虑一个具有启动许多短期运行内核的函数的应用程序，例如：

tight_loop(); //function containing many small kernels 

如果每次遇到这个函数时都执行相同的操作，那么很容易使用流捕获将其转换为 CUDA 图。 您必须引入一个开关（在本例中为捕获的布尔值）来指示是否已经创建了图。 将此开关的声明和初始化放在源代码中，使其范围包括函数tight_loop 的每次调用。

cudaGraphExec_t instance; static bool captured = false;

接下来，使用代码包装函数的任何实际调用以创建其对应的 CUDA 图（如果它尚不存在），然后启动该图。

if (!captured) {// you haven’t captured the graph yet   cudaGraph_t graph;    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);    tight_loop(); //function containing many small kernels   //aggregate all info about the stream capture into “graph”   cudaStreamEndCapture(stream, &graph);   //turn this info into executable CUDA Graph “instance”   cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);   //no need to hang on to the stream info   cudaGraphDestroy(graph);   captured = true; } cudaGraphLaunch(instance, stream);//launch the executable graph 

对tight_loop 函数的调用实际上并不执行任何内核启动或其他CUDA 操作。 它仅记录所有此类操作并将它们存储在数据结构中。

专注于启动内核的功能。 在实际应用中，它看起来像以下代码：

void tight_loop(int first_step, MyStruct params, int delta, dim3 grid_dim, dim3 block_dim, cudaStream_t stream) {   for (int step = first_step; step >= 0; --step, params.size -= delta)   {     tiny_kernel1<<>>(params);     tiny_kernel2<<>>(params);   } } 

显然，如果函数的参数在连续调用时发生变化，则表示 GPU 内部工作的 CUDA 图也应该发生变化。 您不能重复使用原始图表。 但是，假设多次遇到相同的函数参数集，您至少可以通过两种不同的方式处理这种情况：保存和识别图或更新图。

保存和识别图

第一种方法从 C++ 标准模板库中引入一个容器来存储参数集。 每当你遇到一个新的参数集唯一地定义了函数tight_loop，把它和它对应的可执行图一起添加到容器中。

当您遇到容器中已经存在的参数集时，启动相应的 CUDA 图。 假设在这种情况下，变量 first、params.size 和 delta 唯一地定义了tight_loop。 这个三元组是用来区分图形的关键。 您定义它和要在源代码中使用的容器，使其范围包括函数tight_loop 的每次调用。

typedef struct  { //define the fields of a key    int first;    double delta;    int size; } Key; //define the container (map) containing (key,value) pairs map MapOfGraphs; 

无论函数tight_loop 出现在哪里，您都使用填充Key的代码将其包装起来，并在您的容器中查找它。 如果找到Key，代码将启动相应的可执行 CUDA 图。 否则，它会创建一个新图，将其添加到容器中，然后启动它（下图所示）。

Key triplet = {first_step, delta, params.size}; map::iterator it = MapOfGraphs.find(triplet); if (it == MapOfGraphs.end()) { // new parameters, so need to capture new graph   cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);   tight_loop(first_step, params, delta, grid_dim, block_dim, stream);   cudaStreamEndCapture(stream, &graph);   cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);   cudaGraphDestroy(graph);   // add this graph to the container of saved graphs   MapOfGraphs.insert(make_pair(trio,instance));   cudaGraphLaunch(instance, stream); } else {// recognized parameters, so can launch previously captured graph   cudaGraphLaunch(it->second, stream); } 

这种方法通常效果很好，但它有一些内在的危险。 在本例中，您确定只需要三个参数来定义容器中的Key。 对于不同的工作负载，这可能会有所不同，或者其他开发团队成员可能会默默地添加字段来构造 MyStruct。 这会影响非平凡函数 cmpKeys 的编写方式。 该函数是容器所需要的，用于判断某个Key是否小于另一个Key。

为 STL 容器编写一个重要的比较函数通常并不困难，但当一个Key由多个重要的实体组成时可能会很乏味。 一种普遍适用的方法是使用字典比较。 对于此示例，以下代码示例有效：

struct cmpKeys {   bool operator()(const Key& a, const Key& b) const {     if (a.first != b.first) return(a.first < b.first);     else      { if (a.delta != b.delta ) return (a.delta < b.delta); else return(a.size < b.size);     }   } } 

请记住，要重用以前捕获的可执行 CUDA 图，它必须与调用上下文完全匹配：

• 相同的拓扑
  • 相同数量和类型的图形节点
  • 图节点之间的相同依赖关系
• 相同的节点参数

但是，如果其拓扑保持不变，则可以调整 CUDA 图以使其符合新需求。 存在一种方便的机制来确认拓扑等价，同时调整节点参数以返回修改后的可执行图。 它由 cudaGraphExecUpdate 提供，它通过将现有的可执行图与新派生的图（例如，通过流捕获方便地获得）进行比较来工作。 如果可能，差值用于进行更改。

这种方法的好处是双重的。 首先，当更新足够时，您可以避免昂贵的新 CUDA 图实例化。 其次，您不必知道是什么让图表与众不同。 任何图形比较都由更新函数隐式执行。 下面的代码示例实现了此方法。 和以前一样，它从声明和初始化开关开始，以指示先前创建的图形。

static bool captured = false;

将tight_loop的调用点修改如下：

 cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); tight_loop(first_step, params, delta, grid_dim, block_dim, stream); cudaStreamEndCapture(stream, &graph); if (!captured) {    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);    captured = true; } else {   // graph already exists; try to apply changes   if (cudaGraphExecUpdate(instance, graph, NULL, &update) != cudaSuccess)    {// only instantiate a new graph if update fails     cudaGraphExecDestroy(instance);     cudaGraphInstantiate(&instance, graph, NULL, NULL, 0)   { } cudaGraphDestroy(graph); cudaGraphLaunch(instance, stream); 

在这种情况下，您总是进行流捕获以收集有关紧密循环中的 CUDA 操作的信息。 这是一个相对便宜的操作，完全在主机上执行，而不是在 GPU 上。 它可以与之前的 CUDA 图启动重叠，后者本身就是异步操作（如下图所示）。

需要注意的是。 cudaGraphExecUpdate 的复杂性大致与对 CUDA 图节点所做的更改数量成正比，因此如果大部分节点发生更改，它的效率就会降低。

结果

促使这两种方法以灵活方式管理 CUDA 图的应用程序具有两种不同的工作负载大小，其行为略有不同。 所有涉及的内核都需要 2-8 微秒才能在单个 NVIDIA A100 GPU 上执行。 报告的加速是针对可以转换为 CUDA 图的代码部分。

结论

具有许多小型 CUDA 内核的应用程序通常可以使用 CUDA 图加速，即使内核启动模式在整个应用程序中发生变化。 考虑到这种动态环境，最佳方法取决于您的应用程序的具体情况。 希望您发现本文中描述的两个示例易于理解和实现。