CUDA Graphs学习与实验
- 一.参考链接
- 二.测试方案
- 三.测试代码
CUDA图(CUDA Graphs)为CUDA引入了一种全新的工作提交模型。它允许将一系列操作(如内核启动)以图的形式表示,并通过依赖关系将这些操作连接起来。这种图的定义过程与其执行过程是分开的,这意味着我们可以提前定义好一个图,然后多次重复执行它。
这种定义与执行的分离带来了多方面的优化:
- 降低CPU启动开销:相比传统的流(streams)方式,由于大量的设置工作已经在图的定义和实例化阶段完成,实际执行时的CPU开销明显减少。
- 全局优化机会:通过将整个工作流程以图的形式呈现给CUDA,CUDA有机会对整个流程进行优化。这在逐步提交工作的流机制中是无法实现的,因为流机制只能看到局部的、片段式的工作提交。
流机制中的问题:
在传统的流中,当你向流中放置一个内核时,主机驱动程序需要执行一系列操作来准备在GPU上执行该内核。这些操作包括设置内核参数、配置执行环境等。对于执行时间较短的GPU内核,这些准备工作的开销可能占到总执行时间的很大一部分,从而降低了整体效率。
CUDA图的工作提交分为三个阶段:
-
定义(Definition):
在这个阶段,程序创建一个包含操作及其依赖关系的图。开发者描述需要执行的操作(如内核函数)以及这些操作之间的先后顺序或并行关系。
-
实例化(Instantiation):
在定义完成后,CUDA对图进行实例化。实例化过程包括:
- 快照:对图模板进行捕获,生成一个具体的可执行图结构。
- 验证:检查图的正确性,确保所有的操作和依赖关系都是有效的。
- 预处理:执行大部分的设置和初始化工作,目的是尽可能减少在实际执行时需要完成的工作量。
实例化的结果是一个可执行图(executable graph)。
-
执行(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