CUDA同一个Warp中的线程执行顺序

最新推荐文章于 2022-10-26 13:54:24 发布

小肥兔

最新推荐文章于 2022-10-26 13:54:24 发布

阅读量793

点赞数

分类专栏： CUDA/GPU 文章标签： c++ c语言

原文链接：http://cn.voidcc.com/question/p-esqsbeeh-bnk.html

版权

CUDA/GPU 专栏收录该内容

4 篇文章 2 订阅

订阅专栏

#include <stdio.h> 

#define NUM_BLOCKS 4 
#define THREADS_PER_BLOCK 4 

__global__ void hello() 
{ 

    printf("Hello. I'm a thread %d in block %d\n", threadIdx.x, blockIdx.x); 

} 


int main(int argc,char **argv) 
{ 
    // launch the kernel 
    hello<<<NUM_BLOCKS, THREADS_PER_BLOCK>>>(); 

    // force the printf()s to flush 
    cudaDeviceSynchronize(); 

    return 0; 
}

，其中每个线程，将打印threadIdx.x和blockIdx.x。该计划的一个可能的输出是这样的：

Hello. I'm a thread 0 in block 0 
Hello. I'm a thread 1 in block 0 
Hello. I'm a thread 2 in block 0 
Hello. I'm a thread 3 in block 0 
Hello. I'm a thread 0 in block 2 
Hello. I'm a thread 1 in block 2 
Hello. I'm a thread 2 in block 2 
Hello. I'm a thread 3 in block 2 
Hello. I'm a thread 0 in block 3 
Hello. I'm a thread 1 in block 3 
Hello. I'm a thread 2 in block 3 
Hello. I'm a thread 3 in block 3 
Hello. I'm a thread 0 in block 1 
Hello. I'm a thread 1 in block 1 
Hello. I'm a thread 2 in block 1 
Hello. I'm a thread 3 in block 1

运行程序几次，我得到了类似的结果，除了块顺序是随机的。例如，在上面的输出中，我们有这个块顺序0,2,3,1。再次运行问题我得到1,2,3,0。这是预期的。但是，每个块中的线程顺序始终为0,1,2,3。这是为什么发生？我认为这也是随机的。

我试图改变我的代码，强制每个块中的线程0花费更长的时间执行。我这样做是这样的：

__global__ void hello() 
{ 

    if (threadIdx.x == 0) 
    { 
     int k = 0; 
     for (int i = 0; i < 1000000; i++) 
     { 
      k = k + 1; 
     } 
    } 

    printf("Hello. I'm a thread %d in block %d\n", threadIdx.x, blockIdx.x); 

}

我希望线程为了1,2,3，0。然而，我得到了类似的结果，在一个我在上面，其中线顺序总是0显示，1 ，2，3.为什么会发生这种情况？

回答：

由于每块4个线程你只推出每块一个经。 A warp是CUDA中的执行单元（以及调度和资源分配），而不是线程。目前，一个warp包含32个线程。

这意味着每个块的所有4个线程（因为在这种情况下没有条件行为）正在执行锁步。当它们到达printf函数调用时，它们都在同一行代码,中以锁步执行对该函数的调用。

因此，在这种情况下，CUDA运行时如何调度这些“同时”函数调用？这个问题的答案没有具体说明，但它不是“随机的”。因此，对于一个warp内的操作的调度顺序不会随着运行而改变是合理的。

如果启动足够的线程为每个块创建多个warp，并且可能还包含一些其他代码以分散或“随机化”warp之间的行为，则应该能够看到printf操作是从“随机顺序。

问：

谢谢，这是对我的问题的非常准确的答案。我将其标记为答案，因为我学到了很多东西。但是，如果我将代码分散或“随机化”，会发生什么？我有兴趣进一步了解锁步部分。所以假设线程0进入一个if语句，并且线程1到31不（在同一个warp中）。假设线程1-31已准备好继续，它们是否会等待线程0完成if语句并以锁步方式继续？再次感谢您的回答。

回答：

是的，对于if-then-else构造，满足if条件的warp中的所有线程将遵循if路径，而其余线程为“inactive”。然后，满足if条件的线程将变为“不活动”，并且不满足if条件的线程将被激活，并且这些线程将遵循else路径。在完成这个序列时，warp中的所有线程都会隐式地重新同步，然后继续以lockstep执行。（事实上，if路径首先发生还是其他路径首先发生的确切顺序是未指定的。