我把它發佈到了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 位時，性能會略有下降。

哈希表狀態

哈希表中的每個鍵值對可以具有以下四種狀態之一：

• 鍵和值均為空。 在此狀態下，哈希表被初始化。
• 鍵已被記錄，但值尚未被記錄。 如果另一個線程當前正在讀取數據，則它返回空。 這是正常的，如果另一個執行線程早一點工作，也會發生同樣的事情，我們正在討論並發數據結構。
• 鍵和值都會被記錄。
• 該值可供其他執行線程使用，但鍵尚不可用。 發生這種情況是因為 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 哈希和布穀鳥哈希。

來源： www.habr.com