blockIdx 是否与块执行顺序相关？

Question

blockIdx线程块在 GPU 设备上执行的顺序和顺序有什么关系吗？

我的动机是我有一个内核，其中多个块将从全局内存中的同一位置读取，如果这些块可以并发运行会很好（因为 L2 缓存命中很好）。在决定如何将这些块组织到网格中时，是否可以肯定地说blockIdx.x=0与 并发运行的可能性blockIdx.x=1大于与blockIdx.x=200？并且我应该尝试为从全局内存中的相同位置读取的块分配连续索引？

需要明确的是，我不是在询问块间依赖关系（如在这个问题中），并且从程序正确性的角度来看，线程块是完全独立的。我已经在使用共享内存在一个块内广播数据，我不能让这些块变得更大。

编辑：再次，我很清楚

线程块需要独立执行：必须可以以任何顺序并行或串行执行它们。

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

Answer 1

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

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

我修改了他的代码以在我的 GPU 设备（GTX 1080，使用 Pascal GP104 GPU）上运行并随机化运行时间。

方法

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

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

结果

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

我们有 40 个块立即开始（每个 SM 2 个块 * 20 个 SM），随后的块在前一个块结束时开始。

对于二维网格，我发现了相同的线性序列顺序，blockIdx.x分别是快维度和blockIdx.y慢维度：

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

对于 3-d 块网格：

结论

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

当然，基准测试的常见警告适用，并且不能保证这不是特定于特定处理器模型的。

附录

作为参考，这是一个显示随机与固定运行时间的结果的图：

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