CUDA Graphs学习与实验

CUDA图（CUDA Graphs）为CUDA引入了一种全新的工作提交模型。它允许将一系列操作（如内核启动）以图的形式表示，并通过依赖关系将这些操作连接起来。这种图的定义过程与其执行过程是分开的，这意味着我们可以提前定义好一个图，然后多次重复执行它。

这种定义与执行的分离带来了多方面的优化：

1. 降低CPU启动开销：相比传统的流（streams）方式，由于大量的设置工作已经在图的定义和实例化阶段完成，实际执行时的CPU开销明显减少。
2. 全局优化机会：通过将整个工作流程以图的形式呈现给CUDA，CUDA有机会对整个流程进行优化。这在逐步提交工作的流机制中是无法实现的，因为流机制只能看到局部的、片段式的工作提交。

流机制中的问题：

在传统的流中，当你向流中放置一个内核时，主机驱动程序需要执行一系列操作来准备在GPU上执行该内核。这些操作包括设置内核参数、配置执行环境等。对于执行时间较短的GPU内核，这些准备工作的开销可能占到总执行时间的很大一部分，从而降低了整体效率。

CUDA图的工作提交分为三个阶段：

1. 定义（Definition）：

   在这个阶段，程序创建一个包含操作及其依赖关系的图。开发者描述需要执行的操作（如内核函数）以及这些操作之间的先后顺序或并行关系。

2. 实例化（Instantiation）：

   在定义完成后，CUDA对图进行实例化。实例化过程包括：

   • 快照：对图模板进行捕获，生成一个具体的可执行图结构。
   • 验证：检查图的正确性，确保所有的操作和依赖关系都是有效的。
   • 预处理：执行大部分的设置和初始化工作，目的是尽可能减少在实际执行时需要完成的工作量。

   实例化的结果是一个可执行图（executable graph）。

3. 执行（Execution）：

   已实例化的可执行图可以像普通的CUDA工作一样被提交到流中执行。重要的是，这个可执行图可以被多次执行，而无需每次都重新实例化。这大大提高了执行的效率，特别是在需要重复执行相同操作的情况下。

CUDA图的优势：

• 性能提升：通过减少CPU的启动开销，特别是在需要频繁启动小型内核的情况下，CUDA图能够显著提升性能。
• 优化执行：由于CUDA能够提前知道整个工作流程，它可以进行全局优化。例如，它可以重新排列操作以提高并行性，或者优化内存传输以减少延迟。
• 简化编程模型：开发者可以以更直观的方式描述计算任务，而无需手动管理复杂的依赖关系和同步机制。

举例说明：

假设我们有一系列需要按特定顺序执行的内核操作。在传统的流机制中，我们需要：

• 为每个内核启动，都要进行一次完整的设置和启动过程。
• 手动管理这些内核之间的依赖关系，确保它们按正确的顺序执行。

使用CUDA图后，我们可以：

• 一次性地定义所有的内核操作和它们的依赖关系。
• 实例化后，CUDA会处理好所有的设置和依赖关系。
• 执行时，只需简单地启动可执行图即可。

结论：

CUDA图为GPU计算提供了更高效、更灵活的工作提交方式。通过预先定义和实例化计算图，CUDA能够减少CPU的开销，并利用全局信息对执行进行优化。这对于需要高性能计算的应用，尤其是包含大量小型、短时内核的应用，具有重要意义

一.参考链接

• graph management functions of the low-level CUDA driver api
• CUDA Runtime Graph API

二.测试方案

三.测试代码

tee cuda_graph.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>#define CHECK_CUDA(call)                                            \do {                                                            \cudaError_t err = call;                                    \if (err != cudaSuccess) {                                 \std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \exit(EXIT_FAILURE);                                    \}                                                         \} while (0)#define CHECK_CUDA_DRV_API(call)                      \do {                              \CUresult err = call;                  \if (err != CUDA_SUCCESS) {                 \char *error_str=new char[1024];  \cuGetErrorString(err,(const char**)&error_str); \printf("[%s:%d] %s Error :%s!\n",__FILE__,__LINE__,#call,error_str); \}                                      \} while (0)__global__ void Kernel1(float *a,float *b,float *c,float *d)
{unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;a[tid]=1;b[tid]=2;c[tid]=3;d[tid]=0;if(tid==0){printf("Kernel1\n");}
}__global__ void Kernel2(float *a,float *b,float *c,float *d)
{unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;a[tid]+=1;if(tid==0){printf("Kernel2\n");}
}__global__ void Kernel3(float *a,float *b,float *c,float *d)
{unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;b[tid]+=2;if(tid==0){printf("Kernel3\n");}
}__global__ void Kernel4(float *a,float *b,float *c,float *d)
{unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;c[tid]+=3;if(tid==0){printf("Kernel4\n");}
}__global__ void Kernel5(float *a,float *b,float *c,float *d)
{unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;d[tid]=a[tid]+b[tid]+c[tid];if(tid==0){printf("Kernel5\n");}
}struct HostFuncParam
{float *a;float *b;float *c;float *d;int thread_size;
};void CUDART_CB HostFunc(void *data){HostFuncParam *pstParam=(HostFuncParam*)data;for(int i=0;i<pstParam->thread_size;i++){pstParam->d[i]+=1;}printf("HostFunc\n");
}int run(bool graph_mode)
{int deviceid=0;    int block_count=1;int block_size=8;int thread_size=block_count*block_size;int total_count=thread_size*sizeof(float);cudaStream_t stream[3];cudaEvent_t event[3];CHECK_CUDA(cudaSetDevice(deviceid)); for(int i=0;i<3;i++){CHECK_CUDA(cudaStreamCreate(&stream[i]));CHECK_CUDA(cudaEventCreate(&event[i]));}float *a,*b,*c,*d;CHECK_CUDA(cudaMallocManaged(&a, total_count));CHECK_CUDA(cudaMallocManaged(&b, total_count));CHECK_CUDA(cudaMallocManaged(&c, total_count));CHECK_CUDA(cudaMallocManaged(&d, total_count));cudaGraph_t graph;if(graph_mode){CHECK_CUDA_DRV_API(cuGraphCreate(&graph, 0));CHECK_CUDA(cudaStreamBeginCapture(stream[0],cudaStreamCaptureModeGlobal));}    Kernel1<<<block_count, block_size,0,stream[0]>>>(a,b,c,d);CHECK_CUDA(cudaEventRecord(event[0], stream[0]));CHECK_CUDA(cudaStreamWaitEvent(stream[1], event[0]));CHECK_CUDA(cudaStreamWaitEvent(stream[2], event[0]));Kernel2<<<block_count, block_size,0,stream[0]>>>(a,b,c,d);Kernel3<<<block_count, block_size,0,stream[1]>>>(a,b,c,d);CHECK_CUDA(cudaEventRecord(event[1], stream[1]));Kernel4<<<block_count, block_size,0,stream[2]>>>(a,b,c,d);CHECK_CUDA(cudaEventRecord(event[2], stream[2]));CHECK_CUDA(cudaStreamWaitEvent(stream[0], event[1]));CHECK_CUDA(cudaStreamWaitEvent(stream[0], event[2]));Kernel5<<<block_count, block_size,0,stream[0]>>>(a,b,c,d);HostFuncParam stParam;stParam.d=d;stParam.thread_size=thread_size;CHECK_CUDA(cudaLaunchHostFunc(stream[0], HostFunc, (void*)&stParam));if(graph_mode){CHECK_CUDA(cudaStreamEndCapture(stream[0], &graph));cudaGraphExec_t graphExec;CHECK_CUDA(cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0));CHECK_CUDA(cudaGraphLaunch(graphExec, 0));CHECK_CUDA(cudaDeviceSynchronize());CHECK_CUDA_DRV_API(cuGraphDebugDotPrint(graph,"graph.dot",0));CHECK_CUDA(cudaGraphExecDestroy(graphExec));CHECK_CUDA(cudaGraphDestroy(graph));}else{CHECK_CUDA(cudaStreamSynchronize(stream[0]));}for(int i=0;i<thread_size;i++){printf("%6.2f\n",d[i]);}CHECK_CUDA(cudaFree(a));CHECK_CUDA(cudaFree(b));CHECK_CUDA(cudaFree(c));CHECK_CUDA(cudaFree(d));return 0;
}int main(int argc,char *argv[])
{int mode=atoi(argv[1]);if(mode==0){printf("normal mode\n");run(0);}else{printf("graph mode\n");run(1);}
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo \-o cuda_graph cuda_graph.cu \-I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./cuda_graph 0
./cuda_graph 1
dot -Tpng graph.dot -o graph.png

• 输出

normal mode
Kernel1
Kernel2
Kernel3
Kernel4
Kernel5
HostFunc13.0013.0013.0013.0013.0013.0013.0013.00
graph mode
Kernel1
Kernel2
Kernel3
Kernel4
Kernel5
HostFunc13.0013.0013.0013.0013.0013.0013.0013.00