CUDA 원자 및 비 원자 메모리 액세스

Question

전역 메모리에서 연결된 목록을 조작하는 두 가지 CUDA 함수가 있습니다. pmalloc 함수는 목록 중 하나의 head 요소를 제거합니다. 먼저 목록을 선택한 다음 pmallocBucket을 호출하여 실제로 헤드 요소를 제거합니다. 선택한 목록이 비어 있으면 pmalloc이 다른 목록을 시도합니다. 반면에 pfree 함수는 새로운 헤드 요소를 목록에 삽입합니다.

상호 배타는 각 연결된 목록마다 하나씩 세마포를 통해 이루어집니다. 세마포어에 대한 구현은 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; 
} 

나는 매우 실망합니다. 나는 실제로 잘못되고있는 것이 무엇인지 모릅니다. 그리고 그것을 제거하기위한 지금까지의 모든 시도는 실패했습니다. 어떤 도움이라도 대단히 감사합니다.

P. S .: 그렇습니다. 원자 적 연산과 세마포어의 사용이 CUDA 어플리케이션에 이상적이지 않다는 것을 알고 있습니다. 그러나이 경우 어떻게 다른 방식으로 구현할 수 있을지 전혀 모르겠으며 프로젝트는 매우 빠르게 접근하는 마감일이 정해져 있으므로이 작업을 수행해야합니다.

Answer 1

세마포를 획득하기 전에 부실 데이터를 사용하지 않고 세마포어로 보호되는 중요 섹션 내에서 목록 조작이 완전히 수행되는지 확인해야합니다.

휘발성으로 l->next 및 mmd.freeCells을 선언하거나 원자 기능 (atomicExch())을 통해 조작하십시오.

또는 적절한 캐시 연산자로 인라인 어셈블리를 사용할 수 있습니다. 로드에 mov.cg을 사용하면 로컬 캐시 값이 사용되지 않도록하고 과 함께 signal() 바로 전에 세마포어가 해제되기 전에 쓰기가 전역 메모리에 도달했는지 확인해야합니다. asm volatile(...)을 사용하십시오. 그렇지 않으면 컴파일러가 전체 인라인 asm을 중요한 섹션 밖으로 자유롭게 이동할 수 있습니다.