CUDA线程执行顺序

Question

我有以下 CUDA 程序代码：

#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。这是为什么发生了什么？

Answer 1

However, the thread order in every block is always 0,1,2,3. Why is this happening? I thought it would be random too

每个块有 4 个线程，您每个块只启动一个 warp。 warp 是 CUDA 中的执行单元（以及调度和资源分配），而不是线程。目前，一个 warp 由 32 个线程组成。

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

所以问题就变成了，在这种情况下，CUDA 运行time 如何调度这些 "simultaneous" 函数调用？该问题的答案未指定，但不是 "random"。因此，warp 内操作的调度顺序不会从运行运行.

改变是合理的

如果您启动足够多的线程来为每个块创建多个 warp，并且可能还包含一些其他代码来分散和/或 "randomize" warp 之间的行为，您应该能够看到 printf 操作由按 "random" 顺序发生的单独扭曲发出。

Answer 2

要回答问题的第二部分，当控制流在 if 语句处发生分歧时，threadIdx.x != 0 所在的线程只需在 if 语句之后的汇合点等待.在线程 0 完成 if 块之前，它们不会继续执行 printf 语句。

CUDA线程执行顺序

CUDA thread execution order

cuda