cuda
ブロック間通信

備考

ラストブロックガード

いくつかのタスク、例えば並列削減に取り組んでいるグリッドを考えてみましょう。最初に、各ブロックは独立して作業を行うことができ、部分的な結果が得られます。しかし、最終的に部分的な結果を結合して併合する必要があります。典型的な例は、大きなデータに対する削減アルゴリズムです。

典型的なアプローチは、2つのカーネルを呼び出すことです.1つは部分計算用、もう1つはマージ用です。ただし、1つのブロックで効率的にマージを行うことができれば、カーネルコールは1回だけ必要です。これはlastBlockガードによって定義されます：

__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);
}

__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 ++の意味でのテンプレートではありません

グローバル作業キュー

一連の作業項目を考えてみましょう。各作業項目の完了に必要な時間は大きく異なります。ブロック間の作業分布を均衡させるために、各ブロックが前のものが完了したときにのみ次のアイテムをフェッチすることが賢明かもしれない。これは、先験的にブロックに項目を割り当てることとは対照的です。

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

前提条件：

• WorkQueueオブジェクトとgItem配列は、グローバルメモリに常駐しています
• そこからフェッチしているカーネルのWorkQueueオブジェクトには、新しい作業項目は追加されません
• WorkItemは、作業割り当ての小さな表現です。たとえば、別のオブジェクトへのポインタ
• WorkItem::none() staticメンバ関数は、何の作業も全く表さないWorkItemオブジェクトを作成しWorkItem
• WorkQueue::fetch()は、ブロック内のすべてのスレッドによって一様に呼び出される必要があります
• WorkQueue::fetch()別の__syncthreads()持たずにWorkQueue::fetch()を2回呼び出すことはありません。それ以外の場合は、競合状態が表示されます！

この例には、 WorkQueue初期化するWorkQueueや作成する方法は含まれていません。これは別のカーネルやCPUコードによって実行され、かなり簡単です。