在CUDA内核中填充数组或列表,但不在每个线程中填充

Question

基本上,我的内核中有一个if(),如果条件已经验证,我想在动态列表或数组中存储一个新值.问题是我不能使用threadIdx,因为它不会被填充到每个内核中.

就像是 :

__global__ void myKernel(customType *c)
{
    int i = threadIdx.x;
    //whatever
    if(condition)
        c->pop(newvalue)
}

实际上我想避免ac [i] = newvalue,因为最后我需要检查每个c [i]是否插入了一个值或者没有在主机代码中使用for循环并且正确填充另一个结构.我想到了推力但似乎对我的"简单"问题来说太过分了.

希望你能帮我找到解决方法.

Answer 1

如果我正确理解了你的问题,你有两个选择.

第一种方法是为每个线程预先分配一个输出位置,并且只有一些线程写入其输出.这将为您留下一个带有间隙的输出.您可以使用流压缩消除间隙,这是CUDA中已解决的问题 - 快速谷歌搜索将提供许多选项,并且Thrust和CUDPP都具有您可以使用的压缩功能.

第二种选择是使用全局内存计数器,并让每个线程在输出流中使用位置时以原子方式递增计数器,如下所示:

unsigned int opos; // set to zero before call

__global__ void myKernel(customType *c)
{
    //whatever
    if(condition) {
        unsigned int pos = atomicAdd(&opos, 1);
        c[pos] = newval;
    }
}

如果你有一个Kepler卡,并且预期发出输出的线程数很少,第二个选项可能会更快.如果不是这种情况,流压缩可能是更好的选择.