我把它发布到了Github上 新项目一个简单的 GPU 哈希表.

它是一个简单的 GPU 哈希表，每秒能够处理数亿次插入。 在我的 NVIDIA GTX 1060 笔记本电脑上，代码在大约 64 毫秒内插入 210 万个随机生成的键值对，并在大约 32 毫秒内删除 64 万对。

也就是说，笔记本电脑的速度约为每秒 300 亿次插入和 500 亿次删除/秒。

该表是用 CUDA 编写的，尽管相同的技术可以应用于 HLSL 或 GLSL。 为了确保显卡的高性能，该实现有几个限制：

• 仅处理32位密钥和相同的值。
• 哈希表具有固定的大小。
• 并且这个大小必须等于XNUMX的次方。

对于键和值，您需要保留一个简单的分隔符标记（在上面的代码中，这是 0xffffffff）。

无锁哈希表

哈希表使用开放寻址 线性探测，也就是说，它只是一个存储在内存中的键值对数组，具有优越的缓存性能。 对于链接则不然，链接涉及在链表中搜索指针。 哈希表是一个存储元素的简单数组 KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

表的大小是 2 的幂，而不是素数，因为一条快速指令足以应用 powXNUMX/AND 掩码，但模运算符要慢得多。 这在线性探测的情况下很重要，因为在线性表查找中，槽索引必须包含在每个槽中。 结果，操作成本在每个时隙中以模数相加。

该表仅存储每个元素的键和值，而不存储键的哈希值。 由于该表仅存储 32 位密钥，因此哈希计算速度非常快。 上面的代码使用 Murmur3 哈希，它只执行一些移位、异或和乘法。

哈希表使用与内存顺序无关的锁定保护技术。 即使某些写入操作打乱了其他此类操作的顺序，哈希表仍将保持正确的状态。 我们下面会谈到这一点。 该技术非常适合同时运行数千个线程的显卡。

哈希表中的键和值被初始化为空。

代码也可以修改为处理 64 位键和值。 键需要原子读取、写入以及比较和交换操作。 并且值需要原子读写操作。 幸运的是，在 CUDA 中，只要自然对齐，32 位和 64 位值的读写操作都是原子的（见下文）。 这里），现代显卡支持 64 位原子比较和交换操作。 当然，当迁移到 64 位时，性能会略有下降。

哈希表状态

哈希表中的每个键值对可以具有以下四种状态之一：

• 键和值均为空。 在此状态下，哈希表被初始化。
• key已经被写下来了，但是value还没有被写入。 如果另一个线程当前正在读取数据，则它返回空。 这是正常的，如果另一个执行线程早一点工作，也会发生同样的事情，我们正在讨论并发数据结构。
• 键和值都会被记录。
• 该值可供其他执行线程使用，但键尚不可用。 发生这种情况是因为 CUDA 编程模型具有松散排序的内存模型。 这是正常的；在任何事件期间，键仍然为空，即使值不再是空的。

一个重要的细微差别是，一旦密钥被写入插槽，它就不再移动 - 即使密钥被删除，我们将在下面讨论这一点。

哈希表代码甚至适用于松散排序的内存模型，其中内存读取和写入的顺序未知。 当我们查看哈希表中的插入、查找和删除时，请记住每个键值对都处于上述四种状态之一。

插入到哈希表中

将键值对插入哈希表的 CUDA 函数如下所示：

void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
        if (prev == kEmpty || prev == key)
        {
            hashtable[slot].value = value;
            break;
        }
        slot = (slot + 1) & (kHashTableCapacity-1);
    }
}

要插入键，代码会从插入键的哈希值开始迭代哈希表数组。 数组中的每个槽都执行原子比较和交换操作，将该槽中的键与空进行比较。 如果检测到不匹配，则使用插入的密钥更新槽中的密钥，然后返回原始槽密钥。 如果这个原始键为空或与插入的键匹配，则代码找到合适的插入槽并将插入的值插入到该槽中。

如果在一个内核调用中 gpu_hashtable_insert() 有多个元素具有相同的键，那么它们的任意值都可以写入键槽。 这被认为是正常的：调用期间的键值写入之一将成功，但由于所有这一切都是在多个执行线程内并行发生的，因此我们无法预测哪个内存写入将是最后一个。

哈希表查找

搜索键的代码：

uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
        uint32_t slot = hash(key);

        while (true)
        {
            if (hashtable[slot].key == key)
            {
                return hashtable[slot].value;
            }
            if (hashtable[slot].key == kEmpty)
            {
                return kEmpty;
            }
            slot = (slot + 1) & (kHashTableCapacity - 1);
        }
}

为了查找存储在表中的键的值，我们从要查找的键的哈希开始遍历数组。 在每个槽中，我们检查该键是否是我们要查找的键，如果是，则返回其值。 我们还检查键是否为空，如果是，则中止搜索。

如果我们找不到该键，代码将返回一个空值。

所有这些搜索操作都可以通过插入和删除同时执行。 表中的每一对都将具有上述流程的四种状态之一。

在哈希表中删除

删除键的代码：

void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        if (hashtable[slot].key == key)
        {
            hashtable[slot].value = kEmpty;
            return;
        }
        if (hashtable[slot].key == kEmpty)
        {
            return;
        }
        slot = (slot + 1) & (kHashTableCapacity - 1);
    }
}

删除键是以一种不寻常的方式完成的：我们将键保留在表中并将其值（而不是键本身）标记为空。 这段代码非常类似于 lookup()，除了当在键上找到匹配项时，它会将其值设为空。

如上所述，一旦密钥被写入插槽，它就不再被移动。 即使从表中删除元素，键仍保留在原位，其值只是变为空。 这意味着我们不需要对槽值使用原子写入操作，因为当前值是否为空并不重要 - 它仍然会变成空。

调整哈希表的大小

您可以通过创建更大的表并向其中插入旧表中的非空元素来更改哈希表的大小。 我没有实现此功能，因为我想让示例代码保持简单。 此外，在 CUDA 程序中，内存分配通常在主机代码中完成，而不是在 CUDA 内核中完成。

文章 无锁免等待哈希表 描述了如何修改这种受锁保护的数据结构。

竞争力

在上面的函数代码片段中 gpu_hashtable_insert(), _lookup() и _delete() 一次处理一个键值对。 并且更低 gpu_hashtable_insert(), _lookup() и _delete() 并行处理一组对，每对在一个单独的 GPU 执行线程中：

// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);

// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
    unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
    if (threadid < numkvs)
    {
        gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
    }
}

抗锁哈希表支持并发插入、查找和删除。 由于键值对始终处于四种状态之一，并且键不会移动，因此即使同时使用不同类型的操作，该表也能保证正确性。

然而，如果我们并行处理一批插入和删除，并且输入的对数组包含重复的键，那么我们将无法预测哪些对将“获胜”——将最后写入哈希表。 假设我们使用输入数组对来调用插入代码 A/0 B/1 A/2 C/3 A/4。 当代码完成后，配对 B/1 и C/3 保证出现在表中，但同时任何对都会出现在表中 A/0, A/2 или A/4。 这可能是问题，也可能不是问题——这完全取决于应用程序。 您可能事先知道输入数组中没有重复的键，或者您可能不关心最后写入的是哪个值。

如果这对您来说是个问题，那么您需要将重复对分成不同的 CUDA 系统调用。 在 CUDA 中，任何调用内核的操作总是在下一次内核调用之前完成（至少在一个线程内。在不同的线程中，内核是并行执行的）。 在上面的示例中，如果您使用以下命令调用一个内核 A/0 B/1 A/2 C/3，另一个与 A/4，那么关键 A 将得到该值 4.

现在我们来讨论一下函数是否应该 lookup() и delete() 使用普通或易失​​性指针指向哈希表中的对数组。 CUDA 文档 指出：

编译器可能会选择优化对全局或共享内存的读写...可以使用关键字禁用这些优化 volatile: ...对此变量的任何引用都会编译为实际内存读取或写入指令。

正确性考虑不需要应用 volatile。 如果执行线程使用先前读取操作的缓存值，那么它将使用稍微过时的信息。 但是，这仍然是来自内核调用某个时刻的哈希表的正确状态的信息。 如果需要使用最新的信息，可以使用索引 volatile，但随后性能会略有下降：根据我的测试，删除32万个元素时，速度从500亿删除/秒下降到450亿删除/秒。

Производительность

在插入64万个元素和删除32万个元素的测试中，之间的竞争 std::unordered_map 并且 GPU 几乎没有哈希表：

std::unordered_map 插入和删除元素然后释放它们花费了 70 毫秒 unordered_map （摆脱数百万个元素需要很多时间，因为里面 unordered_map 进行了多次内存分配）。 老实说， std:unordered_map 完全不同的限制。 它是单CPU线程执行，支持任意大小的key-value，在高利用率下表现良好，多次删除后表现稳定。

GPU 和程序间通信的哈希表的持续时间为 984 毫秒。 这包括将表放入内存并删除它（一次分配 1 GB 内存，这在 CUDA 中需要一些时间）、插入和删除元素以及迭代它们所花费的时间。 所有进出视频卡内存的副本也会被考虑在内。

哈希表本身需要 271 毫秒才能完成。 这包括视频卡插入和删除元素所花费的时间，并且不考虑复制到内存和迭代结果表所花费的时间。 如果 GPU 表存在很长一段时间，或者哈希表完全包含在显卡的内存中（例如，创建一个将由其他 GPU 代码而不是中央处理器使用的哈希表），则测试结果是相关的。

由于高吞吐量和主动并行化，视频卡的哈希表表现出高性能。

限制

哈希表架构有几个问题需要注意：

• 线性探测受到集群的阻碍，这会导致表中的键放置得不太完美。
• 使用该功能不会删除按键 delete 随着时间的推移，它们会把桌子弄得乱七八糟。

因此，哈希表的性能会逐渐下降，特别是当它存在很长时间并且有大量插入和删除时。 缓解这些缺点的一种方法是重新哈希到利用率相当低的新表中，并在重新哈希期间过滤掉已删除的键。

为了说明所描述的问题，我将使用上面的代码创建一个包含 128 亿个元素的表，并循环遍历 4 万个元素，直到填满 124 亿个槽（利用率约为 0,96）。 这是结果表，每一行都是 CUDA 内核调用，将 4 万个新元素插入到一个哈希表中：

使用率
插入持续时间 4 个元素

0,00
11,608448 毫秒（361,314798 万键/秒）

0,03
11,751424 毫秒（356,918799 万键/秒）

0,06
11,942592 毫秒（351,205515 万键/秒）

0,09
12,081120 毫秒（347,178429 万键/秒）

0,12
12,242560 毫秒（342,600233 万键/秒）

0,16
12,396448 毫秒（338,347235 万键/秒）

0,19
12,533024 毫秒（334,660176 万键/秒）

0,22
12,703328 毫秒（330,173626 万键/秒）

0,25
12,884512 毫秒（325,530693 万键/秒）

0,28
13,033472 毫秒（321,810182 万键/秒）

0,31
13,239296 毫秒（316,807174 万键/秒）

0,34
13,392448 毫秒（313,184256 万键/秒）

0,37
13,624000 毫秒（307,861434 万键/秒）

0,41
13,875520 毫秒（302,280855 万键/秒）

0,44
14,126528 毫秒（296,909756 万键/秒）

0,47
14,399328 毫秒（291,284699 万键/秒）

0,50
14,690304 毫秒（285,515123 万键/秒）

0,53
15,039136 毫秒（278,892623 万键/秒）

0,56
15,478656 毫秒（270,973402 万键/秒）

0,59
15,985664 毫秒（262,379092 万键/秒）

0,62
16,668673 毫秒（251,627968 万键/秒）

0,66
17,587200 毫秒（238,486174 万键/秒）

0,69
18,690048 毫秒（224,413765 万键/秒）

0,72
20,278816 毫秒（206,831789 万键/秒）

0,75
22,545408 毫秒（186,038058 万键/秒）

0,78
26,053312 毫秒（160,989275 万键/秒）

0,81
31,895008 毫秒（131,503463 万键/秒）

0,84
42,103294 毫秒（99,619378 万键/秒）

0,87
61,849056 毫秒（67,815164 万键/秒）

0,90
105,695999 毫秒（39,682713 万键/秒）

0,94
240,204636 毫秒（17,461378 万键/秒）

随着利用率的增加，性能会下降。 在大多数情况下这是不可取的。 如果应用程序将元素插入表中然后丢弃它们（例如，在计算书中的单词数时），那么这不是问题。 但是，如果应用程序使用长期存在的哈希表（例如，在图形编辑器中存储用户频繁插入和删除信息的图像的非空部分），则这种行为可能会出现问题。

并测量了 64 万次插入后的哈希表探测深度（利用率0,5）。 平均深度为 0,4774，因此大多数键要么位于最佳位置，要么位于距最佳位置一处的位置。 最大探测深度为60。

然后，我在具有 124 亿次插入的表上测量了探测深度（利用率为 0,97）。 平均深度已经是 10,1757，最大深度是 - 6474 （！！）。 高利用率下线性传感性能显着下降。

最好保持该哈希表的利用率较低。 但随后我们以牺牲内存消耗为代价来提高性能。 幸运的是，对于 32 位键和值，这是合理的。 如果在上面的例子中，在一个有 128 亿个元素的表中，我们保持利用率为 0,25，那么我们可以在其中放置不超过 32 万个元素，剩余的 96 万个槽位将丢失 - 每对 8 个字节, 768 MB 丢失内存。

请注意，我们谈论的是显卡内存的损失，它是比系统内存更有价值的资源。 尽管大多数支持 CUDA 的现代桌面显卡至少有 4 GB 内存（在撰写本文时，NVIDIA 2080 Ti 有 11 GB），但失去如此大的内存仍然不是最明智的决定。

稍后我将写更多关于为没有探测深度问题的显卡创建哈希表，以及重用已删除插槽的方法。

测深深度测量

为了确定键的探测深度，我们可以从其实际表索引中提取键的哈希（其理想表索引）：

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

由于二进制补码的魔力以及哈希表的容量是 1 的 3 次方这一事实，即使将键索引移动到表的开头，这种方法也将起作用。 让我们取一个散列为 4 的键，但插入到槽 XNUMX 中。然后对于容量为 XNUMX 的表，我们得到 (3 — 1) & 3，相当于 2。

结论

如果您有疑问或意见，请给我发电子邮件： Twitter 或打开一个新主题 储存库.

这段代码是在优秀文章的启发下编写的：

• 世界上最简单的无锁哈希表
• 无锁免等待哈希表

以后我会继续写显卡的哈希表实现，并分析其性能。 我的计划包括在 GPU 友好的数据结构中使用原子操作进行链接、Robin Hood 哈希和布谷鸟哈希。

来源： habr.com