You need to enable JavaScript to run this app.
最新活动
大模型
产品
解决方案
定价
生态与合作
支持与服务
开发者
了解我们

CUDA Graph执行耗时高于原核函数启动循环的问题求助

CUDA Graph执行耗时高于原核函数启动循环的问题求助

各位大佬好,我最近在做CUDA程序的性能优化,遇到了一个特别费解的问题,想过来请教下大家。

先给大家说下我的原始代码逻辑:我写了一个循环,里面会启动多个有依赖关系的核函数,用了两个流(stream1和stream2)加上事件来做同步,具体代码如下:

for (int i= 1; i<= 1024 ; i++) {
    // 主流操作
    kernel1<<<1,512,0,stream1>>>(i , /*其他参数*/);

    // 分流到stream2执行
    cudaEventRecord(event1, stream1);
    cudaStreamWaitEvent(stream2, event1, 0);

    kernel2<<<1,512,0,stream1>>>(i , /*其他参数*/);
    kernel3<<<gridDim, blockDim, 0, stream2>>>(i , /*其他参数*/);

    // 将stream2的结果同步回主流
    cudaEventRecord(event2, stream2);
    cudaStreamWaitEvent(stream1, event2, 0);
}

因为循环里要反复启动核函数,开销有点高,我就想着用CUDA Graph来优化——毕竟CUDA Graph就是用来减少这类重复启动开销的对吧?

但我的核函数里有动态参数(就是循环里的i,每次循环都会变),而且还有跨流的依赖关系,之前我就纠结过怎么用Graph处理这种场景。现在我好不容易把Graph搭起来跑通了,结果发现Graph的执行耗时居然比原来的循环还高!这完全和我预期的优化方向相反,实在搞不懂哪里出问题了。

有没有大佬能帮我分析下可能的原因?比如是不是我捕获Graph的时候没处理好跨流的依赖逻辑?或者动态参数的传递方式不对,反而引入了额外的开销?麻烦大家给我支支招,谢谢啦!

备注:内容来源于stack exchange,提问作者Photos

火山引擎 最新活动