cuda - blockIdx 是否与 block 执行顺序相关？

Question

blockIdx与线程 block 在GPU设备上的执行顺序有关系吗？

我的动机是我有一个内核，其中多个 block 将从全局内存中的同一位置读取，如果这些 block 能够同时运行就更好了(因为 L2 缓存命中很好)。在决定如何将这些 block 组织成网格时，是否可以安全地说 blockIdx.x=0 比 blockIdx.x=1 更有可能同时运行使用 blockIdx.x=200？我应该尝试将连续的索引分配给从全局内存中的相同位置读取的 block ？

需要明确的是，我不是在询问 block 间依赖性(如 this question 中所示)，并且从程序正确性的角度来看，线程 block 是完全独立的。我已经在使用共享内存在一个 block 内广播数据，但我无法让这些 block 变大。

编辑:同样，我很清楚

Thread blocks are required to execute independently: It must be possible to execute them in any order, in parallel or in series.

并且 block 是完全独立的——它们可以以任何顺序运行并产生相同的输出。我只是问我将 block 排列到网格中的顺序是否会影响哪些 block 最终同时运行，因为这确实会通过 L2 缓存命中率影响性能。

Answer 1

我发现了一篇文章，其中一位 CS 研究人员使用微基准测试对 Fermi 设备上的 block 调度程序进行逆向工程:

http://cs.rochester.edu/~sree/fermi-tbs/fermi-tbs.html

我调整了他的代码以在我的 GPU 设备(GTX 1080，使用 Pascal GP104 GPU)上运行并随机化运行时间。

方法

每个 block 只包含 1 个线程，并且启动时有足够的共享内存，每个 SM 只能驻留 2 个 block 。内核记录它的开始时间(通过 clock64() 获得)，然后运行一段随机时间(该任务足够恰当，是使用进位乘法算法生成随机数)。

GTX 1080 由 4 个图形处理集群 (GPC) 组成，每个集群有 5 个流式多处理器 (SM)。每个 GPC 都有自己的时钟，所以我使用链接中描述的相同方法来确定哪些 SM 属于哪些 GPC，然后减去一个固定的偏移量以将所有时钟值转换为相同的时区。

结果

对于一维 block 网格，我发现 block 确实是按连续顺序启动的:

我们有 40 个 block 立即开始(每个 SM 2 个 block * 20 个 SM)，后续 block 在前一个 block 结束时开始。

对于二维网格，我发现了相同的线性顺序，即 blockIdx.x作为快速维度和blockIdx.y慢维度:

注意:我在标记这些图时犯了一个严重的错字。 “threadIdx”的所有实例都应替换为“blockIdx”。

对于 3 维 block 网格:

结论

对于一维网格，这些结果与 Pai 博士在链接的文章中报告的结果相符。然而，对于 2-D 网格，我没有发现 block 执行顺序中空间填充曲线的任何证据，因此这可能在费米和帕斯卡之间的某个地方发生了变化。

当然，基准测试的常见注意事项也适用，并且不能保证这不是针对特定处理器型号的。

附录

作为引用，下面的图表显示了随机运行时间与固定运行时间的结果:

我们在随机运行时看到这种趋势这一事实让我更有信心这是一个真实的结果，而不仅仅是基准测试任务的一个怪癖。