最后一名后卫

考虑在某个任务上工作的网格，例如并行缩减。最初，每个块可以独立完成其工作，产生一些部分结果。然而，最后，部分结果需要组合并合并在一起。一个典型的例子是大数据的约简算法。

一种典型的方法是调用两个内核，一个用于部分计算，另一个用于合并。但是，如果可以通过单个块有效地完成合并，则只需要一个内核调用。这是通过 lastBlock 防护来实现的，定义如下：

Version >= 2.0

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

Version >= 1.1

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

使用这样的保护，最后一个块可以保证看到所有其他块产生的所有结果，并且可以执行合并。

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

假设：

• 计数器必须是全局内存指针，在调用内核之前初始化为 0。
• lastBlock 函数由所有块中的所有线程统一调用
• 内核在一维网格中调用（为了简化示例）
• T 命名你喜欢的任何类型，但该示例并非旨在成为 C++意义上的模板