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

Question

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

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

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

编辑：再一次，我很清楚

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

并且这些块是完全独立的——它们可以按任何顺序运行并产生相同的输出。我只是在问我将块排列到网格中的顺序是否会影响最终同时运行的块，因为这确实会通过 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 博士在链接文章中报告的结果相符。然而，对于二维网格，我没有发现任何证据表明在块执行顺序中存在空间填充曲线，因此这可能在费米和帕斯卡之间发生了变化。

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

附录

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

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

Answer 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.

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