blockIdx 是不是与块执行的顺序相关？

Posted 2023-04-15

【中文标题】blockIdx 是不是与块执行的顺序相关？

【英文标题】：Is blockIdx correlated to the order of block execution?blockIdx 是否与块执行的顺序相关？

【发布时间】：2018-03-21 11:11:18

【问题描述】：

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

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

明确地说，我不是在询问块间依赖关系（如this question），从程序正确性的角度来看，线程块是完全独立的。我已经在使用共享内存来广播块内的数据，我无法让块变得更大。

编辑：再一次，我很清楚

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

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

【问题讨论】：

来自CUDA Programming Guide：“线程块需要独立执行：必须能够以任何顺序并行或串行执行它们。”但是您可以使用一个全局变量，为刚开始的每个新块增加一个全局变量，并将该变量用作您的“块 id”。我很确定在 SO 上至少有一个关于如何在此处执行此操作的问题。

这不是我要问的，我已经编辑了这个问题以更清楚地说明这一点。正如您所建议的那样，我并没有试图确定实际的执行顺序；我在问内置变量blockIdx 是否与发出执行块的顺序有任何关系。

许多硬件特性让我相信无法预测任何连接。例如，在某个特定 sm 上调度块没有规则，它还取决于其他进程、可能同时运行的内核，甚至可能是内存操作。

【参考方案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 博士在链接文章中报告的结果相符。然而，对于二维网格，我没有发现任何证据表明在块执行顺序中存在空间填充曲线，因此这可能在费米和帕斯卡之间发生了变化。

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

附录

作为参考，下面的图表显示了随机运行时与固定运行时的结果：

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

【讨论】：

这个答案比我的要全面得多，这是我两个小时前输入的，但直到与同事共进晚餐后才提交 - 赞成。

所有这些努力都是为了确定明天可能会改变的事情。或者你可以接受@Shadow 给出的提示，并简单地保证块执行顺序，用你自己的块索引代替机器交给你的那个。

@RobertCrovella 老实说，直到你的评论让我以不同的眼光再次阅读它之前，我不明白影子想说什么。如果您发布一个原子块计数器作为答案，我会接受它。

【参考方案2】：

是的，肯定存在相关性（当然不能保证）。

您最好在您的设备上试用一下。您可以使用 %globaltimer 和 %smid 特殊 PTX 寄存器和一些内联汇编：

#include <stdio.h>

__managed__ unsigned long long starttime;

__device__ unsigned long long globaltime(void)

    unsigned long long time;
    asm("mov.u64  %0, %%globaltimer;" : "=l"(time));
    return time;


__device__ unsigned int smid(void)

    unsigned int sm;
    asm("mov.u32  %0, %%smid;" : "=r"(sm));
    return sm;


__global__ void logkernel(void)

    unsigned long long t = globaltime();
    unsigned long long t0 = atomicCAS(&starttime, 0ull, t);
    if (t0==0) t0 = t;
    printf("Started block %2u on SM %2u at %llu.\n", blockIdx.x, smid(), t - t0);



int main(void)

    starttime = 0;
    logkernel<<<30, 1, 49152>>>();
    cudaDeviceSynchronize();

    return 0;

我使用了 48K 的共享内存来使结果更有趣 - 你应该用它的实际启动配置替换你感兴趣的内核。

如果我在装有 GTX 1050 的笔记本电脑上运行此代码，我会得到如下输出：

Started block  1 on SM  1 at 0.
Started block  6 on SM  1 at 0.
Started block  8 on SM  3 at 0.
Started block  0 on SM  0 at 0.
Started block  3 on SM  3 at 0.
Started block  5 on SM  0 at 0.
Started block  2 on SM  2 at 0.
Started block  7 on SM  2 at 0.
Started block  4 on SM  4 at 0.
Started block  9 on SM  4 at 0.
Started block 10 on SM  3 at 152576.
Started block 11 on SM  3 at 152576.
Started block 18 on SM  1 at 153600.
Started block 16 on SM  1 at 153600.
Started block 17 on SM  0 at 153600.
Started block 14 on SM  0 at 153600.
Started block 13 on SM  2 at 153600.
Started block 12 on SM  2 at 153600.
Started block 19 on SM  4 at 153600.
Started block 15 on SM  4 at 153600.
Started block 20 on SM  0 at 210944.
Started block 21 on SM  3 at 210944.
Started block 22 on SM  0 at 211968.
Started block 23 on SM  3 at 211968.
Started block 24 on SM  1 at 214016.
Started block 26 on SM  1 at 215040.
Started block 25 on SM  2 at 215040.
Started block 27 on SM  2 at 215040.
Started block 28 on SM  4 at 216064.
Started block 29 on SM  4 at 217088.

所以你看到确实有很强的相关性。

【讨论】：

您的代码比我的要优雅得多，我很高兴我们得到了相同的结果。现在我了解到您可以从内核内部printf！感谢您发布代码。