人工智能/深度学习

在动态环境中使用 CUDA 图

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

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

在 ToolkitVersion10 中引入 CUDA 图形时,可以对其进行更新,以反映其实例化中的一些细微变化。此后,此类更新操作的覆盖范围和效率显著提高。在这篇文章中,我描述了一些通过使用 CUDA 图来提高实际应用程序性能的场景,其中一些场景包括图更新功能。

上下文

考虑一个应用程序,该函数具有启动许多短运行内核的功能,例如:

tight_loop(); //function containing many small kernels 

如果每次遇到此函数时都以相同的方式执行,则可以使用流捕获将其转换为 CUDA 图。在本例中,必须引入一个开关布尔值captured,以指示是否已经创建了图形。将此开关的解除 Clara 操作和初始化放在源代码中,使其范围包括对函数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 _循环函数的调用实际上并不执行任何内核启动或其他 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<<<grid_dim, block_dim, 0, stream>>>(params);
     tiny_kernel2<<<grid_dim, block_dim, 0, stream>>>(params);
   }
 } 

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

保存并识别 CUDA 图形

第一种方法从 C ++标准模板库中引入容器来存储参数集。每当您遇到一个新的参数集来唯一地定义函数tight_loop,请将它连同相应的可执行图形一起添加到容器中。

当您遇到容器中已经存在的参数集时,启动相应的 CUDA 图形。假设在本例中,变量firstparams.sizedelta唯一地定义了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<Key, cudaGraphExec_t, cmpKeys> MapOfGraphs; 

无论函数tight_loop出现在何处,都要用填充键的代码将其包装起来,并在容器中查找。如果找到键,代码将启动相应的可执行 CUDA 图。否则,它将创建一个新图形,将其添加到容器中,然后启动它(图 1 )。

Key triplet = {first_step, delta, params.size};
 map<Key, cudaGraphExec_t, cmpKeys>::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);
 } 
Flowchart starting with defining the container for pairs, encountering tightloop, deciding whether the argument set is in the container already, and either relaunching or capturing and launching the graph.
图 1 。保存和识别图形。

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

为 STL 容器编写一个非平凡的比较函数通常并不困难,但当一个键由多个非平凡的实体组成时,可能会很乏味。一种普遍适用的方法是使用词典比较。对于本例,以下代码示例有效:

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 图,它必须与调用上下文完全匹配:

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

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

这种方法的好处是双重的。首先,当更新足够时,可以避免昂贵的新 CUDA 图实例化。第二,你不必知道是什么让图形独一无二。任何图形比较都由 update 函数隐式执行。下面的代码示例实现了此方法。与之前一样,它从开关的解除 Clara 和初始化开始,以指示先前创建的图形。

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); 

在这个场景中,您总是执行流捕获来收集关于tight_loop中 CUDA 操作的信息。这是一个相对便宜的操作,完全在主机上执行,而不是 GPU 。它可以与以前的 CUDA 图形启动重叠,这些启动本身就是异步操作(图 2 )。

Flowchart that starts with encountering tightloop, capturing a graph, trying to update the graph, launching the graph if the update was successful, and instantiating a new graph if the update was not.
图 2 。更新图形

一句警告的话已经准备好了。cudaGraphExecUpdate的复杂性大致与 CUDA 图形节点的更改数量成正比,因此如果大部分节点发生更改,则效率会降低。

后果

推动这两种方法以灵活方式管理 CUDA 图的应用程序有两种不同的工作负载大小,但行为有所不同(表 1 )。所有涉及的内核在单个 NVIDIA A100 GPU 上执行需要 2 – 8 微秒。报告的加速是针对代码中可以转换为 CUDA 图形的部分。

Table 1. Speedups achieved by employing CUDA graphs
  Workload size
  Small Large
Kernels per Graph 504 2520
Graph launches 70 1540
Method Update Recognize Update Recognize
Stream captures 70 5 1540 3
Graph updates 69 N.A 1539 N.A
Graph instantiations 3 5 1 3
Speedup 16% 19% 26% 24%
 

结论

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

 

标签