cuda - 在关于warp同步线程执行工作原理的直觉上挣扎

Question

我是CUDA的新手。我正在研究基本的并行算法，例如归约，以了解线程执行的工作方式。我有以下代码:

__global__ void
Reduction2_kernel( int *out, const int *in, size_t N )
{
    extern __shared__ int sPartials[];
    int sum = 0;
    const int tid = threadIdx.x;
    for ( size_t i = blockIdx.x*blockDim.x + tid;
          i < N;
          i += blockDim.x*gridDim.x ) {
        sum += in[i];
    }
    sPartials[tid] = sum;
    __syncthreads();

    for ( int activeThreads = blockDim.x>>1;
              activeThreads > 32;
              activeThreads >>= 1 ) {
        if ( tid < activeThreads ) {
            sPartials[tid] += sPartials[tid+activeThreads];
        }
        __syncthreads();
    }
    if ( threadIdx.x < 32 ) {
        volatile int *wsSum = sPartials;
        if ( blockDim.x > 32 ) wsSum[tid] += wsSum[tid + 32]; // why do we need this statement, any exampele please?
        wsSum[tid] += wsSum[tid + 16];  //how these statements are executed in paralle within a warp
        wsSum[tid] += wsSum[tid + 8];
        wsSum[tid] += wsSum[tid + 4];
        wsSum[tid] += wsSum[tid + 2];
        wsSum[tid] += wsSum[tid + 1];
        if ( tid == 0 ) {
            volatile int *wsSum = sPartials;// why this statement is needed?
            out[blockIdx.x] = wsSum[0];
        }
    }
}

不幸的是，我不清楚 if ( threadIdx.x < 32 )条件和之后的代码是如何工作的。有人可以给出一个带有线程ID以及如何执行语句的直观示例吗？我认为理解这些概念很重要，因此任何帮助都将是有帮助的!

Answer 1

让我们看一下代码块，并一路回答您的问题:

int sum = 0;
const int tid = threadIdx.x;
for ( size_t i = blockIdx.x*blockDim.x + tid;
      i < N;
      i += blockDim.x*gridDim.x ) {
    sum += in[i];
}

上面的代码遍历大小为 N的数据集。为了理解的目的，我们可以做一个假设，即 N> blockDim.x*gridDim.x，这最后一项只是网格中线程的总数。由于 N大于线程总数，因此每个线程都将数据集中的多个元素相加。从给定线程的角度来看，它是求和元素的间隔，这些元素由线程的网格尺寸( blockDim.x*gridDim.x)隔开。每个线程将其总和存储在名为 sum的局部(可能是寄存器)变量中。

sPartials[tid] = sum;
__syncthreads();

当每个线程结束时(即，因为它的for循环超过 N)，它将其中间 sum存储在共享内存中，然后等待该块中的所有其他线程结束。

for ( int activeThreads = blockDim.x>>1;
          activeThreads > 32;
          activeThreads >>= 1 ) {
    if ( tid < activeThreads ) {
        sPartials[tid] += sPartials[tid+activeThreads];
    }
    __syncthreads();
}

到目前为止，我们还没有讨论块的尺寸-无关紧要。假设每个块都有32个线程的整数倍。下一步将开始将存储在共享内存中的各种中间和收集到越来越小的变量组中。上面的代码从选择线程块( blockDim.x>>1)中的一半线程开始，并使用这些线程中的每一个在共享内存中合并两个部分和。因此，如果我们的线程块从128个线程开始，则仅使用其中的64个线程将128个部分和减少为64个部分和。此过程在for循环中重复进行，每次将线程切成两半并合并部分和，每个线程每次两个。只要 activeThreads> 32，该过程就会继续。因此，如果 activeThreads为64，则这64个线程会将128个部分和合并为64个部分和。但是，当 activeThreads变为32时，for循环终止，而没有将64个部分和合并为32。因此，在此代码块完成时，我们采用了(32个线程的任意倍数)线程块，并减少了许多部分和我们首先开始，直到64个。将256个部分和，128个部分和64个部分和相结合的过程必须在每次迭代时等待所有线程(在多个线程中)才能完成工作，因此 __syncthreads();语句在for循环的每次遍历中执行。

请记住，此时，我们已将线程块减少到64个部分和。

if ( threadIdx.x < 32 ) {

在此之后的内核其余部分，我们将仅使用前32个线程(即第一个线程)。所有其他线程将保持空闲状态。请注意，在这一点之后也没有 __syncthreads();，因为这将违反使用它的规则(所有线程都必须参与 __syncthreads();)。

    volatile int *wsSum = sPartials;

现在，我们正在创建一个指向共享内存的 volatile指针。从理论上讲，这告诉编译器不应进行各种优化，例如，将特定值优化到寄存器中。我们为什么以前不需要这个？因为 __syncthreads();也带有 a memory-fencing function。 __syncthreads();调用除了使所有线程彼此等待屏障之外，还强制所有线程更新返回共享或全局内存。但是，我们不再依赖于此功能，因为从现在开始，我们将不再使用 __syncthreads();，因为我们已将自己(对于内核的其余部分)限制为一次扭曲。

    if ( blockDim.x > 32 ) wsSum[tid] += wsSum[tid + 32]; // why do we need this

前一个约简块为我们提供了64个部分和。但是我们此时限制自己只能使用32个线程。因此，我们必须再进行一次组合，以将64个部分和汇总为32个部分和，然后才能进行余下的余数。

    wsSum[tid] += wsSum[tid + 16];  //how these statements are executed in paralle within a warp

现在，我们终于开始进行一些扭曲同步编程。这行代码取决于32个线程正在同步执行的事实。要了解原因(以及它如何工作)，将其分解为完成此代码行所需的操作序列将很方便。它看起来像:

    read the partial sum of my thread into a register
    read the partial sum of the thread that is 16 higher than my thread, into a register
    add the two partial sums
    store the result back into the partial sum corresponding to my thread

所有32个线程将按照上述顺序进行锁定。所有32个线程将通过将 wsSum[tid]读入(本地线程)寄存器开始。这意味着线程0读取 wsSum[0]，线程1读取 wsSum[1]等。此后，每个线程将另一部分和读取到不同的寄存器中:线程0读取 wsSum[16]，线程1读取 wsSum[17]，等等。的确，我们不在乎 wsSum[32] (及更高)值；我们已经将它们折叠到前32个 wsSum[]值中。但是，正如我们将看到的，在此步骤中，只有前16个线程将有助于最终结果，因此前16个线程将把32个部分和组合成16个。接下来的16个线程也将起作用，但他们只是在做垃圾工作-它将被忽略。

上面的步骤将32个部分和合并到 wsSum[]的前16个位置中。下一行代码:

    wsSum[tid] += wsSum[tid + 8];

以8的粒度重复此过程。同样，所有32个线程都处于 Activity 状态，并且微序列如下所示:

    read the partial sum of my thread into a register
    read the partial sum of the thread that is 8 higher than my thread, into a register
    add the two partial sums
    store the result back into the partial sum corresponding to my thread

因此，前8个线程将前16个部分和( wsSum[0..15])组合成8个部分和(包含在 wsSum[0..7]中)。接下来的8个线程也将 wsSum[8..23]合并为 wsSums[8..15]，但是对8..15的写入是在线程0..8读取这些值之后发生的，因此有效数据不会被破坏。这只是额外的垃圾工作。对于经线内的其他8个线程块也是如此。因此，在这一点上，我们已将部分利息和合并为8个位置。

    wsSum[tid] += wsSum[tid + 4];  //this combines partial sums of interest into 4 locations
    wsSum[tid] += wsSum[tid + 2];  //this combines partial sums of interest into 2 locations
    wsSum[tid] += wsSum[tid + 1];  //this combines partial sums of interest into 1 location

这些代码行遵循与前两行相似的模式，将warp分成8组，每组4个线程(只有第一个4线程组有助于最终结果)，然后将warp分为16组，每组2个线程，只有第一个2线程组有助于最终结果。最后，分成32组，每组1个线程，每个线程生成一个部分和，仅关注第一个部分和。

    if ( tid == 0 ) {
        volatile int *wsSum = sPartials;// why this statement is needed?
        out[blockIdx.x] = wsSum[0];
    }

最后，在上一步中，我们将所有部分和减少为单个值。现在是时候将单个值写到全局内存中了。我们完成减少了吗？也许，但是可能不会。如果上述内核仅用1个线程块启动，那么我们将可以完成-我们的最终“部分”和实际上是数据集中所有元素的和。但是，如果我们启动了多个块，则每个块的最终结果仍然是“部分”总和，并且所有块的结果必须加在一起(以某种方式)。

并回答您的最后一个问题？

我不知道为什么需要那句话。

我的猜测是它是在还原内核的先前迭代中遗留下来的，程序员忘记删除它，或者没有注意到它是不需要的。也许其他人会知道答案。

最后， cuda reduction sample提供了很好的参考代码供研究，随附的 pdf document很好地描述了可以在此过程中进行的优化。