CUDA递增内核中的全局设备计数器

Question

我想找出一个简单的CUDA设计问题的解决方案。假设我有一个处理数据的内核。如果当前处理的数据符合指定条件，则outputArray中的相应元素将获得当前计数器值，并且计数器将递增。

它看起来像：

__global__ void setTags(INDATA* inputData, int* tags) 
{ 
    int blockId = blockIdx.x + blockIdx.y * gridDim.x; 
    int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x; 
    if(threadId < N) 
    { 
     INDATA current = inputData[threadId]; 
     if(/* current meets some criteria */) 
     { 
      tags[threadId] = /*current counter value */ 
      /* increment counter value */ 
     } 
    } 
} 

符合条件的元素数量多于所有元素显著少。重点在于处理停留在GPU上，并且我的案例使用唯一整数标记，范围从0到符合条件的案例数量 - 1.是否有更快的方法来执行此操作，而不是单线程增量执行？

Answer 1

这听起来像你想要的是atomicAdd函数增加一些全局计数器，可以由许多线程同时访问。你可能有这样的事情：

__device__ int counter; // initialise before running kernel 

__global__ void setTags(INDATA* inputData, int* tags) 
{ 
    int blockId = blockIdx.x + blockIdx.y * gridDim.x; 
    int threadId = blockId * (blockDim.x * blockDim.y) + 
        (threadIdx.y * blockDim.x) + threadIdx.x; 
    if(threadId < N) 
    { 
     INDATA current = inputData[threadId]; 
     if(/* current meets some criteria */) 
     { 
      int current_val = atomicAdd(&counter, 1); 
      tags[threadId] = current_val; 
     } 
    } 
} 

这里，atomicAdd将读取的counter值原子和单个原子操作由一个增加它。如果总增量操作的数量不是太大，这将是高性能的。但是，由于操作将序列化执行，因此如果您启动的大量总线程将访问并递增全局计数器，则可能会寻找另一种方法。