配列からのカウント値CUDA

Question

私はfloat値の配列、すなわちlifeを持っていますが、その中でCUDAの値が0より大きいエントリの数を数えたいと思います。 CPUオン

、コードは次のようになります。

int numParticles = 0; 
for(int i = 0; i < MAX_PARTICLES; i++){ 
    if(life[i]>0){ 
     numParticles++; 
    } 
} 

を今CUDAに、私はこのようなものを試してみた：

__global__ void update(float* life, int* numParticles){ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    if (life[idx]>0){ 
     (*numParticles)++; 
    } 
} 
//life is a filled device pointer 
int launchCount(float* life) 
{ 
    int numParticles = 0; 
    int* numParticles_d = 0; 
    cudaMalloc((void**)&numParticles_d, sizeof(int)); 
    update<<<MAX_PARTICLES/THREADS_PER_BLOCK,THREADS_PER_BLOCK>>>(life, numParticles_d); 
    cudaMemcpy(&numParticles, numParticles_d, sizeof(int), cudaMemcpyDeviceToHost); 
    std::cout << "numParticles: " << numParticles << std::endl; 
} 

をしかし、何らかの理由でCUDAの試みは常にリターンnumParticlesの場合は0です。どうして？

Answer 1

実際にコードはMAX_PARTICLESスレッドを起動しており、複数のスレッドブロックが同時に実行しています(*numParticles)++;。競合状態です。だから、結果0を持っている、またはあなたがすべてのiためlife[i]>0 ? 1 : 0を総括するためにあなたの試みとして0

よりも、時には少し大きい運、している場合、あなたはあなたのカーネルを実装するためにCUDA parallel reductionに従うか、あなたを簡素化するためにThrust reductionを使用することができます生活。

Answer 2

この：

if (life[idx]>0){ 
    (*numParticles)++; 
} 

は、読み取り後の書き込みの危険性です。複数のスレッドが同時にnumParticlesからの読み取りと書き込みを試みます。 CUDA実行モデルは、同時トランザクションの順序について何も保証しません。

次の例のように、アトミックメモリトランザクションを使用して、この作品を作ることができる：

if (life[idx]>0){ 
    atomicAdd(numParticles, 1); 
} 

これは、メモリトランザクションをシリアル化し、計算が正しいようになります。また、パフォーマンスに大きな悪影響を及ぼします。

各ブロックに還元型の計算を使用してローカルの合計を計算させてから、ブロックローカルの合計をアトミックに、またはホスト上で、または2番目のカーネルで合計したい場合があります。