ポスターはすでに彼自身の問題に対する答えを見つけています。それでも、以下のコードでは、CUDAにクリティカルセクションを実装するための一般的なフレームワークを提供しています。より詳細には、コードはブロックカウントを実行しますが、クリティカルセクションで実行される他の操作をホストするように簡単に変更できます。以下では、コードの説明と、CUDAのクリティカルセクションの実装における「典型的な」間違いについても報告します。
コード
#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS 512
#define NUMTHREADS 512 * 2
/***************/
/* LOCK STRUCT */
/***************/
struct Lock {
int *d_state;
// --- Constructor
Lock(void) {
int h_state = 0; // --- Host side lock state initializer
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state
}
// --- Destructor
__host__ __device__ ~Lock(void) {
#if !defined(__CUDACC__)
gpuErrchk(cudaFree(d_state));
#else
#endif
}
// --- Lock function
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
// --- Unlock function
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
/*************************************/
/* BLOCK COUNTER KERNEL WITHOUT LOCK */
/*************************************/
__global__ void blockCountingKernelNoLock(int *numBlocks) {
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
}
/**********************************/
/* BLOCK COUNTER KERNEL WITH LOCK */
/**********************************/
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) {
if (threadIdx.x == 0) {
lock.lock();
numBlocks[0] = numBlocks[0] + 1;
lock.unlock();
}
}
/****************************************/
/* BLOCK COUNTER KERNEL WITH WRONG LOCK */
/****************************************/
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) {
lock.lock();
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
lock.unlock();
}
/********/
/* MAIN */
/********/
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
// --- Unlocked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelNoLock << <NUMBLOCKS, NUMTHREADS >> >(d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the unlocked case: %i\n", h_counting);
// --- Locked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelLock << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the locked case: %i\n", h_counting);
gpuErrchk(cudaFree(d_counting));
}
コードの説明
クリティカルセクションは、CUDAスレッドによって順番に実行される必要がある一連の操作です。
スレッドグリッドのスレッドブロックの数を計算するタスクを持つカーネルを構築するとします。考えられるアイデアの1つは、各ブロックの各スレッドにthreadIdx.x == 0グローバルカウンターを増やすことです。競合状態を防ぐために、すべての増加は順番に発生する必要があるため、クリティカルセクションに組み込む必要があります。
上記のコードには、との2つのカーネル関数がblockCountingKernelNoLockありblockCountingKernelLockます。前者は、カウンターを増やすためにクリティカルセクションを使用せず、ご覧のとおり、間違った結果を返します。後者は、クリティカルセクション内のカウンターの増加をカプセル化するため、正しい結果を生成します。しかし、クリティカルセクションはどのように機能しますか？
クリティカルセクションは、グローバルステートによって管理されd_stateます。最初の状態は0です。さらに、2つの__device__メソッド、、はこの状態を変更できますlock。andメソッドunlockは、各ブロック内の単一のスレッド、特にローカルスレッドインデックスを持つスレッドによってのみ呼び出すことができます。lockunlockthreadIdx.x == 0
実行中にランダムに、たとえば、ローカルスレッドインデックスthreadIdx.x == 0とグローバルスレッドインデックスを持つスレッドの1つがt、メソッドを最初に呼び出すことになりlockます。特に、起動しatomicCAS(d_state, 0, 1)ます。最初d_state == 0から、d_stateはに更新され1、atomicCASはに戻り0、スレッドはlock関数を終了して、更新命令に渡されます。その間、そのようなスレッドは前述の操作を実行しますが、他のすべてのブロックの他のすべてのスレッドはメソッドthreadIdx.x == 0を実行しlockます。d_stateただし、に等しい値が見つかるため、1更新atomicCAS(d_state, 0, 1)は実行されずに返さ1れるため、これらのスレッドはwhileループを実行したままになります。そのスレッドの後t更新を完了すると、unlock関数、つまりatomicExch(d_state, 0)、が実行され、に復元さd_stateれ0ます。この時点で、ランダムに、別のスレッドthreadIdx.x == 0が状態を再びロックします。
上記のコードには、3番目のカーネル関数、つまり。も含まれていますblockCountingKernelDeadlock。ただし、これはクリティカルセクションの別の誤った実装であり、デッドロックにつながります。実際、ワープはロックステップで動作し、すべての命令の後に同期することを思い出します。したがって、を実行するblockCountingKernelDeadlockと、ワープ内のスレッドの1つ、たとえばローカルスレッドインデックスを持つスレッドt≠0が状態をロックする可能性があります。この状況では、と同じワープのスレッドをt含む、の同じワープ内の他のスレッドは、スレッドthreadIdx.x == 0と同じwhileループステートメントを実行しtます。これは、ロックステップで実行される同じワープのスレッドの実行です。したがって、すべてのスレッドは誰かが状態のロックを解除するのを待ちますが、他のスレッドはそれを行うことができず、コードはデッドロックに陥ります。