CUDA线程执行顺序

我有一个 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 and 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。为什么会发生这种情况？

然而，每个块中的线程顺序始终是0,1,2,3。为什么会发生这种情况？我也以为会是随机的

每块有 4 个线程，您只需启动一个warp每块。 Awarp http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture是 CUDA 中的执行（以及调度和资源分配）单元，而不是线程。目前，一个 warp 由 32 个线程组成。

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

那么问题就变成了，在这种情况下，CUDA 运行时如何调度这些“同时”的函数调用？这个问题的答案是未指定的，但它不是“随机的”。因此，warp 内操作的调度顺序在不同的运行中不会改变是合理的。

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