在 CUDA 中同步多个变量

Posted 2023-02-21

【中文标题】在 CUDA 中同步多个变量

【英文标题】：more than one variable to synchronize in CUDA

【发布时间】：2016-02-11 04:05:21

【问题描述】：

我的程序有很多 4 字节字符串，例如“aaaa”“bbbb”“cccc”...我需要收集通过 crc 检查的特定字符串。

因为字符串通过 crc 检查的可能性很小，所以我不想使用非常大的缓冲区来保存所有结果。我更喜欢一一连接的结果，就像输入一样。例如，如果输入为“aaaabbbbcccc”且“bbbb”未通过crc检查，则输出字符串应为“aaaacccc”且output_count应为2。

代码如下：

__device__
bool is_crc_correct(char* str, int len) 
    return true; // for simplicity, just return 'true';


// arguments:
// input: a sequence of 4-bytes-string, eg: aaaabbbbccccdddd....
__global__
void func(char* input, int* output, int* output_count) 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x;

    if(is_crc_correct(input + 4*index)) 
        // copy the string
        memcpy(output + (*output_count)*4,
               input + 4*index,
               4);
        // increase the counter
        (*output_count)++;
    

显然内存拷贝不是线程安全的，我知道atomicAdd函数可以用于++操作，但是如何让output和output_count线程安全呢？

【问题讨论】：

我相信您正在尝试重新发明 Stream 压缩，尤其是 Gather 操作，效率非常低。并行编程通常需要不同的思考。例如，您避免竞争，而不是尝试用原子和锁来解决它们（序列化有点违背并行化的目的）。你可能会使用thrust::copy_if。

【参考方案1】：

您正在寻找的是无锁线性分配器。这样做的常用方法是使用原子增加的累加器，用于索引缓冲区。例如，在您的情况下，以下应该有效：

__device__
char* allocate(char* buffer, int* elements) 
    // Here, the size of the allocated segment is always 4.
    // In a more general use case you would atomicAdd the requested size.
    return buffer + atomicInc(elements) * 4;

然后可以这样使用：

__global__
void func(char* input, int* output, int* output_count) 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x;

    if(is_crc_correct(input + 4*index)) 
        // Reserve the output buffer.
        char* dst = allocate(output, output_count);
        memcpy(dst, input + 4 * index, 4);
    

虽然这是完全线程安全的，但不保证保留输入顺序。例如，“ccccaaaa”将是一个有效的输出。

---

正如 Drop 在他们的评论中提到的那样，您正在尝试做的实际上是流压缩（并且 Thrust 可能已经提供了您需要的东西）。

我上面发布的代码可以通过首先聚合输出字符串通过warp来进一步优化，而不是直接分配到全局缓冲区中。这将减少全局原子争用并可能带来更好的性能。有关如何执行此操作的说明，我邀请您阅读以下文章：CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics。

【讨论】：

这很聪明。

【参考方案2】：

我可能会因为这个建议而下地狱，但是在内核内部动态分配内存怎么样？有关示例，请参阅此问题/答案：CUDA allocate memory in __device__ function

然后您将向每个内核传递一个共享内存数组，并且在内核运行之后，该数组的每个元素要么指向一块动态分配的内存，要么指向 NULL。因此，在您的线程块运行后，您将在单个线程上运行最终的清理内核来构建最终的字符串。

【讨论】：

谢谢，我什至不知道内核可以使用malloc/free。我记得我见过一些编译错误，比如“不能在设备代码中使用主机函数”，可能是旧的 CUDA 3.x。现在我尝试了它编译的 CUDA 6.5。