CUDA同核函数执行时长异常及流启动时机问题咨询

测试代码 #

// default stream, 1000 blocks
cudaEventRecord(start1);
kernel_gpu<<<1000, dim3{32, 32, 1}>>>();
cudaEventRecord(stop1);

// stream1, 100 blocks
cudaEventRecord(start2, stream1);
kernel_gpu<<<100, dim3{32, 32, 1}, 0, stream1>>>();
cudaEventRecord(stop2, stream1);

// stream2, 1000 blocks
cudaEventRecord(start3, stream2);
kernel_gpu<<<1000, dim3{32, 32, 1}, 0, stream2>>>();
cudaEventRecord(stop3, stream2);

cudaDeviceSynchronize();

float t;

// first kernel time
cudaEventElapsedTime(&t, start1, stop1);
std::cout << t << std::endl;

// second kernel time
cudaEventElapsedTime(&t, start2, stop2);
std::cout << t << std::endl;

// third kernel time
cudaEventElapsedTime(&t, start3, stop3);
std::cout << t << std::endl;

// start time of the second kernel relative to the first
cudaEventElapsedTime(&t, start1, start2);
std::cout << t << std::endl;

// start time of the second kernel relative to the first
cudaEventElapsedTime(&t, start1, start3);
std::cout << t << std::endl << std::endl;

测试场景 #

在多帧场景下三次运行同一kernel：

• 第一次：使用默认流，包含1000个block
• 第二次：使用stream1，包含100个block
• 第三次：使用stream2，包含1000个block

同步后测量以下时长：

1. 第一个kernel的执行时长
2. 第二个kernel的执行时长
3. 第三个kernel的执行时长
4. 第二个kernel相对第一个的启动时间
5. 第三个kernel相对第一个的启动时间

测试结果 #

测试结果随帧变化，最终收敛为两种模式：

模式一 #

25.1095
3.23891
25.8181
25.1116
25.1136

模式二 #

23.074
27.5548
26.0465
23.0779
23.0758

技术问题 #

1. 为何block数量仅为另外两个1/10的第二个kernel，有时执行时长反而更长？
2. 为何第二个和第三个kernel均在第一个kernel结束后才启动，既不重叠执行也不顺序启动？

问题解答 #

问题1解答 #

第二个kernel执行时长偶尔更长，核心原因是GPU任务调度延迟与资源竞争：

• 默认流的大kernel（1000个block）会占用几乎所有SM（流式多处理器）资源，stream1的小kernel需要等待SM释放资源才能启动。而cudaEventRecord记录的是从事件触发到执行结束的完整时间，包含了等待资源的耗时，最终测量时长自然被拉长。
• 多帧场景下前一帧的残留任务、系统级GPU资源波动等因素，也会放大小kernel的调度延迟，进一步推高测量时长。

问题2解答 #

这是因为CUDA默认流的隐式同步特性：

• 在未开启特定调度策略的情况下，默认流（空流）会与所有显式流进行隐式同步——所有显式流中的任务，必须等待默认流已提交的任务完成后才能启动。
• 代码中先提交了默认流的kernel，后续stream1和stream2的任务受隐式同步限制，只能等默认流kernel执行完毕后才启动，因此无法与默认流任务重叠，也不会提前顺序启动。

内容的提问来源于stack exchange，提问作者nikitablack