數(shù)十年的計(jì)算機(jī)科學(xué)史致力于設(shè)計(jì)有效存儲(chǔ)和檢索信息的解決方案。哈希圖（或哈希表）是一種流行的信息存儲(chǔ)數(shù)據(jù)結(jié)構(gòu)，因?yàn)樗鼘?duì)元素的插入和檢索具有攤銷(xiāo)、恒定的時(shí)間保證。

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

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

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

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

哈希圖基礎(chǔ)知識(shí)

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

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

單值與多值比較

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

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

存儲(chǔ)和檢索

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

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

哈希沖突

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

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

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

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

打開(kāi)尋址

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

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

圖 2 :開(kāi)放尋址通過(guò)探測(cè)方案在不同位置存儲(chǔ)沖突條目，探測(cè)方案以確定性順序遍歷一系列備選存儲(chǔ)桶

這種方法具有緩存效率，因?yàn)樗L問(wèn)內(nèi)存中的連續(xù)位置。如果 負(fù)載因子（已填充的存儲(chǔ)桶與總存儲(chǔ)桶的比率）較高，則可能會(huì)導(dǎo)致性能下降，因?yàn)檫@會(huì)導(dǎo)致額外的內(nèi)存讀取。

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

隨機(jī)存儲(chǔ)器訪問(wèn)

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

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

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

表 1 . 帶寬計(jì)算為訪問(wèn)大小乘以訪問(wèn)次數(shù)除以時(shí)間

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

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

GPU 哈希圖實(shí)現(xiàn)

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

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

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

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

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

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

插入后不能刪除鍵值對(duì)

請(qǐng)注意，這些不是基本的限制，可以通過(guò) cuCollections 庫(kù)中提供的更高級(jí)的實(shí)現(xiàn)來(lái)克服。

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

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

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

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

__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)值的方式類(lèi)似。沿鑰匙探測(cè)順序檢查每個(gè)位置，直到找到包含所需鑰匙的存儲(chǔ)桶或空存儲(chǔ)桶，表明鑰匙無(wú)法駐留在表中。

合作團(tuán)體

最初，為每個(gè)輸入元素分配一個(gè)工作線程似乎是一個(gè)合理的比率。但是，請(qǐng)考慮以下事項(xiàng)：

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

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

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

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

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

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

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的實(shí)際實(shí)現(xiàn)的簡(jiǎn)化版本。

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

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

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

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

長(zhǎng)探測(cè)序列也出現(xiàn)在具有高密鑰乘數(shù)的多值場(chǎng)景中，因?yàn)橄嗤拿荑€遍歷相同的桶序列。合作探測(cè)也有助于加快這些場(chǎng)景。

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

最后，這兩行并不相同，因此左側(cè)表的第 24 行不匹配。最后，左表的第 25 行與右表的第 200 行匹配，因?yàn)楣Ｖ迪嗤?，并且行相等性檢查（ row_equal(200, 25) ）也通過(guò)。

基準(zhǔn)連接操作是一個(gè)復(fù)雜的主題，因?yàn)橛性S多大小、選擇性等選項(xiàng)。有關(guān)詳細(xì)信息，請(qǐng)參見(jiàn) 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)。這一切都是從高帶寬存儲(chǔ)器架構(gòu)開(kāi)始的，對(duì)于許多小的隨機(jī)讀取和原子更新，該架構(gòu)比 CPU 快一個(gè)數(shù)量級(jí)。這直接轉(zhuǎn)化為 GPU 上高效的哈希表插入和探測(cè)性能。

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