0

我有两个 CUDA 函数来操作全局内存中的链表。该函数pmalloc删除列表之一的头元素。它首先选择一个列表，然后调用pmallocBucket实际删除头元素的列表。如果所选列表为空，pmalloc将尝试其他列表。pfree另一方面，该函数会将一个新的 head 元素插入到列表中。

互斥是通过信号量实现的，每个链表一个信号量。信号量的实现来自《CUDA By Example 》一书。在其他一些测试代码中，信号量可以完美运行。

我对代码的问题如下：有时，多个线程会尝试同时访问同一个链表。这些访问被信号量成功地顺序化，但有时，一个线程会从列表中删除与前一个线程相同的头元素。这可能会立即连续发生，或者中间可能有一个或多个其他线程。然后线程将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->next将and声明mmd.freeCells为 volatile，或通过原子函数 ( atomicExch()) 操作它们。

或者，您可以将内联汇编与合适的缓存运算符一起使用。使用mov.cg负载应该足以确保不使用本地缓存的值，以及在释放信号量之前确保写入已到达全局内存的__threadfence()右侧。signal()确保使用asm volatile(...)，否则编译器可以自由地将整个内联汇编移出关键部分。

于 2012-10-26T11:53:14.747 回答