CUDA Graph(创建静态图加快后续多次启动速度)

博客原文

应用场景:多次迭代(每次迭代执行的任务图都一样)

假设:shortKernel是一个运行时间极短的kernel

1. 初始版本:

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

总共平均耗时9.6μs;kernel执行耗时2.9us;

缺点:启动kernel-->执行kernel-->等待执行完;

2. 优化版本:

// 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;kernel执行耗时2.9us;

优点:启动下一个kernel和执行上一个kernel,能够并行起来;

缺点:每个kernel还得启动一次;

3. Graph优化版本:

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

总共平均耗时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

 

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值