c++ - 填充 CUDA 内核中的数组或列表，但不是在每个线程中

Question

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

类似的东西:

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

事实上，我想避免使用 c[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 卡，并且预期发出输出的线程数量很少，则第二个选项可能会更快。如果不是这种情况，流压缩可能是更好的选择。