CUDAとの共有メモリミューテックス - アイテムリストへの追加

Question

私の問題は次のとおりです。GPUを使用していくつかの興味のあるポイントを検出したイメージがあります。検出は処理の点では重量テストですが、平均でテストに合格するのは約25ポイントにすぎません。アルゴリズムの最終段階は、ポイントのリストを構築することです。 CPU上では、これは次のように実装されます：

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

私は各CUDAブロックで16x16ピクセルを処理しています。問題は、最終的にグローバルメモリ内のポイントの単一の統合リストを持つためには特別なことをする必要があることです。現時点では、最終的にグローバルメモリに書き込まれるブロックごとの共有メモリ内のポイントのローカルリストを生成しようとしています。私はCPUに何かを返すことを避けようとしています。そのあとでCUDAのステージが増えています。

私は、共有メモリ上でpush_back関数を実装するためにアトミック操作を使用できることを期待していました。しかし、私はこれを得ることができません。 2つの問題があります。最初の厄介な問題は、次のようなコンパイラクラッシュに常に遭遇していることです。「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は配列（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ピクセルブロックを使用します。私は現在、データを分割せずに1つのグローバル配列に書き出すことに注意してください。グローバルなatomicAdd関数を使用して、各結果セットのベースアドレスを計算しました。これはブロックごとに1回しか呼び出されないので、私はこれを行うことでより多くの利便性を得ていましたが、スローダウンはあまり見つけられませんでした。私はまた、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; 
    } 
}