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原子的值并在单个原子操作中将其递增1.如果总增量操作的数量不是太大,则这将是高效的.但是,由于操作序列化执行,如果您启动的大量线程将访问并递增全局计数器,您可能会寻找另一种方法.