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

Fill an array or a list in CUDA kernel but not in every thread

本文关键字：列表 线程 一个 内核 CUDA 填充 数组      更新时间：2023-10-16

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

类似于：

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

事实上，我希望避免使用c[I]=newvalue，因为在最后，我需要检查每个c[I]是否在宿主代码中插入了一个for循环，并正确填充另一个结构。我考虑过推力，但对于我的"简单"问题来说，这似乎有些过头了。

希望你能帮我找到解决办法。

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

第一种方法是为每个线程预先分配一个输出位置，并且只让一些线程写入其输出。这给你留下了一个有间隙的输出。你可以使用流压缩来消除间隙，这是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;
    }
}

如果你有一个开普勒卡，并且预期发出输出的线程数量很小，那么第二种选择可能会更快。如果不是这样的话，流压缩可能是更好的选择。

如果我理解正确，您描述的是流压缩。某些线程（并非所有线程）将创建一个值，并且您希望将这些值存储在一个没有任何间隙的数组中。

实现这一点的一种方法是使用Thrust中提供的流压缩算法（查看此示例）。请注意，这确实需要您分两次执行操作。

如果您是在一个线程块内（而不是整个网格）执行此操作，那么您也可以查看CUB。每个线程将计算一个标志，指示是否要存储值，对标志进行前缀求和，以确定列表中每个线程的偏移量，然后进行存储。

相关文章：