为什么有时相同的内核执行速度要慢 10 倍？

Question

这是代码：

// 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;

它运行相同的内核 3 次，多个帧。第二个内核的块数量减少了。 同步后，我测量以下时间：

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

各帧的结果不同，但它收敛到 2 种模式：

25.1095
3.23891
25.8181
25.1116
25.1136

和

23.074
27.5548
26.0465
23.0779
23.0758

或以图形方式：

有2个问题：

1. 为什么有时第二个内核（其工作量比其他两个内核少 10 倍）执行时间更长？
2. 为什么第二个和第三个内核都在第一个内核结束后开始？不是一个接着一个按顺序，也不是所有 3 个都重叠。

Answer 1

发生这种情况是因为您将第二次和第三次启动放在不同的（非默认）流中。这告诉 CUDA 运行时它们是独立的，并且由运行时按照它想要的任何顺序（取决于流优先级）来安排两者。该功能允许您同时运行多个网格。如果您希望两者按顺序运行，您可以将它们放在同一个流上，或者使用事件同步这两个流。

我猜想，较小的网格有时似乎执行时间更长，因为运行时随机决定完成第三次启动的运行块，然后再继续第二次启动。但看起来总运行时间保持不变。

这是使用事件进行基准测试时的一大缺点。当多个内核在不同的流上并行运行时，单次启动的时间测量将失去意义。也许您更应该查看来自 Nsight Systems 的跟踪。

为了完整性：您可以通过分配优先级来决定运行时确定哪个流的优先级。但这并没有改变手头的基准测试问题。