我在Github上发布了一个名为A Simple GPU Hash Table的新项目。这是用于GPU的简单哈希表，每秒可以处理数亿个插入。在装有NVIDIA GTX 1060的笔记本电脑上，该代码在大约210毫秒内插入了6400万随机生成的键值对，并在大约64毫秒内删除了3200万对键值对。也就是说，笔记本电脑的速度约为每秒3亿次插入和每秒5亿次清除。该表是用CUDA编写的，尽管可以将相同的技术应用于HLSL或GLSL。该实现有几个限制，以确保视频卡具有高性能：- 仅处理32位密钥和相同的值。
- 哈希表具有固定大小。
- 并且此大小应等于2度。
对于键和值，您需要保留一个简单的定界标记（在上面的代码中为0xffffffff）。没有锁的哈希表
哈希表使用带有线性感应的开放式寻址，也就是说，它只是存储在内存中的一组键值对数组，并具有出色的缓存性能。链接不是这种情况，这意味着在链接列表中查找指针。哈希表是存储元素的简单数组KeyValue：struct KeyValue
{
uint32_t key;
uint32_t value;
};
该表的大小等于2的幂，而不是质数，因为使用pow2 / AND-mask，一条快速指令就足够了，而模块运算符要慢得多。这在线性感应的情况下很重要，因为在表的线性搜索中，插槽索引必须包装在每个插槽中。结果，在每个时隙中模运算增加了操作成本。该表仅存储每个项目的键和值，而不存储键哈希。由于该表仅存储32位密钥，因此哈希计算非常快。上面的代码使用Murmur3哈希，该哈希仅执行一些移位，XOR和乘法。哈希表使用的阻塞保护技术不依赖于内存分配顺序。即使某些写入操作违反了其他此类操作的顺序，哈希表仍将保持正确的状态。我们将在下面讨论。该技术非常适合视频卡，其中有数千个线程在竞争。哈希表中的键和值被初始化为空。可以修改该代码，以便它可以处理64位键和值。密钥需要原子的读，写和交换操作（比较和交换）。值需要原子读取和写入操作。幸运的是，在CUDA中，对32位和64位值的读写操作是原子的，只要它们自然对齐即可（请参见此处）），现代的视频卡支持64位原子操作与交换的比较。当然，切换到64位时，性能会略有下降。哈希表状态
哈希表中的每个键值对可以具有以下四种状态之一：- 键和含义为空。在这种状态下，哈希表被初始化。
- 密钥已被记录，但尚未值。如果此时另一个执行线程正在读取数据，则它将返回一个空值。这很正常，如果另一个执行线程早一点起作用，也会发生同样的事情，而我们正在谈论的是竞争性数据结构。
- 键和值都被记录。
- 该值可用于其他执行线程，但该键尚未提供。之所以会发生这种情况，是因为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执行线程中：
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);
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，但是性能会略有下降：根据我的测试，当删除3200万个项目时，速度从5亿个删除/秒降低到4.5亿个删除/秒。性能
在插入6400万个元素并删除3200万个元素的测试中，std::unordered_mapGPU的哈希表与它们之间几乎没有竞争：std::unordered_map在插入和删除元素以及随后的释放时花费了70691毫秒unordered_map（从数百万个元素中释放会花费很多时间，因为内部unordered_map执行了大量内存分配）。老实说，存在std:unordered_map完全不同的限制。这是一个执行的CPU线程，它支持任何大小的键值，在高利用率下工作良好，并且在多次删除后显示稳定的性能。GPU和程序间通信的哈希表的持续时间为984毫秒。这包括将表放置到内存中并将其删除所需的时间（一次性分配1 GB的内存，这在CUDA中需要花费一些时间），插入和删除元素以及对其进行迭代。还考虑到了视频卡内存中的所有复制。哈希表本身花费了271毫秒。这包括视频卡插入和删除项目所花费的时间，并且没有考虑复制到内存并遍历结果表所花费的时间。如果GPU表寿命很长，或者哈希表完全包含在视频卡的内存中（例如，创建一个哈希表将由另一个GPU代码而不是中央处理器使用），则测试结果是有意义的。由于其高带宽和主动并行化，该视频卡的哈希表显示出高性能。缺点
哈希表体系结构要牢记几个问题：- 聚类会干扰线性探测，因此表中的键远非理想。
- 使用此功能不会删除键，
delete并且随着时间的推移，表会变得混乱。
结果，哈希表的性能可能会逐渐降低，特别是如果哈希表存在很长时间并且在其中执行许多插入和删除操作时。减轻这些缺点的一种方法是以较低的利用率重新哈希新表并在重新哈希时过滤远程键。为了说明所描述的问题，我使用上面的代码创建了一个包含1.28亿个元素的表，我将循环插入400万个元素，直到填充1.24亿个插槽（利用率约为0.96）。这是结果表，每一行都是对CUDA核心的调用，其中在一个哈希表中插入了400万个新元素：随着利用率的提高，生产率下降。在大多数情况下，这是不希望的。如果应用程序将元素插入表中，然后将其丢弃（例如，当计数书中的单词时），那么这不是问题。但是，如果应用程序使用寿命很长的哈希表（例如，在用户经常插入和删除信息时在图形编辑器中存储图像的非空部分），则此行为可能会很麻烦。然后，他测量了6400万次插入后探测哈希表的深度（利用率0.5）。平均深度为0.4774，因此大多数按键位于最佳位置或距最佳位置一个插槽中。最大探测深度为60。然后，我在表中测量了1.24亿个插件的发声深度（利用率0.97）。平均深度已经是10.1757，最大深度是6474（!!）。在高利用率下，线性探测性能会急剧下降。最好将此哈希表保持在较低水平。但是随后我们通过消耗内存来提高生产力。幸运的是，对于32位键和值，这是合理的。如果在上面的示例中，表中保存了1.28亿个元素，利用率系数为0.25，那么我们可以在其中放置不超过3200万个元素，并且剩余的9600万个插槽将丢失-每对8个字节，丢失768 MB的内存。请注意，我们正在谈论视频卡内存的丢失，这是比系统内存更有价值的资源。尽管大多数支持CUDA的现代台式机图形卡至少具有4 GB的内存（在撰写本文时，NVIDIA 2080 Ti具有11 GB），但是丢失此类容量将不是最明智的决定。稍后，我将写更多有关为视频卡创建哈希表的信息，这些哈希表在探测深度以及重用远程插槽的方式方面都没有问题。感应深度测量
要确定按键听起来的深度，我们可以从其实际表索引中提取键哈希（表中的理想索引）：
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
由于附加代码中的两个二进制数具有魔力，并且哈希表的容量等于2的幂，因此即使将键索引移到表的开头，该方法也可以使用。取一个散列为1但插入插槽3的键。然后对于容量为4的表，我们得到的(3 — 1) & 3值等于2。结论
如果您有任何疑问或意见，请在Twitter上给我写信，或在资源库中打开一个新主题。此代码的灵感来自一些出色的文章：将来，我将继续写有关视频卡哈希表的实现，并将分析其性能。我计划使用在方便于视频卡的数据结构中使用原子操作来链接，哈希Robin Hood和布谷鸟哈希。