最後一名後衛

考慮在某個任務上工作的網格，例如並行縮減。最初，每個塊可以獨立完成其工作，產生一些部分結果。然而，最後，部分結果需要組合並合併在一起。一個典型的例子是大資料的約簡演算法。

一種典型的方法是呼叫兩個核心，一個用於部分計算，另一個用於合併。但是，如果可以通過單個塊有效地完成合並，則只需要一個核心呼叫。這是通過 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++意義上的模板