CUDA와 공유 메모리 뮤텍스 - 항목 목록에 추가

Question

내 문제는 다음과 같습니다. GPU를 사용하여 관심 지점을 감지 한 이미지가 있습니다. 탐지는 처리면에서 중량 테스트이지만 25 점 중 약 1 점이 평균 테스트에 합격합니다. 알고리즘의 마지막 단계는 포인트 목록을 작성하는 것입니다. I CUDA 각 블록의 처리를 16 × 16 픽셀들을 갖고 GPU에

forall pixels x,y 
{ 
    if(test_this_pixel(x,y)) 
     vector_of_coordinates.push_back(Vec2(x,y)); 
} 

다음에 CPU가 구현 될 것이다. 문제는 결국 글로벌 메모리에서 하나의 통합 포인트 목록을 갖기 위해 특별한 것을해야한다는 것입니다. 현재 전역 메모리에 쓰여지는 블록 당 공유 메모리의 로컬 지점 목록을 생성하려고합니다. 이후 CUDA 스테이지가 더 많아지기 때문에 CPU로 아무것도 보내지 않으려합니다.

나는 원자 연산을 사용하여 공유 메모리에 push_back 함수를 구현할 수있을 것으로 기대하고있었습니다. 그러나 나는이 일을 얻을 수 없습니다. 두 가지 문제가 있습니다. 첫 번째 성가신 문제는 다음과 같은 컴파일러 크래시가 지속적으로 발생한다는 것입니다. "nvcc 오류 : 'ptxas'상태가 0xC0000005 (ACCESS_VIOLATION) 상태로 사망했습니다. 내가 뭔가를 컴파일 할 수 있는지 여부가 맞았거나 놓쳤습니다. 아무도이 원인을 알고 있습니까? 다음 커널은 오류 재현

:

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pCounts) 
{ 
    __shared__ unsigned int test; 
    atomicInc(&test, 1000); 
} 

둘째, 공유 메모리에 뮤텍스 잠금을 포함 내 코드는 GPU를 중단하고 이해 해달라고 왜 : 예에서

__device__ void lock(unsigned int *pmutex) 
{ 
    while(atomicCAS(pmutex, 0, 1) != 0); 
} 

__device__ void unlock(unsigned int *pmutex) 
{ 
    atomicExch(pmutex, 0); 
} 

__global__ void gpu_kernel_non_max_suppress(int w, int h, RtmPoint *pPoints, int *pCounts) 
{ 
    __shared__ RtmPoint localPoints[64]; 
    __shared__ int localCount; 
    __shared__ unsigned int mutex; 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 

    int threadid = threadIdx.y * blockDim.x + threadIdx.x; 
    int blockid = blockIdx.y * gridDim.x + blockIdx.x; 

    if(threadid==0) 
    { 
     localCount = 0; 
     mutex = 0; 
    } 

    __syncthreads(); 

    if(x<w && y<h) 
    { 
     if(some_test_on_pixel(x,y)) 
     { 
      RtmPoint point; 
      point.x = x; 
      point.y = y; 

      // this is a local push_back operation 
      lock(&mutex); 
      if(localCount<64) // we should never get >64 points per block 
       localPoints[localCount++] = point; 
      unlock(&mutex); 
     } 
    } 

    __syncthreads(); 

    if(threadid==0) 
     pCounts[blockid] = localCount; 
    if(threadid<localCount) 
     pPoints[blockid * 64 + threadid] = localPoints[threadid]; 
} 

this site의 코드를 작성한 저자는 공유 메모리에서 원자 연산을 성공적으로 사용하므로 내 사건이 작동하지 않는 이유에 대해 혼란 스럽습니다. 잠금 및 주석 해제 행을 주석 처리하면 코드가 정상적으로 실행되지만 분명히 목록에 잘못 추가됩니다.

어쨌든 원자 연산이나 뮤텍스 잠금을 사용하여 성능 문제가 걱정되므로이 문제가 발생하는 이유와 목표를 달성하는 데 더 나은 솔루션이있는 경우에 대한 조언을 주시면 감사하겠습니다.

Answer 1

병렬 처리를 높이기 위해 prefix-sum을 사용하여 해당 부분을 구현하는 것이 좋습니다. 그렇게하려면 공유 배열을 사용해야합니다. 기본적으로 prefix-sum은 배열 (1,1,0,1)을 (0,1,2,2,3)으로 바꿀 것입니다. 즉, 인 - 스턴스 실행시의 배타적 합계를 계산하여 스레드 당 색인을 작성하십시오.

__shared__ uint8_t vector[NUMTHREADS]; 

.... 

bool emit = (x<w && y<h); 
    emit = emit && some_test_on_pixel(x,y); 
__syncthreads(); 
scan(emit, vector); 
if (emit) { 
    pPoints[blockid * 64 + vector[TID]] = point; 
} 

접두사 합 예 :

template <typename T> 
__device__ uint32 scan(T mark, T *output) { 
#define GET_OUT (pout?output:values) 
#define GET_INP (pin?output:values) 
    __shared__ T values[numWorkers]; 
    int pout=0, pin=1; 
    int tid = threadIdx.x; 

    values[tid] = mark; 

    syncthreads(); 

    for(int offset=1; offset < numWorkers; offset *= 2) { 
    pout = 1 - pout; pin = 1 - pout; 
    syncthreads(); 
    if (tid >= offset) { 
     GET_OUT[tid] = (GET_INP[tid-offset]) +(GET_INP[tid]); 
    } 
    else { 
     GET_OUT[tid] = GET_INP[tid]; 
    } 
    syncthreads(); 
    } 

    if(!pout) 
    output[tid] =values[tid]; 

    __syncthreads(); 

    return output[numWorkers-1]; 

#undef GET_OUT 
#undef GET_INP 
} 

Answer 2

여기 권고 사항을 바탕으로, 나는 내가 마지막에 사용 된 코드를 포함한다. 그것은 16x16 픽셀 블록을 사용합니다. 현재 데이터를 분할하지 않고 하나의 전역 배열에 쓰고 있습니다. 전역 atomicAdd 함수를 사용하여 각 결과 집합에 대한 기본 주소를 계산했습니다. 이 방법은 블록 당 한 번만 호출되기 때문에 느린 속도를 너무 많이 찾지는 못했지만이 작업을 통해 더 많은 편의를 얻었습니다. 나는 또한 prefix_sum의 입력과 출력을위한 공유 버퍼를 피합니다. GlobalCount는 커널 호출 전에 0으로 설정됩니다.

#define BLOCK_THREADS 256 

__device__ int prefixsum(int threadid, int data) 
{ 
    __shared__ int temp[BLOCK_THREADS*2]; 

    int pout = 0; 
    int pin = 1; 

    if(threadid==BLOCK_THREADS-1) 
     temp[0] = 0; 
    else 
     temp[threadid+1] = data; 

    __syncthreads(); 

    for(int offset = 1; offset<BLOCK_THREADS; offset<<=1) 
    { 
     pout = 1 - pout; 
     pin = 1 - pin; 

     if(threadid >= offset) 
      temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid] + temp[pin * BLOCK_THREADS + threadid - offset]; 
     else 
      temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid]; 

     __syncthreads(); 
    } 

    return temp[pout * BLOCK_THREADS + threadid]; 
} 

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pGlobalCount) 
{ 
    __shared__ int write_base; 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 

    int threadid = threadIdx.y * blockDim.x + threadIdx.x; 
    int valid = 0; 

    if(x<w && y<h) 
    { 
     if(test_pixel(x,y)) 
     { 
      valid = 1; 
     } 
    } 

    int index = prefixsum(threadid, valid); 

    if(threadid==BLOCK_THREADS-1) 
    { 
     int total = index + valid; 
     if(total>64) 
      total = 64; // global output buffer is limited to 64 points per block 
     write_base = atomicAdd(pGlobalCount, total); // get a location to write them out 
    } 

    __syncthreads(); // ensure write_base is valid for all threads 

    if(valid) 
    { 
     RtmPoint point; 
     point.x = x; 
     point.y = y; 
     if(index<64) 
      pPoints[write_base + index] = point; 
    } 
}