數(shù)十年的計算機科學(xué)史致力于設(shè)計有效存儲和檢索信息的解決方案。哈希圖（或哈希表）是一種流行的信息存儲數(shù)據(jù)結(jié)構(gòu)，因為它對元素的插入和檢索具有攤銷、恒定的時間保證。

然而，盡管哈希圖很流行，但很少在 GPU 加速計算的上下文中討論哈希圖。雖然 GPU 以其大量線程和計算能力而聞名，但其極高的內(nèi)存帶寬使許多數(shù)據(jù)結(jié)構(gòu)（如哈希圖）得以加速。

這篇文章介紹了哈希圖的基本原理，以及它們的內(nèi)存訪問模式如何使其非常適合 GPU 加速。我們將介紹 cuCollections ，這是一個用于并發(fā)數(shù)據(jù)結(jié)構(gòu)（包括哈希圖）的新的開源 CUDA C ++庫。

最后，如果您對在應(yīng)用程序中使用 GPU 加速哈希映射感興趣，我們將提供多列關(guān)系連接算法的示例實現(xiàn)。 RAPIDS cuDF 集成了 GPU 哈希圖，這有助于實現(xiàn) 數(shù)據(jù)科學(xué)工作負(fù)載的驚人加速。要了解更多信息，請參閱 GitHub 上的 rapidsai/cudf 和 使用 Dask 和 RAPIDS 為自然語言處理加速 TF-IDF。

您還可以將 cuCollections 用于表格數(shù)據(jù)處理之外的許多用例，例如推薦系統(tǒng)、流壓縮、圖形算法、基因組學(xué)和稀疏線性代數(shù)操作。

哈希圖基礎(chǔ)知識

哈希映射是 associative 容器，這意味著它們存儲 key – value 對，其中 key 映射到關(guān)聯(lián)的 value ，從而通過查找其鍵來檢索值。例如，您可以使用哈希圖來實現(xiàn)電話簿，方法是使用個人的姓名作為密鑰，將其電話號碼作為關(guān)聯(lián)值。

哈希映射與其他關(guān)聯(lián)容器的不同之處在于，插入或檢索等操作的平均成本是恒定的。 std::map 在 C ++標(biāo)準(zhǔn)模板庫中不是哈希表，但通常實現(xiàn)為二進制搜索樹。 std::unordered_map 更類似于與此討論相關(guān)的哈希表類型。在本文中，哈希表和哈希圖之間沒有區(qū)別。這兩個術(shù)語將在整個過程中互換使用。

單值與多值比較

討論哈希表時的一個重要區(qū)別是是否允許重復(fù)鍵。單值哈希表或哈希映射要求鍵是唯一的（例如，std::unordered_map），而多值哈希表和哈希多映射允許重復(fù)鍵（例如，std::unordered_multimap）。

使用電話簿類比，后者指的是一個人可以擁有多個電話號碼的情況。例如，電話簿可以具有 (k=Alice, v=408-555-0148) 和具有另一個值 (k=Alice, v=408-555-3847) 的重復(fù)密鑰。

存儲和檢索

從概念上講，哈希映射由一個桶數(shù)組組成，其中每個桶可以包含一個或多個鍵值對。要在映射中插入新的對，將向鍵應(yīng)用哈希函數(shù)以生成哈希值。然后使用該哈希值選擇其中一個桶。如果存儲桶可用，則該對存儲在該存儲桶中。

例如，要插入對 (Alice, 408-555-0148) ，您對鍵 hash(Alice)=4, 進行散列以獲取其散列值，并選擇位置 4 處的存儲桶來存儲對。稍后，要檢索與 Alice 關(guān)聯(lián)的值，可以使用相同的哈希函數(shù) hash(Alice), 再次選擇位置 4 處的存儲桶并檢索先前存儲的值。

哈希沖突

如果表中桶的數(shù)量等于可能的鍵的數(shù)量，則可以使用哈希桶和鍵之間的一對一關(guān)系，其中每個鍵正好映射到表中的一個桶。

然而，這在大多數(shù)情況下是不切實際的，因為潛在密鑰的數(shù)量事先不知道，或者為每個密鑰保留存儲桶所需的存儲將超過可用的存儲容量。想象一下，如果你的電話簿必須為宇宙中每個可能的名字保留一個條目！

因此，哈希函數(shù)通常是不完美的，并可能導(dǎo)致哈希沖突，其中兩個不同的鍵映射到相同的哈希值（圖 1 ）。好的哈希函數(shù)尋求最小化沖突的可能性，但在大多數(shù)情況下它們是不可避免的。

圖 1 。兩個不同的鍵， Alice 和 Bob ，具有相同的哈希值，導(dǎo)致桶 4 處的哈希沖突

打開尋址

在文獻中可以找到許多解決哈希沖突的策略，但本文重點介紹了一種名為 open addressing with linear probing 的策略。

開放尋址哈希表使用內(nèi)存中的連續(xù)存儲桶數(shù)組。使用線性探測，如果在位置 i, 處遇到已占用的鏟斗，則移動到下一個相鄰位置 i+1 。如果這個桶也被占用了，你就轉(zhuǎn)到 i+2, ，依此類推。當(dāng)你到達最后一個桶時，你就繞回到開頭。這個所謂的 probing scheme 對于每個密鑰都是確定性的（圖 2 ）。

圖 2 :開放尋址通過探測方案在不同位置存儲沖突條目，探測方案以確定性順序遍歷一系列備選存儲桶

這種方法具有緩存效率，因為它訪問內(nèi)存中的連續(xù)位置。如果 負(fù)載因子（已填充的存儲桶與總存儲桶的比率）較高，則可能會導(dǎo)致性能下降，因為這會導(dǎo)致額外的內(nèi)存讀取。

從地圖中檢索關(guān)鍵字 Bob 的工作方式相同：從位置 hash(Bob)=4 開始遵循關(guān)鍵字的探測順序，直到在位置 6 找到所需的桶。如果在給定鍵的探測序列中的任何一點上遇到空桶，則知道所查詢的鍵不在映射中。

隨機存儲器訪問

精心設(shè)計的哈希函數(shù)通過最大化哈希任意兩個鍵將導(dǎo)致不同哈希值的可能性來最小化沖突次數(shù)。這意味著對于任何給定的兩個鍵，它們對應(yīng)的桶可能位于不同的內(nèi)存位置。

因此，大多數(shù)哈希表操作的內(nèi)存訪問模式實際上是隨機的。為了理解哈希表的性能，了解隨機內(nèi)存訪問的性能非常重要。

表 1 比較了理論峰值帶寬與在現(xiàn)代 GPU s 和 GPU s 上通過 GUPS benchmark 測量的隨機 64 位讀取的實現(xiàn)帶寬。

表 1 . 帶寬計算為訪問大小乘以訪問次數(shù)除以時間

如果您有興趣在系統(tǒng)上運行 GUPS GPU 基準(zhǔn)測試，請參閱 NVIDIA developer blog code samples GitHub 存儲庫。您可以訪問 ParRes/Kernels GitHub 存儲庫中的 CPU 代碼。

如您所見，隨機內(nèi)存訪問大約比理論峰值帶寬慢 10 倍。這是因為內(nèi)存子系統(tǒng)針對順序訪問進行了優(yōu)化。更重要的是， NVIDIA GPU s 的隨機訪問吞吐量比現(xiàn)代 CPU s 的高一個數(shù)量級。這些結(jié)果表明，性能最好的 CPU 哈希表可能比性能最好的 GPU 哈希表慢一個數(shù)量級。

GPU 哈希圖實現(xiàn)

隨機內(nèi)存訪問在哈希表實現(xiàn)中是不可避免的， GPU 在隨機訪問方面優(yōu)于 CPU 。這很有希望，因為它暗示 GPU 應(yīng)該擅長哈希表操作。為了測試這一理論，本節(jié)討論 GPU 哈希表的實現(xiàn)和優(yōu)化，并將性能與 CPU 實現(xiàn)進行比較。

目標(biāo)不是開發(fā)一個標(biāo)準(zhǔn) C ++容器（如std::unordered_map）的替代品，而是專注于實現(xiàn)一個哈希表，該哈希表適用于 GPU 加速應(yīng)用程序中出現(xiàn)的大規(guī)模并行、高吞吐量問題。

本示例使用以下簡化假設(shè)：

表的容量是固定的，不能添加超出初始容量的其他鍵值對

需要將其中一個可能的鍵值設(shè)置為哨兵值，以指示空桶

鍵和值類型的大小之和必須小于或等于 8 個字節(jié)

插入后不能刪除鍵值對

請注意，這些不是基本的限制，可以通過 cuCollections 庫中提供的更高級的實現(xiàn)來克服。

首先，示例哈希表使用開放尋址，由一組桶組成。每個存儲桶可以保存一個鍵 – 值對，并使用鍵/值標(biāo)記進行初始化，以表示當(dāng)前為空。對于沖突解決，使用線性探測 。

GPU 加速哈希表需要支持來自多個線程的并發(fā)更新，例如，如果兩個線程試圖在同一位置插入，則需要采取步驟避免數(shù)據(jù)競爭。為了避免昂貴的鎖定，示例哈希表通過 cuda::std::atomic 其中每個鏟斗定義為cuda::std::atomic>。

為了插入新的密鑰，實現(xiàn)根據(jù)其哈希值計算第一個存儲桶，并執(zhí)行原子比較和交換操作，期望存儲桶中的密鑰等于empty_sentinel。如果是，則插槽為空，插入成功。否則，它前進到下一個桶，直到最終找到一個空桶。

下面的代碼顯示了哈希表插入函數(shù)的簡化版本。

__device__ bool insert(Key k, Value v) {
// get initial probing position from the hash value of the key
auto i = hash(k) % capacity;
while (true) {
  // load the content of the bucket at the current probe position
  auto [old_k, old_v] = buckets[i].load(memory_order_relaxed);
  // if the bucket is empty we can attempt to insert the pair
  if (old_k == empty_sentinel) {
    // try to atomically replace the current content of the bucket with the input pair
    bool success = buckets[i].compare_exchange_strong(
                    {old_k, old_v}, {k,v}, memory_order_relaxed);
    if (success) {
      // store was successful
      return true;
    }
  } else if (old_k == k) {
    // input key is already present in the map
    return false;
  }
  // if the bucket was already occupied move to the next (linear) probing position
  // using the modulo operator to wrap back around to the beginning if we     
  // go beyond the capacity
  i = ++i % capacity;
}
}

在映射中查找特定鍵的關(guān)聯(lián)值的方式類似。沿鑰匙探測順序檢查每個位置，直到找到包含所需鑰匙的存儲桶或空存儲桶，表明鑰匙無法駐留在表中。

合作團體

最初，為每個輸入元素分配一個工作線程似乎是一個合理的比率。但是，請考慮以下事項：

輸入中的相鄰鍵與其在存儲器中的相關(guān)探測位置之間沒有關(guān)系。這意味著扭曲中的每個線程都可能訪問哈希圖的完全不同的區(qū)域。在最壞的情況下，每個探測步驟需要從全局內(nèi)存中的 32 個不同位置加載每個扭曲。（回想隨機內(nèi)存訪問。）

利用線性探測，每個線程可以從其初始探測位置開始訪問多個相鄰?fù)?。這種本地訪問模式將允許使用單個聯(lián)合負(fù)載預(yù)取多個探測位置，不幸的是，這無法通過單個線程實現(xiàn)。

我們能做得更好嗎？對 CUDA cooperative groups 模型可以輕松地重新配置工作分配的粒度。每個輸入元素不使用單個 CUDA 線程，而是將一個元素分配給同一經(jīng)線內(nèi)的一組連續(xù)線程。

對于給定的輸入鍵，不是按順序遍歷其關(guān)聯(lián)的探測序列，而是用單個合并負(fù)載預(yù)取多個相鄰?fù)暗拇翱凇Ｈ缓?，該組使用有效的ballot和shuffle內(nèi)部函數(shù)協(xié)同確定窗口內(nèi)的候選桶。

圖 3 。密鑰 Bob 的組協(xié)作探測步驟及其中間步驟

下面的代碼擴展了之前引入的 insert 函數(shù)，以使用扭曲中的四個連續(xù)線程協(xié)同插入一個鍵。cg::thread_block_tile<4>表示子 rp 中的四個線程。

enum class probing_state { SUCCESS, DUPLICATE, CONTINUE };


__device__ bool insert(cg::thread_block_tile<4> group, Key k, Value v) {
// get initial probing position from the hash value of the key
auto i = (hash(k) + group.thread_rank()) % capacity;
auto state = probing_state::CONTINUE;
while (true) {
  // load the contents of the bucket at the current probe position of each rank in a coalesced manner
  auto [old_k, old_v] = buckets[i].load(memory_order_relaxed);
  // input key is already present in the map
  if(group.any(old_k == k)) return false;
  // each rank checks if its current bucket is empty, i.e., a candidate bucket for insertion
  auto const empty_mask = group.ballot(old_k == empty_sentinel);
  // it there is an empty buckets in the group's current probing window
  if(empty_mask) {
    // elect a candidate rank (here: thread with lowest rank in mask)
    auto const candidate = __ffs(empty_mask) - 1;
    if(group.thread_rank() == candidate) {
      // attempt atomically swapping the input pair into the bucket
      bool const success = buckets[i].compare_exchange_strong(
                      {old_k, old_v}, {k, v}, memory_order_relaxed);
      if (success) {
        // insertion went successful
        state = probing_state::SUCCESS;
      } else if (old_k == k) {
        // else, re-check if a duplicate key has been inserted at the current probing position
        state = probing_state::DUPLICATE;
      }
    }
    // broadcast the insertion result from the candidate rank to all other ranks
    auto const candidate_state = group.shfl(state, candidate);
    if(candidate_state == probing_state::SUCCESS) return true;
    if(candidate_state == probing_state::DUPLICATE) return false;
  } else {
    // else, move to the next (linear) probing window
    i = (i + group.size()) % capacity;
  }
}
}

哈希表插入函數(shù)的前面的代碼示例是 cuCollections cuco::static_map的實際實現(xiàn)的簡化版本。

圖 4 顯示了在 NVIDIA A100 80 GB GPU 上測量的不同組大小和表占用率的非協(xié)作和協(xié)作探測方法的性能，沒有具體化。

圖 4 。通過協(xié)作探測，吞吐量以 GB / s 為單位（越高越好）。紅色虛線顯示了峰值 GUPS 結(jié)果，它提供了在該系統(tǒng)上可以實現(xiàn)的吞吐量上限。

如果負(fù)載系數(shù)較低，則非合作（非 CG ）表現(xiàn)出接近最佳性能。然而，如果負(fù)載因子增加，則吞吐量會由于沖突次數(shù)增加和探測序列較長而急劇下降。這是有問題的，因為較高的表加載系數(shù)對應(yīng)于更好的內(nèi)存利用率。

協(xié)作探測提高了此類高負(fù)載因數(shù)場景的性能。與非合作方法相比，當(dāng)負(fù)載系數(shù)較高時，如果組大小為 4 ，可以觀察到插入吞吐量高出 13% ，查找吞吐量高出 40% 。

長探測序列也出現(xiàn)在具有高密鑰乘數(shù)的多值場景中，因為相同的密鑰遍歷相同的桶序列。合作探測也有助于加快這些場景。

有關(guān)組協(xié)作哈希表探測的更多信息，請參見 Parallel Hashing on Multi-GPU Nodes 和 WarpCore: A Library for Fast Hash Tables on GPUs 。

現(xiàn)有 CPU 和 GPU 哈希圖比較

多年來，已經(jīng)提出了多種 C ++哈希圖實現(xiàn)。其中最流行的是libstdc++/libc++ std::unordered_ map 和 Abseil absl::flat_hash_map。這些是順序?qū)崿F(xiàn)，從多個線程使用它們需要額外的同步。

TBB 中的tbb::concurrent_hash_map和 Folly 中的folly::AtomicHashMap是并發(fā)多線程 CPU 數(shù)據(jù)結(jié)構(gòu)的示例。 GPU s 中可用的少數(shù)實現(xiàn)之一是 Kokkos 庫中的kokkos::UnorderedMap。

將上面提供的映射實現(xiàn)的性能與 cuCollection cuco::static_map進行比較?；鶞?zhǔn)設(shè)置如下。

首先，插入 227（ 1GB ）唯一的 4 字節(jié)密鑰/ 4 字節(jié)值對，然后查詢同一組密鑰以檢索它們的相關(guān)值。每次運行的目標(biāo)工作臺負(fù)載系數(shù)為 50% 。性能以內(nèi)存吞吐量衡量（每秒 GB ；越高越好）。

結(jié)果如圖 5 所示。cuco::static_map在單個 NVIDIA H100-80GB-SXM 上實現(xiàn)了 87.5 GB / s 的插入吞吐量和 134.6 GB / s 的查找吞吐量，這意味著與最快的 CPU 單線程和多線程實現(xiàn)相比，速度提高了超過數(shù)量級。此外， cuCollections 在本測試中的性能優(yōu)于其他 GPU 實現(xiàn)kokkos::UnorderedMap，插入的性能為 3.8 倍，查找的性能為 2.6 倍。

請注意，在這個基準(zhǔn)設(shè)置中，對于 CPU 側(cè)的實現(xiàn)，每個操作的 I / O 向量都駐留在 CPU memory 中，而對于 GPU 側(cè)的實施，則駐留在 GPU memory 中。如果數(shù)據(jù)向量需要駐留在 GPU 哈希映射的 CPU 存儲器中，這將要求首先將輸入數(shù)據(jù)移動到 GPU ，然后將結(jié)果移回 CPU 存儲器。

這可以通過顯式（異步批處理）復(fù)制或使用 CUDA 的 unified memory 概念自動頁面遷移來實現(xiàn)。結(jié)果表明，我們實現(xiàn)的吞吐量始終遠遠高于 H100 上 PCIe Gen4 甚至 PCIe Gen5 的實際可用帶寬。這意味著該方法能夠使 CPU 和 GPU 之間的鏈路完全飽和。

換句話說，即使數(shù)據(jù)不在 GPU 內(nèi)存中， cuCollections 也能以系統(tǒng) PCIe 帶寬的速度構(gòu)建和查詢哈希表。此外，由于 GPU 和 GPU 之間的快速 NVLink-C2C 互連， NVIDIA Grace Hopper Superchip 可以提供額外的加速，釋放哈希表的全部吞吐量。相比之下， CPU 哈希映射通常實現(xiàn)比 PCIe 低得多的吞吐量。

圖 5 。流行 CPU 和 GPU 哈希圖實現(xiàn)的性能比較

多列關(guān)系聯(lián)接示例

本節(jié)以 GPU 哈希表如何用于實現(xiàn)復(fù)雜算法的真實示例為特色。

cuDF 是用于數(shù)據(jù)分析的 GPU 加速庫。它為數(shù)據(jù)操作（如加載、連接和聚合）提供原語。通過利用 cuCollections 哈希表，它使用哈希連接算法來執(zhí)行連接操作。

圖 6 。 RAPID cuDF 中內(nèi)部連接實現(xiàn)的構(gòu)建和探測階段

圖 6 顯示了 cuDF 連接實現(xiàn)如何為內(nèi)部連接工作。 cuDF 提供了一個內(nèi)置的哈希函數(shù)，用于將任意類型的行哈希為哈希值。不同的行可以具有相同的哈希值，因此需要行相等性檢查來確定兩行是否真正相同。

左側(cè)的表格用于填充 cuco::static_multimap 其中鍵是行的哈希值，有效載荷是關(guān)聯(lián)的行索引。行 24 插在鏟斗 47 上，行 25 插在鏟斗 48 上。在探測階段，右表中的行 200 的哈希值為 47 ，這與哈希表中的桶 47 的哈希值（或相同密鑰）相同。

為了最終確定兩行是否相等，將右側(cè)表中的{ Andr é -Marie ， Amp è re }的行索引 200 與左側(cè)表中的｛ Alessandro ， Volta ｝的行索引 24 傳遞給行相等函數(shù) row_equal(200, 24) 。

最后，這兩行并不相同，因此左側(cè)表的第 24 行不匹配。最后，左表的第 25 行與右表的第 200 行匹配，因為哈希值相同，并且行相等性檢查（ row_equal(200, 25) ）也通過。

基準(zhǔn)連接操作是一個復(fù)雜的主題，因為有許多大小、選擇性等選項。有關(guān)詳細(xì)信息，請參見 How to Get the Most out of GPU Accelerated Database Operators 和 Effective, Scalable Multi-GPU Joins 。

如何在代碼中使用 GPU 哈希圖

GPU s 非常適合于哈希圖等并發(fā)數(shù)據(jù)結(jié)構(gòu)。這一切都是從高帶寬存儲器架構(gòu)開始的，對于許多小的隨機讀取和原子更新，該架構(gòu)比 CPU 快一個數(shù)量級。這直接轉(zhuǎn)化為 GPU 上高效的哈希表插入和探測性能。

這篇文章介紹了設(shè)計大規(guī)模并行哈希圖時的幾個重要考慮事項： 1 ）具有開放尋址的哈希桶的平面內(nèi)存布局，以解決沖突。作為 cuCollections 庫的一部分，您可以在 GitHub 上找到快速靈活的哈希圖實現(xiàn)。