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




