algorithm - 在 GPU 上并行写入位集(数组的数组)

Question

这是一个与我试图加速的问题相关的 GPU 算法问题:

假设我在概念上有一个如下所示的数据字段，其中 512 是 block 中的线程数:

bool is_a_foo[131072][512];

此结构中的 bool 表示其他地方的数据(恰好具有相似的维度......但这是不相关的)是否是 foo。为简单起见，假设我只是在一个 GPU block 上运行，每个线程都通过(通过 __syncwarp() 锁定步骤...但请不要让它太分散注意力，实际上，我正在做一些更有意义的事情)位置 0->131071。换句话说，每个线程的代码如下所示:

// assume is_a_foo is initialized earlier to 0's by some sort of memset call
// assume that the values for is_a_foo can go from false->true but never from true->false
for (int i = 0; i < 131072; ++i) {
    if (something_kind_of_expensive_but_not_the_bottleneck()) {
        is_a_foo[ i ][thread] = true;
    }
}

每个 bool 都表示为 8 位，没有数据丢失。但是，假设我想减少内存/缓存占用和带宽消耗。我们可以将上述数据结构表示为:

unsigned int is_a_foo[131072][512 / (sizeof(unsigned int) * 8)];

我们可以执行位算术将感兴趣的特定位设置为 1。

问题在于，如果不进行任何特殊处理，对is_a_foo 的写入将相互粉碎，并不是每个应该设置为 1 的位都必须设置为 1。

如果我们愿意做一些特别的事情，我们可以使用atomicCAS 来确保没有丢失任何写入。不幸的是，这似乎有点贵。事实上，在我的应用程序中，内核启动大约需要 30 毫秒，内核执行时间增加了约 33%。目前还不清楚额外的时间是由于原子操作还是额外的指令，但我怀疑这是原子操作。

如果我能够对 unsigned char 而不是 unsigned int 进行操作，可以减轻损害的一件事。不幸的是，CUDA 没有提供这样的接口(interface)。而且，当我对 unsigned short 进行操作时，我收到一个编译器错误，提示该函数不适用于 unsigned short(可根据要求提供详细信息)。

所有这些都是要问，是否有任何算法/数据结构非常适合 GPU 上的此类操作？

Answer 1

我不知道任何支持 CUDA 的 GPU 的 warp 大小为 512，所以我假设你打算写 block 大小和 __syncthreads() 而不是 warp 大小和 __syncwarp()(目前每个 CUDA 架构的 warp 大小都是 32)。我还可能会提醒您注意存在一个 atomicOr() 的事实。功能。

为了最小化原子数(或一般的全局内存流量)，典型的方法是执行 parallel reduction在你的 block 中(使用共享内存)为整个 block 建立结果，然后只在最后使用一堆线程将结果移出到全局内存。总的来说，我强烈建议您查看CUB。对于提供各种并行编程原语(例如归约)的 CUDA 实现的库。但是，在您的特定情况下，同一 warp 中的线程可以使用 __ballot() 简单地执行有问题的减少。 warp vote 函数(映射到单个指令)。由于数字在您的情况下有效，因此结果恰好是每个扭曲(32 个线程)一个 32 位位掩码，您可以只做一个 __ballot() 然后有一个(例如，第一个) 每个 warp 的线程写入结果。如果我正确理解你的问题，那么你甚至不需要原子，因为结果似乎是每个 block 每个扭曲一个位掩码，这意味着只要你只有一个线程访问全局内存，就不会并发访问同一位置扭曲……