0

グローバル メモリ内のリンク リストを操作する 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 アプリケーションにとって理想的ではないことは認識しています。しかし、この場合、これを別の方法で実装する方法はまだわかりません。私のプロジェクトは、非常に急速に近づいている完全に固定された期限にあるため、これを行う必要があります。

4

1 に答える 1

1

セマフォが取得される前の古いデータを使用せずに、セマフォによって保護されたクリティカル セクション内でリスト操作が完全に実行されるようにする必要があります。

l->nextandmmd.freeCellsを volatile として宣言するか、アトミック関数 ( atomicExch()) で操作します。

または、適切なキャッシュ オペレータを使用してインライン アセンブリを使用することもできます。ロードに使用mov.cgするだけで、ローカルにキャッシュされた値が使用されないことを確認できます。また、セマフォが解放される前に書き込みがグローバル メモリに到達したことを確認するため__threadfence()の直前に使用することもできます。signal()を必ず使用してください。そうしないasm volatile(...)と、コンパイラはインライン asm 全体をクリティカル セクションから自由に移動できます。

于 2012-10-26T11:53:14.747 に答える