與CUDA共享內存互斥 - 增加項目

Question

我的問題的列表如下：我有，我發現使用GPU的興趣幾點的圖像。檢測是一個重量級的測試，但是隻有25個點中有1個平均通過測試。該算法的最後階段是建立一個點的列表。在CPU上，這將實現爲：

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

在GPU上，我有每個CUDA塊處理16x16像素。問題是我需要做一些特殊的事情，最終在全局內存中有一個統一的點列表。目前我正在嘗試在每個塊的共享內存中生成一個本地列表，這些列表最終將被寫入全局內存。我試圖避免發送任何回到CPU，因爲在此之後有更多的CUDA階段。

我期待，我可以用原子操作來實現對共享內存的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

我建議使用前綴和落實這部分增加的並行性。要做到這一點，你需要使用共享數組。基本上，前綴總和會將一個數組（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被設置爲零。

#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; 
    } 
}