synchronization - GPU 上的高效桶排序

Question

对于当前的 OpenCL GPGPU 项目，我需要根据具有 64 个可能值的某个键对数组中的元素进行排序。我需要最后一个数组让所有具有相同键的元素是连续的。有一个关联数组就足够了 new_index[old_index]作为这个任务的输出。

我把任务分成两部分。首先，我为每个可能的键(桶)计算具有该键(进入该桶)的元素数量。我扫描这个数组(生成一个前缀和)，它指示每个存储桶的元素的新索引范围，比如每个存储桶的“开始”索引。

然后，第二步必须为每个元素分配一个新索引。如果我要在 CPU 上实现它，算法将是这样的:

for all elements e:
    new_index[e] = bucket_start[bucket(e)]++

当然，这不适用于 GPU。每个项目都需要访问bucket_start读写模式下的数组本质上是所有工作项之间的同步，这是我们能做的最坏的事情。

一个想法是在工作组中进行一些计算。但我不确定这应该如何完成，因为我在 GPGPU 计算方面没有经验。

在全局内存中，我们用上面的前缀 sum 初始化了存储桶起始数组。对这个数组的访问是用原子整数“互斥”的。 (我是新手，所以可能在这里混了一些词。)

每个工作组都被隐式分配了输入元素数组的一部分。它使用一个包含新索引的本地桶数组，相对于我们还不知道的(全局)桶开始。在这些“本地缓冲区”之一已满后，工作组必须将本地缓冲区写入全局数组。为此，它锁定对全局存储区起始数组的访问，按当前本地存储区大小递增这些值，解锁，然后可以将结果写入全局 new_index。数组(通过添加相应的偏移量)。重复此过程，直到处理完所有分配的元素。

出现两个问题:

这是一个好方法吗？ 我知道读取和写入全局内存很可能是这里的瓶颈，特别是因为我试图获得对全局内存(至少只有一小部分)的同步访问。但也许有更好的方法来做到这一点，也许使用内核分解。请注意，我尽量避免在内核期间将数据从 GPU 读回 CPU(以避免 OpenCL 命令队列刷新，这也很糟糕，正如我所认为的那样)。

在上面的算法设计中，我如何实现锁定机制 ?像下面的代码会起作用吗？特别是，当硬件在 SIMD 组中“真正并行”执行工作项时，我预计会出现问题，例如 Nvidia“扭曲”。在我当前的代码中，工作组的所有项目都将尝试以 SIMD 方式获取锁。我应该仅限于第一个工作项目吗？并使用障碍使它们在本地保持同步？

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

__kernel void putInBuckets(__global uint *mutex,
                           __global uint *bucket_start,
                           __global uint *new_index)
{
    __local bucket_size[NUM_BUCKETS];
    __local bucket[NUM_BUCKETS][LOCAL_MAX_BUCKET_SIZE]; // local "new_index"

    while (...)
    {
        // process a couple of elements locally until a local bucket is full
        ...

        // "lock"
        while(atomic_xchg(mutex, 1)) {
        }

        // "critical section"
        __local uint l_bucket_start[NUM_BUCKETS];
        for (int b = 0; b < NUM_BUCKETS; ++b) {
            l_bucket_start[b] = bucket_start[b]; // where should we write?
            bucket_start[b] += bucket_size[b];   // update global offset
        }

        // "unlock"
        atomic_xchg(mutex, 0);

        // write to global memory by adding the offset
        for (...)
            new_index[...] = ... + l_bucket_start[b];
    }
}

Answer 1

首先，永远不要尝试在 GPU 上实现锁定算法。它会陷入僵局和停滞。
这是因为 GPU 是 SIMD 设备，线程不像在 CPU 上那样独立执行。 GPU 同步执行一组称为 WARP/WaveFront 的线程。因此，如果波前中的一个线程停止，它会停止波前中的所有其他线程。如果解锁线程处于停滞的波前，它将不会执行和解锁互斥锁。

原子操作没问题。

您应该考虑的是一种无锁方法。有关解释和示例 CUDA 代码，请参阅本文:
http://www.cse.iitk.ac.in/users/mainakc/pub/icpads2012.pdf/

它用一些示例 CUDA 代码描述了无锁哈希表、链表和跳过列表。

建议的方法是创建一个两级数据结构。

第一级是无锁跳过列表。每个跳过列表条目具有用于重复值的无锁链表的二级结构。以及条目数的原子计数。

插入方法

1) 生成64桶key
2) 在跳过列表中查找键
3) 如果未找到，则插入跳过列表
4) 向链表中插入数据
5) 增加这个桶的原子计数器

在插入前缀和跳过列表桶的所有计数器之后，你找到了
输出。