CUDA Graph(创建静态图加快后续多次启动速度)
博客原文
应用场景:多次迭代(每次迭代执行的任务图都一样)
假设:shortKernel是一个运行时间极短的kernel
1. 初始版本:
define NSTEP 1000
#define NKERNEL 20// start CPU wallclock timer
for(int istep=0; istep>>(out_d, in_d);cudaStreamSynchronize(stream);}
}
//end CPU wallclock time
总共平均耗时9.6μs;kernel执行耗时2.9us;
缺点:启动kernel-->执行kernel-->等待执行完;
2. 优化版本:
// start wallclock timer
for(int istep=0; istep>>(out_d, in_d);}cudaStreamSynchronize(stream);
}
//end wallclock timer
总共平均耗时3.8μs;kernel执行耗时2.9us;
优点:启动下一个kernel和执行上一个kernel,能够并行起来;
缺点:每个kernel还得启动一次;
3. Graph优化版本:
bool graphCreated=false;
cudaGraph_t graph;
cudaGraphExec_t instance;
for(int istep=0; istep>>(out_d, in_d);}cudaStreamEndCapture(stream, &graph);cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);graphCreated=true;}cudaGraphLaunch(instance, stream);cudaStreamSynchronize(stream);
}
总共平均耗时3.4μs;kernel执行耗时2.9us;
优点:整个graph启动一次;头一次构建graph慢,但是后面的迭代就可以复用该graph了;
simpleCUDAGraphs例子里,是DAG图构建后反复启动;
DAG包含节点和依赖;节点支持GPU kernel、CPU<-->GPU内存copy、CPU函数;支持多stream,支持跨多GPU卡;
参考资料:
CUDA Graphs section of the Programming Guide
CUDA: New Features and Beyond
本文来自互联网用户投稿,文章观点仅代表作者本人,不代表本站立场,不承担相关法律责任。如若转载,请注明出处。 如若内容造成侵权/违法违规/事实不符,请点击【内容举报】进行投诉反馈!
