Getting Started with CUDA Graphs

翻译博客：Getting Started with CUDA Graphs
https://developer.nvidia.com/blog/cuda-graphs

GPU架构的性能随着新一代的出现而不断提高。现代GPU速度如此之快，以至于在许多感兴趣的情况下，每个GPU操作（例如内核或内存拷贝） (e.g. kernel or memory copy) 所花费的时间现在都以微秒为单位。然而，将每个操作提交给GPU也会产生微秒级的开销，这在越来越多的情况下变得越来越重要。

实际应用程序执行大量的GPU操作：一个典型的模式涉及许多迭代（或时间步），每个步骤中有多个操作。例如，分子系统的模拟在许多时间步长上迭代，其中每个分子的位置在每个步长都会根据其他分子施加在其上的力进行更新。对于精确建模自然的模拟技术，通常每个时间步长需要对应于多个GPU操作的多个算法阶段。如果这些操作中的每一个都单独启动到GPU，并迅速完成，那么开销可能会结合起来，形成显著的整体性能下降。

CUDA Graphs 的设计允许将工作定义为graph，而不是单个操作。它们通过提供通过单个CPU操作启动多个GPU操作的机制来解决上述问题，从而减少开销。在本文中，我们通过展示如何增强一个非常简单的示例来演示如何开始使用CUDA Graphs。

The Example
考虑一种情况，其中我们在每个时间步长内有一系列短GPU内核：

Loop over timesteps
    …
    shortKernel1
    shortKernel2
    …
    shortKernelN
    …

我们将创建一个模仿这种模式的简单代码。然后，我们将使用它来演示标准启动机制所涉及的开销，并演示如何引入包含多个内核的CUDA图，该图可以在单个操作中从应用程序启动。

首先，让我们编写一个计算内核，如下所示：

#define N 500000 // tuned such that kernel takes a few microseconds

__global__ void shortKernel(float * out_d, float * in_d){
  int idx=blockIdx.x*blockDim.x+threadIdx.x;
  if(idx<N) out_d[idx]=1.23*in_d[idx];
}

这只需从内存中读取浮点数字的输入数组，将每个元素乘以一个常数因子，然后将输出数组写回内存。此内核所花费的时间取决于数组大小，数组大小已设置为500000个元素，因此内核需要几微秒的时间。我们可以使用 profiler 来测量所需的时间为2.9μs，其中我们使用CUDA 10.1在NVIDIA Tesla V100 GPU上运行（我们已将每个块的线程数设置为512个线程）。在本文的剩余部分中，我们将固定这个内核，改变它的调用方式。

First Implementation with Multiple Launches

我们可以使用上述内核在模拟时间步长内模拟每个短内核，如下所示：

#define NSTEP 1000
#define NKERNEL 20

// start CPU wallclock timer
for(int istep=0; istep<NSTEP; istep++){
  for(int ikrnl=0; ikrnl<NKERNEL; ikrnl++){
    shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
    cudaStreamSynchronize(stream);
  }
}
//end CPU wallclock time

上面的代码片段调用内核20次，每次1000次迭代。我们可以使用基于CPU的wallclock timer来测量整个操作所花费的时间，并除以NSTEP*NKERNEL，每个内核的时间为9.6μs（包括开销）：远高于2.9μs的内核执行时间。

请注意，在每次内核启动后都存在cudaStreamSynchronize调用，这意味着在前一个内核完成之前，每个后续内核都不会启动。这意味着与每次启动相关的任何开销都将完全暴露：总时间将是内核执行时间加上任何开销的总和。我们可以使用Nsight Systems profiler直观地看到这一点：

这显示了时间线的一部分（时间从左到右增加），包括8个连续的内核启动。理想情况下，GPU应该以最小的空闲时间保持繁忙，但这里的情况并非如此。每个内核执行都可以在“CUDA（Tesla V100-SXM2-16G）”部分的图像底部看到。可以看出，在GPU空闲的情况下，每个内核执行之间都有很大的间隙。

我们可以通过查看“CUDA API”行来获得更多的见解，该行从CPU的角度显示了与GPU相关的活动。该行中的紫色条目对应于CPU线程在启动内核的CUDA API函数中所花费的时间，绿色条目是在与GPU同步的CUDA API函数中所耗费的时间，即等待内核在GPU上完全启动和完成。因此，内核之间的差距可以归因于CPU和GPU启动开销的组合。

请注意，在这个时间尺度上（我们检查的是非常短的事件），profiler增加了一些额外的启动开销，因此为了准确分析性能，应该使用基于CPU的wallclock计时器（就像我们在本文中所做的那样）。尽管如此，profiler还是有效地提供了对代码行为的深入了解。

Overlapping Kernel Launch and Execution

我们可以对上述代码进行简单但非常有效的改进，方法是将同步移出最内层的循环，使其只发生在每个时间步之后，而不是每次内核启动之后：

// start wallclock timer
for(int istep=0; istep<NSTEP; istep++){
  for(int ikrnl=0; ikrnl<NKERNEL; ikrnl++){
    shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
  }
  cudaStreamSynchronize(stream);
}
//end wallclock timer

内核仍将按顺序执行（因为它们在同一个流中），但这一更改允许在上一个内核完成之前启动内核，从而允许在内核执行之后隐藏启动开销。当我们这样做时，我们测量每个内核所花费的时间（包括开销）为3.8μs（而内核执行时间为2.9μs）。这一点得到了显著改善，但仍存在与多次发射相关的开销。

profiler现在显示：

可以看出，除了时间步结束时的调用之外，我们已经删除了绿色同步API调用。在每个时间步长内，可以看到启动开销现在能够与内核执行重叠，并且连续内核之间的间隙已经减少。但我们仍在为每个内核执行单独的启动操作，其中每个内核都忽略了其他内核的存在。

CUDA Graph Implementation

我们可以通过使用CUDA图在单个操作中启动每次迭代中的所有内核来进一步提高性能。

我们介绍一个图形如下：

bool graphCreated=false;
cudaGraph_t graph;
cudaGraphExec_t instance;
for(int istep=0; istep<NSTEP; istep++){
  if(!graphCreated){
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    for(int ikrnl=0; ikrnl<NKERNEL; ikrnl++){
      shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
    }
    cudaStreamEndCapture(stream, &graph);
    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
    graphCreated=true;
  }
  cudaGraphLaunch(instance, stream);
  cudaStreamSynchronize(stream);
}

新插入的代码允许通过使用CUDA Graph来执行。我们引入了两个新对象：类型为cudaGraph_t的graph包含定义图的结构和内容的信息；cudaGraphExec_t类型的instance是一个“可执行图”：以类似于单个内核的方式启动和执行的形式表示图。

因此，首先我们必须定义graph，并通过捕获在cudaStreamBeginCapture和cudaStreamEndCapture调用之间提交到stream的GPU活动的信息来实现这一点。然后，我们必须通过cudaGraphInstantate调用实例化图，该调用创建并预初始化所有内核工作描述符，以便它们可以尽可能快地重复启动。然后可以通过cudaGraphLaunch调用提交生成的实例以供执行。

至关重要的是，只需要捕获和实例化一次（在第一个时间步上），并在所有后续时间步上重复使用同一实例（此处由graphCreated布尔值上的条件语句控制）。

因此，我们现在有以下步骤：

• 第一步：
  • 创建和实例化图形
  • 启动图（包括20个内核）
  • 等待图形完成
• 对于剩余的999个步骤中的每一个
  • 启动图（包括20个内核）
  • 等待图形完成

测量这个完整过程所花费的时间，除以1000×20得出每个内核的有效时间（包括开销），得出3.4μs（而内核执行时间为2.9μs），因此我们成功地进一步降低了开销。请注意，在这种情况下，创建和实例化graph的时间相对较大，约为400μs，但这只执行了一次，因此这只会对我们的每个内核成本产生约0.02μs的影响。类似地，第一个graph的启动比所有后续的启动慢33%左右，但当多次重复使用同一个graph时，这就变得无关紧要了。初始化开销的严重程度显然取决于问题：通常，为了从图中获益，您需要重复使用相同的graph足够多次。许多现实世界中的问题都涉及大量的重复，因此适合使用graph。

剩余的开销是由于在GPU上启动每个graph所需的必要步骤，我们希望通过未来对CUDA的改进来进一步减少这些开销。我们有意不在这里显示任何配置文件，因为我们仍在研究CUDA Graph与配置文件工具的兼容性。对于当前的CUDA版本，概要文件将类似于“重叠内核启动和执行”中所示的概要文件，不同之处在于，对于每组20个内核执行，CUDA API行中只有一个“cudaGraphLaunch”条目，并且在CUDA API行中会在与graph创建和实例化相对应的一开始有额外条目。这20个内核中的每一个仍将显示为单独的条目，但为了提供这样的图片，profiler当前禁用了一些与图相关的优化。更准确的概要文件不会禁用任何优化，而是通过显示单个graph条目来表示每组20个内核。

Further Information

即使在上述非常简单的演示情况下（其中大部分开销已经通过重叠的内核启动和执行隐藏起来），也很高兴观察到CUDA图的好处，但当然，更复杂的情况提供了更多的节约机会。图支持多个交互流，不仅包括内核执行，还包括内存拷贝和在主机CPU上执行的函数，如CUDA示例中的simpleCUDAGraphs示例中更深入地所示。

本文中的示例使用流捕获机制来定义图，但也可以通过新提供的API调用显式地定义节点和依赖项–simpleCUDAGraphs示例演示了如何使用这两种技术实现相同的问题。此外，图形还可以跨越多个GPU。

在单个图中实现多个活动，而不是单独处理每个活动，最终会为CUDA提供更多信息，从而提供更多优化机会。欲了解更多信息，请参阅《编程指南》的CUDA图表部分，并观看GTC 2019谈话录音《CUDA：新功能及超越》。

补充代码：

#define N 500000 // tuned such that kernel takes a few microseconds

__global__ void shortKernel(double * out_d, double * in_d){
    int idx=blockIdx.x*blockDim.x+threadIdx.x;
    if(idx<N) {
        out_d[idx]=1.23*in_d[idx];
    }
}

#define NSTEP 1000
#define NKERNEL 20

int main(){
    //cuda graph
    const int M=sizeof(double) *N;
    double *out_h=(double*)malloc(M);
    double *in_h=(double*)malloc(M);

    for(int n=0;n<N;++n){
        in_h[n]=1;
    }
    double *out_d,*in_d;
    cudaMalloc((void **)&out_d,M);
    cudaMalloc((void **)&in_d,M);

    cudaMemcpy(out_d,out_h,M,cudaMemcpyHostToDevice);
    cudaMemcpy(in_d,in_h,M,cudaMemcpyHostToDevice);

    const int block_size=32;
    const int grid_size=N/block_size;    

    cudaStream_t stream;
	cudaStreamCreate(&stream);
    
    bool graphCreated=false;
    cudaGraph_t graph;
    cudaGraphExec_t instance;
    for(int istep=0; istep<NSTEP; istep++){
        if(!graphCreated){
            cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
            for(int ikrnl=0; ikrnl<NKERNEL; ikrnl++){
                shortKernel<<<grid_size, block_size, 0, stream>>>(out_d, in_d);
            }
            cudaStreamEndCapture(stream, &graph);
            cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
            graphCreated=true;
        }
    cudaGraphLaunch(instance, stream);
    cudaStreamSynchronize(stream);
    }

    free(out_h);
    free(in_h);
    cudaFree(out_d);
    cudaFree(in_d);

    cudaStreamDestroy(stream);
}