グローバル メモリ内のリンク リストを操作する 2 つの CUDA 関数があります。この関数pmalloc
は、リストの 1 つのヘッド要素を削除します。pmallocBucket
最初にリストを選択してから、 head 要素を実際に削除する呼び出しを行います。選択したリストが空の場合、pmalloc
他のリストを試します。一方pfree
、関数は新しい head 要素をリストに挿入します。
相互排除は、リンクされたリストごとに 1 つずつ、セマフォによって実現されます。セマフォの実装は、本CUDA By Exampleからのものです。他のテスト コードでは、セマフォは完全に機能します。
このコードの問題点は次のとおりです。複数のスレッドが同じリンク リストに同時にアクセスしようとする場合があります。これらのアクセスはセマフォによって正常に順次化されますが、スレッドが前のスレッドと同じヘッド要素をリストから削除することがあります。これは、すぐに連続して発生する場合もあれば、間に 1 つ以上の他のスレッドが存在する場合もあります。スレッドはfree
未割り当てのメモリ領域になり、プログラムがクラッシュします。
言及された機能は次のとおりです。mmd
別の関数から初期化されるグローバル メモリ内の構造体です。
extern __device__ void wait(int* s) {
while(atomicCAS(s, 0, 1) != 0);
}
extern __device__ void signal(int* s) {
atomicExch(s, 0);
}
__device__ void pfree(Expression* node) {
LinkedList* l = (LinkedList*) malloc(sizeof(LinkedList));
l->cell = node;
node->type = EMPTY;
node->funcidx = 0;
node->name = NULL;
node->len = 0;
node->value = 0;
node->numParams = 0;
free(node->params);
int targetBin = (blockIdx.x * mmd.bucketSize + threadIdx.x) / BINSIZE;
/*
* The for loop and subsequent if are necessary to make sure that only one
* thread in a warp is actively waiting for the lock on the semaphore.
* Leaving this out will result in massive headaches.
* See "CUDA by example", p. 273
*/
for(int i = 0; i < WARPSIZE; i++) {
if(((threadIdx.x + blockIdx.x * blockDim.x) % WARPSIZE) == i) {
wait(&mmd.bucketSemaphores[targetBin]);
l->next = mmd.freeCells[targetBin];
mmd.freeCells[targetBin] = l;
signal(&mmd.bucketSemaphores[targetBin]);
}
}
}
__device__ Expression* pmalloc() {
Expression* retval = NULL;
int i = 0;
int bucket = (blockIdx.x * mmd.bucketSize + threadIdx.x) / BINSIZE;
while(retval == NULL && i < mmd.numCellBins) {
retval = pmallocBucket((i + bucket) % mmd.numCellBins);
i++;
}
if(retval == NULL) {
printf("(%u, %u) Out of memory\n", blockIdx.x, threadIdx.x);
}
return retval;
}
__device__ Expression* pmallocBucket(int bucket) {
Expression* retval = NULL;
if(bucket < mmd.numCellBins) {
LinkedList* l = NULL;
for(int i = 0; i < WARPSIZE; i++) {
if(((threadIdx.x + blockIdx.x * blockDim.x) % WARPSIZE) == i) {
wait(&mmd.bucketSemaphores[bucket]);
l = mmd.freeCells[bucket];
if(l != NULL) {
retval = l->cell;
mmd.freeCells[bucket] = l->next;
}
signal(&mmd.bucketSemaphores[bucket]);
free(l);
}
}
}
return retval;
}
私はかなり途方に暮れています。実際に何が問題なのかはわかりません。これまでのところ、問題を解決するためのすべての試みは成功していません。どんな助けでも大歓迎です。
PS: はい、アトミック操作とセマフォの使用が CUDA アプリケーションにとって理想的ではないことは認識しています。しかし、この場合、これを別の方法で実装する方法はまだわかりません。私のプロジェクトは、非常に急速に近づいている完全に固定された期限にあるため、これを行う必要があります。