代码之家 › 专栏 › 技术社区 › KQS

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

cuda

KQS · 技术社区 · 7 年前

这两者之间有什么关系吗 blockIdx 以及在GPU设备上执行线程块的顺序?

我的动机是我有一个内核,其中多个块将从全局内存中的同一位置读取,如果这些块将同时运行就好了(因为二级缓存命中很好)。在决定如何将这些块组织成网格时,可以这样说吗 blockIdx.x=0 更可能与同时运行 blockIdx.x=1 blockIdx.x=200 ? 我应该尝试将连续索引分配给从全局内存中相同位置读取的块?

明确地说,我不是在问块间依赖关系(如 this question )从程序正确性的角度来看,线程块是完全独立的。我已经在使用共享内存在一个块内广播数据,我不能使块再大了。

编辑:再一次,我很清楚

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

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

2 回复 | 直到 7 年前

Community CDub 4 年前

我发现了一篇文章,其中一位CS研究人员使用微基准测试在费米设备上对块调度器进行反向工程:

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

我修改了他的代码,使其在我的GPU设备(GTX 1080,带有Pascal GP104 GPU)上运行,并随机化运行时。

方法

clock64() )然后随机运行一段时间(适当地说,任务是使用带进位乘法算法生成随机数)。

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

后果

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

我们有40个块立即开始(每个短消息2个块*20个短消息),后续块在前一个块结束时开始。

对于二维网格,我发现相同的线性序列顺序 blockIdx.x 成为快速维度和 blockIdx.y

注:我在给这些情节贴标签时打了一个很糟糕的错别字。“threadIdx”的所有实例都应替换为“blockIdx”。

对于三维块网格:

结论

对于一维网格,这些结果与Pai博士在链接的writeup中报告的结果相匹配。然而,对于二维网格,我没有发现任何证据表明在块执行顺序中存在空间填充曲线,因此这可能在费米和帕斯卡之间的某个地方发生了变化。

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

附录

以下是显示随机运行时与固定运行时结果的曲线图,供参考:

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

tera 7 年前

是的,这肯定是有关联的(尽管当然不能保证)。

你最好在你的设备上试用一下。您可以使用 %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的共享内存来让结果更有趣——你应该用它的实际启动配置来代替你感兴趣的内核。

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.

你看,确实有很强的相关性。