GPU無鎖跳步哈希表

張 娟，孫建伶

1.浙江大學(xué) 計算機科學(xué)與技術(shù)學(xué)院,杭州 310027

2.阿里巴巴-浙江大學(xué)前沿技術(shù)聯(lián)合研究中心，杭州 311121

1 引言

圖形處理單元（graphics processing unit，GPU）具有卓越的并行加速能力。將通用內(nèi)存索引結(jié)構(gòu)應(yīng)用到GPU之上成為了一個新的研究方向。目前針對GPU優(yōu)化的內(nèi)存索引結(jié)構(gòu)還較少，只有很少的完全并發(fā)且可動態(tài)更新的結(jié)構(gòu)能夠適應(yīng)GPU。

完全并發(fā)的GPU數(shù)據(jù)結(jié)構(gòu)的應(yīng)用場景更加廣泛，無鎖特性又可以解決傳統(tǒng)基于鎖的方法由于大量駐留線程對資源的爭用而造成的低效率。本文設(shè)計并實現(xiàn)GPU完全并發(fā)且可動態(tài)更新的無鎖跳步哈希表——GPU無鎖跳步哈希表（GPU lock-free hopscotch Hash table，GLHT）。

目前尚未有GPU完全并發(fā)且可動態(tài)更新的跳步哈希表，但是有少許GPU其他哈希表設(shè)計。GPU其他哈希表設(shè)計主要分為兩個方向：靜態(tài)哈希表、完全并發(fā)且可動態(tài)更新的哈希表。據(jù)本文所知，雖然已有多種有效的GPU靜態(tài)哈希表（例如Alcantara等人設(shè)計的杜鵑哈希表[1]），但完全并發(fā)且可動態(tài)更新的GPU哈希表目前只有Misra和Chaudhuri實現(xiàn)的無鎖鏈式哈希表[2]和Ashkiani等人設(shè)計的Slab Hash[3]，并且文獻[2]中的哈希表還不是完全動態(tài)的。

GLHT的基礎(chǔ)數(shù)據(jù)結(jié)構(gòu)是跳步哈希表[4]。跳步哈希表的插入操作保持數(shù)據(jù)的緊湊。當發(fā)生數(shù)據(jù)沖突時，新數(shù)據(jù)插入到哈希槽（哈希槽即指鍵原始應(yīng)該被哈希到的槽）隨后的H個槽，這H個槽稱為當前槽的鄰域，H是用戶設(shè)置的常數(shù)。每個槽關(guān)聯(lián)一個由H+1個bit組成的bitmap，指示當前槽和后續(xù)H個槽中的項是否是最初哈希到當前槽的項。若某個槽的項本來應(yīng)該哈希到前面的槽，則稱這個槽“從屬”于前面的那個槽。圖1是鍵v插入跳步哈希表的過程，白色表示空槽，灰色表示槽中有項，該哈希表的H為3。鍵v本應(yīng)哈希到槽6，但是發(fā)生了數(shù)據(jù)沖突。于是，首先通過線性探測找到距槽6最近的空槽13。如果兩個槽的距離小于等于H，則可以將鍵直接插入到該空槽中，但是槽13到槽6超過了H的范圍，因此需要按照鄰域從屬關(guān)系，置換它們之間的鍵，將空槽移近槽6。觀察槽10(13-H=10)的bitmap，發(fā)現(xiàn)只有槽11從屬于槽10，于是置換槽11的鍵w到槽13并更新槽10的bitmap?，F(xiàn)在空槽為槽11，但它仍然不在槽6的鄰域內(nèi)，于是觀察槽8(11-H=8)的bitmap，發(fā)現(xiàn)槽9從屬于槽8，于是置換鍵z到槽11并更新槽8的bitmap?，F(xiàn)在，槽9在槽6的鄰域范圍內(nèi)了，可以直接將鍵v安排在槽9。通過這一系列的移位操作，跳步哈希表保證了數(shù)據(jù)與原始哈希槽的距離不會大于H，因此查找時只需檢查哈希槽及其鄰域中是否有目標鍵，若無則可確定目標鍵不存在，由此保證任何情況下的查找時間都是O(1)。

Fig.1 Insert key vinto hopscotch Hash table圖1 鍵v插入跳步哈希表

在GPU中，若一個warp內(nèi)的線程請求訪問連續(xù)對齊的內(nèi)存塊，則會進行合并訪問（coalesced access）以便最大化內(nèi)存帶寬。跳步哈希表的所有操作恰好都只需要并行讀取連續(xù)內(nèi)存范圍內(nèi)的哈希槽和鄰域，因此可以使用高效的GPU合并訪問完成讀取請求。而其他哈希表，例如杜鵑哈希表[5]，在插入過程中反而追求項的隨機分布，自然不利于合并訪問的使用。

設(shè)計實現(xiàn)GPU哈希表并不是直接將原有的CPU哈希表簡單地放置到GPU上，不僅需要考慮GPU環(huán)境下的并發(fā)安全問題，還要結(jié)合GPU的硬件特性，實現(xiàn)哈希表在GPU上的并行性能最大化。GLHT的設(shè)計主要圍繞兩方面：

（1）warp內(nèi)并行：采用warp協(xié)同工作共享策略（warp-cooperative work sharing strategy），減少程序控制流中的分支與發(fā)散，以實現(xiàn)對哈希表單個操作的并行加速。

（2）warp間完全并發(fā)：全局內(nèi)存配合CUDA（compute unified device architecture）原子操作atomic-CAS以及特殊的并發(fā)控制策略設(shè)計，在實現(xiàn)完全并發(fā)和無鎖特性的同時，保證了讀操作的無等待特性,以實現(xiàn)哈希表多個操作的并發(fā)執(zhí)行。

本文進行了實驗評估，結(jié)果表明GLHT具有在靈活性和性能上的優(yōu)勢。GLHT與其他GPU靜態(tài)哈希表相比，具有可以接受的構(gòu)建和檢索速度；與現(xiàn)有的CPU跳步哈希表相比，具有4～9倍的性能優(yōu)勢；比采取預(yù)先分配內(nèi)存的GPU無鎖鏈式哈希表[2]更加靈活，并且在寫操作較多的工作負載中獲得了更好的性能。

本文工作安排如下：第2章介紹GPU數(shù)據(jù)結(jié)構(gòu)相關(guān)工作；第3章描述GLHT的總體設(shè)計；第4章介紹GLHT的實現(xiàn)細節(jié)；第5章為實驗評估；第6章對全文進行總結(jié)。

2 相關(guān)工作

目前有多種GPU靜態(tài)哈希表。Alcantara等人的杜鵑哈希表[1]在批量構(gòu)建階段和檢索階段都有很好的性能，但隨著負載因子要求的增加，批量構(gòu)建過程越來越有可能失敗。該哈希表已用于CUDA數(shù)據(jù)并行原語庫（CUDA data parallel primitives library，CUDPP）[6]。García等人[7]提出了一種基于Robin hood的哈希方法，他們專注于更高的負載因子并利用了圖形應(yīng)用程序的空間局部性，但代價是該哈希方法與杜鵑哈希相比性能有所下降。Khorasani等人[8]提出了Stadium Hashing（Stash）技術(shù)，它也是一種杜鵑哈希表設(shè)計，可以擴展為大型哈希表。它解決的重點問題是out-of-core哈希表不能完整地放進單個GPU內(nèi)存中。通過將表容器存儲在CPU內(nèi)存中，Stash消除了將哈希表整個維護在有限的GPU內(nèi)存上的限制。Stash使用了名為ticket-board的緊湊數(shù)據(jù)結(jié)構(gòu)，這個數(shù)據(jù)結(jié)構(gòu)引導(dǎo)了哈希表上的所有操作。在最好的情況下（即空表），Stash的插入操作只需要一個原子操作和一個常規(guī)的內(nèi)存寫操作，查找操作則至少需要兩個內(nèi)存讀取操作。雖然各種靜態(tài)哈希表的側(cè)重有所不同，但文獻[1]似乎是這些設(shè)計中具有最佳性能指標的通用in-core哈希表。

在GPU完全并發(fā)且可動態(tài)更新的哈希表研究方面，Misra和Chaudhuri[2]測試了幾種已知的CPU無鎖數(shù)據(jù)結(jié)構(gòu)移植到GPU后的加速情況。他們實現(xiàn)了一個GPU上的無鎖鏈表，并由此實現(xiàn)了無鎖鏈式哈希表，這個哈希表能夠支持并發(fā)的插入、刪除和查找操作。但該實現(xiàn)實際上仍然不是完全動態(tài)的，因為在它的實驗中，為將來所有的插入操作都預(yù)先分配了一個結(jié)點資源數(shù)組（必須在編譯時知道），并且不能在運行時動態(tài)分配新項和釋放已刪除項，這就是所謂的“預(yù)先分配內(nèi)存”，而本文實現(xiàn)的GLHT則完全不需要這樣的過程，因此更具靈活性。Cederman等人[9]對各種已知的基于鎖和無鎖的Queue實現(xiàn)進行了類似文獻[2]的實驗，他們得出的結(jié)論是：Queue面向GPU的并行優(yōu)化將有利于性能的提升?，F(xiàn)在，人們也開發(fā)出了一些更簡單的、專為GPU設(shè)計的數(shù)據(jù)結(jié)構(gòu)，例如隊列[10]和鏈表[11]。此外，graph-based算法也使用優(yōu)化的GPU實現(xiàn)了速度的加快[12-14]。受文獻[2]的啟發(fā)，Moscovici等人[15]提出了基于細粒度鎖的GPU友好的跳表（GPU-friendly skip list，GFSL），該工作主要考慮的是GPU的優(yōu)選合并內(nèi)存訪問（preferred coalesced memory accesses）。

最近，Ashkiani等人[3]設(shè)計了一種完全并發(fā)的GPU動態(tài)無鎖鏈式哈希表——Slab Hash。他們認為，GFSL無論在插入、刪除還是查找操作中，都無法擊敗Slab Hash的性能峰值。

3 設(shè)計

GLHT通過warp內(nèi)并行實現(xiàn)對單個操作的并行加速，通過warp間并發(fā)實現(xiàn)多個操作的并發(fā)執(zhí)行。

3.1 warp內(nèi)并行：warp協(xié)同工作共享策略

GPU運行時，各個線程塊被分配給不同的流式多處理器（streaming multiprocessors，SM）執(zhí)行。SM會以32個線程為一組執(zhí)行線程塊操作，這稱為warp調(diào)度。一個warp中的線程從相同的程序計數(shù)器開始執(zhí)行，但是也可以獨立地進行分支與發(fā)散（branch and diverge）。如果一個warp內(nèi)的線程由于判斷條件的不同而進行了分支，則warp將依次執(zhí)行每個線程所采用的分支路徑。當所有的分支路徑被執(zhí)行完時，warp中的線程才會重新聚到共同路徑中。

在GPU上執(zhí)行一組獨立操作的傳統(tǒng)方法是讓每個線程都獨立處理一個操作，例如，GPU上經(jīng)典的鏈表操作[2]。圖2描繪了傳統(tǒng)方法的執(zhí)行過程，圖中空白的時間塊表明當線程在處理分支時，其他線程將處于等待狀態(tài)。頻繁的控制流發(fā)散將會嚴重影響執(zhí)行性能，由此可知，這種傳統(tǒng)方法并沒有充分發(fā)揮出GPU線程的并行能力。

Slab Hash[3]和 warp-wide直方圖計算[16]，讓 warp內(nèi)的線程協(xié)同地并行工作，可以指定warp內(nèi)線程，使用一些warp-wide指令，協(xié)同處理同一個操作，也就是將原本分配給不同線程的操作統(tǒng)一分配給整個warp來處理，如圖3這種方法就稱為warp協(xié)同工作共享策略。warp-wide指令指的是NVIDIA GPU支持的一組內(nèi)建函數(shù)，可以協(xié)同warp內(nèi)線程的通信過程以減少分支與發(fā)散。與傳統(tǒng)的單線程獨立處理相比，warp協(xié)同工作共享策略顯著減少GPU程序中的分支與發(fā)散。

3.2 warp間完全并發(fā)：全局內(nèi)存配合CUDA原子操作

如圖4，雖然GLHT讓warp大小的整個線程塊內(nèi)的線程協(xié)同地處理同一個操作任務(wù)，但不同線程塊之間仍然是操作獨立且完全并發(fā)。

Fig.2 Traditional method圖2 傳統(tǒng)方法

Fig.3 warp-cooperative work sharing strategy圖3 warp協(xié)同工作共享策略

Fig.4 Fully concurrent operations between warp圖4 warp間完全并發(fā)的操作

如何做到warp間完全并發(fā)，首先需要考慮操作執(zhí)行在GPU內(nèi)存的哪個層次。GPU的內(nèi)存結(jié)構(gòu)分為三個層次：可以被設(shè)備內(nèi)所有線程訪問的大的全局內(nèi)存；每個線程塊有著的更小但更快的共享內(nèi)存；線程塊中每個線程的本地寄存器。共享內(nèi)存很?。ㄍǔ?6 KB），并且它進行了分區(qū)，因此來自不同塊的線程無法訪問另一個塊的共享內(nèi)存。GPU的全局內(nèi)存容量大，可供所有線程訪問。由于數(shù)以百萬計的線程可以執(zhí)行GPU內(nèi)核函數(shù)，但只有有限數(shù)量的SM存在，因此線程塊需要排隊等待SM。因此，除了內(nèi)核函數(shù)結(jié)束的時候，并沒有辦法可以全局地同步所有線程。為了實現(xiàn)warp間操作的完全并發(fā)，GLHT通過全局內(nèi)存實現(xiàn)各線程對所有數(shù)據(jù)狀態(tài)的共享。

GLHT選擇無鎖樂觀并發(fā)控制。這種控制方法會在訪問內(nèi)存資源時“樂觀地”假設(shè)沒有并發(fā)沖突，對數(shù)據(jù)不加鎖就直接拿來用，在最后真正更新數(shù)據(jù)時再判斷沖突是否發(fā)生。選擇這種并發(fā)控制方法的好處：一是在GPU編程環(huán)境中，鎖的設(shè)計代價非常昂貴；二是它可以減少成千上萬的駐留線程對鎖資源的爭用，從而提高執(zhí)行效率。而這種并發(fā)控制方法的缺點是，當數(shù)據(jù)沖突發(fā)生時，解決沖突的代價較大，除非沖突發(fā)生的幾率很小。

常見的無鎖編程一般基于原子操作。常用的原子操作是比較和設(shè)置（compare-and-set，CAS）操作。CAS操作將內(nèi)存數(shù)據(jù)與給定值進行比較，只有當它們相同時，才會將該內(nèi)存數(shù)據(jù)修改為新值。GLHT就用到了CUDA的CAS原子操作atomicCAS。

4 實現(xiàn)

GLHT的實際數(shù)據(jù)結(jié)構(gòu)是一個在GPU內(nèi)存中的unsigned long long int數(shù)組，而對GLHT查找操作、刪除操作和插入操作的具體實現(xiàn)細節(jié)感興趣的讀者可自行閱讀代碼及注釋（https://github.com/fanny2011/GLHT），本章僅作簡要介紹。

4.1 拆分操作階段

首先將插入操作拆分為不同的階段并區(qū)分不同階段的槽角色。拆分階段是為了細分并發(fā)操作的粒度，而區(qū)分槽角色只是為了描述的方便。

插入操作可以拆分成find、find_empty、update和find_closer_empty四個階段，其中find_closer_empty階段又可以循環(huán)多個swap_value_into_empty階段，如圖5。而GLHT的刪除操作則只用分為find和update兩個階段。

Fig.5 Phase decomposition of insert operation圖5 插入操作的階段分解

下面描述插入操作不同階段的槽角色，先將哈希槽以角色hash_pos表示。find階段找出哈希表中是否已有相等的鍵，沒有則執(zhí)行find_empty階段。插入操作和刪除操作的find階段與查找操作做的是相同的事情。

find_empty階段返回正好為空的hash_pos或后方最靠近hash_pos的空槽，若空槽為hash_pos或在hash_pos鄰域內(nèi)，則將此空槽以角色target表示，并執(zhí)行update階段，否則執(zhí)行find_closer_empty階段。

插入操作的update階段將目標鍵通過atomic-CAS放進hash_pos，而刪除操作的update階段將target通過atomicCAS置為空，update階段的槽角色如圖6，注意target可能與hash_pos重合。

Fig.6 Slot role in update phase圖6 update階段的槽角色

find_closer_empty階段的目標是將找到的空槽向前移動一次，find_closer_empty階段循環(huán)多個swap_value_into_empty階段，直到移動成功。swap_value_into_empty階段每次對一塊置換區(qū)域操作，置換區(qū)域的第一個槽以角色swap_head表示，置換區(qū)域的最后一個槽即前面找到的那個空槽。從前到后在置換區(qū)域中尋找一個“從屬”于swap_head的槽，將這個槽以角色swap表示，并置換target和swap的項，置換完成則find_closer_empty階段也完成了；但若沒有找到swap，則find_closer_empty階段將角色swap_head向后推動一個位置，并循環(huán)swap_value_into_empty階段。find_closer_empty階段最初將swap_head定在target前的第H個位置。swap_value_into_empty階段的槽角色如圖7所示，注意swap_head與swap可能重合。

Fig.7 Slot role in swap_value_into_empty phase圖7 swap_value_into_empty階段的槽角色

4.2 鎖標記

GLHT采用樂觀并發(fā)控制，在數(shù)據(jù)項上設(shè)置鎖標記，操作時使用原子操作來更改這些鎖標記，以達到使用原子操作鎖定數(shù)據(jù)項的目的。對應(yīng)于上一節(jié)所述的四個角色（hash_pos、target、swap_head和swap），GLHT設(shè)計了兩種鎖標記：multiple_lock和swap_lock。規(guī)定鎖標記間的互斥關(guān)系及它們對并發(fā)讀寫操作的互斥性質(zhì)，就能保證warp間操作的完全并發(fā)安全性。

multiple_lock的含義如下：

（1）當其標記在非空槽時，表示該槽正處于插入或刪除操作的update階段的hash_pos角色。

（2）當其標記在空槽時，有兩種可能：①該槽正處于swap_value_into_empty階段的target角色；②該槽正處于插入操作的update階段的target角色。

swap_lock的含義如下：

表示該槽正處于swap_head角色，或表示該槽正處于swap角色。

兩種鎖標記均為排他鎖標記，即當槽帶上上述標記后，不能再帶上另外的鎖標記，也不能重復(fù)帶上相同的標記。

GLHT查找操作不涉及對鎖標記的操作，只讀取鎖標記的狀態(tài)，根據(jù)hash_pos中的項是否帶有multiple_lock鎖標記（項帶有的swap_lock可以忽略），決定重讀或繼續(xù)下一個步驟。

刪除操作在update階段，若發(fā)現(xiàn)hash_pos帶有任何鎖標記，就需要從頭重試整個操作；否則，為hash_pos帶上multiple_lock，以表明本操作對該hash_pos及其領(lǐng)域擁有了操作權(quán)，其他操作發(fā)現(xiàn)鎖標記的狀態(tài)改變后只能重試。然后，刪除操作將target改變?yōu)榭詹?。最后，操作收尾，取消hash_pos上的multiple_lock。期間，任何一個原子操作失敗后，都需要清理鎖標記并從頭重試整個操作。

Fig.8 Operations on lockflag during insert operation圖8 插入操作過程中對鎖標記的操作

插入操作過程中對鎖標記執(zhí)行的操作與刪除操作類似，但更為復(fù)雜，如圖8。首先，在find_empty階段開始前，需要在hash_pos上帶上multiple_lock。發(fā)現(xiàn)target后，也要為它帶上multiple_lock，這是因為后續(xù)可能伴隨著find_closer_empty階段，這個階段持續(xù)時間較長，所以需要提前搶占這個槽，保持它只能被本操作讀寫。在swap_value_into_empty階段，首先，為swap_head帶上swap_lock；然后，為swap帶上swap_lock；接著，在target中填入swap的項，同時取消target的multiple_lock；將swap變?yōu)榭詹?，同時取消swap_lock并帶上multiple_lock，在下一階段，這個槽就成為了新的target；最后，取消swap_head的swap_lock。在插入操作期間，任何一個原子操作失敗后，都需要按帶上時的倒序清理鎖標記并從特定階段的開頭重試操作。

之所以設(shè)計了兩種互斥鎖，原因在于multiple_lock與swap_lock在讀-寫互斥關(guān)系的表達上是不同的：multiple_lock用作寫-寫互斥和讀-寫互斥，即當發(fā)現(xiàn)槽帶有multiple_lock時，這對該槽及其“從屬”槽的操作，無論讀操作還是寫操作，都需要不斷重試直到multiple_lock消失。swap_lock只用作寫-寫互斥，即對帶有這個標記的槽，任何針對該槽及其“從屬”槽的寫操作都需要重試，但讀操作可以忽略它。

4.3 warp內(nèi)并行：warp-wide指令的使用

應(yīng)用warp協(xié)同工作共享策略的GLHT使用了warp-wide指令shuffle、ballot和ffs。shuffle指令允許線程直接讀取同一個warp內(nèi)的其他線程的寄存器值，這種通信方式比通過訪問共享內(nèi)存進行線程間通信的效果更好、延遲更低，同時也不用消耗額外的內(nèi)存資源來執(zhí)行數(shù)據(jù)交換。ballot指令的作用是在warp內(nèi)線程間進行投票，也常用于讓線程根據(jù)同名變量了解其他線程所處的狀態(tài)。每個線程將同名變量作為輸入，ballot指令將判斷這些變量是否等于零，比較結(jié)果將統(tǒng)一廣播給每一個線程，若比較結(jié)果的第N位被置為1，則表示該warp內(nèi)的第N個線程處于活動狀態(tài)且它的變量非零。ffs指令返回輸入的最低有效位（即最低為1的bit）的下標，下標從1開始，減去1即得到真正的最低有效位下標，這個指令通常會搭配ballot指令。

warp-wide指令以32個線程為一組執(zhí)行操作，因此，GLHT設(shè)置線程塊大小為32，運用warp協(xié)同工作共享策略對槽數(shù)組進行操作。相應(yīng)的，設(shè)置常數(shù)H=31，使得GLHT可以以線程塊為單位對整個鄰域進行操作。

4.3.1 查找操作

1.__device__void Find(LLkey,LL*result,Slot*position){

2.hash=Hash(key);

3.do{

4.*position=table[hash+threadIdx.x];

5.*hash_pos=__shfl(*position,0);

6.}while((*hash_pos& MULTIPLE_LOCK_MASK)!=0);

7.bitmap=getBitmap(*hash_pos);

8.if(isValid(*position,bitmap)

9.&&(((*position)&EMP_FLAG_MASK)==0)

10.&&isEqual(*position,key)

11.&&(getHash(*position)==hash)){

12.predict=1;

13.}

14.ans=__ffs(__ballot(predict));

15.if(ans==0){

16.*result=WRONG_POS;

17.}else{

18.*result=hash+(ans-1);

19.}

20.}

以上是查找操作Find的偽代碼，其中MULTIPLE_LOCK_MASK用來判斷槽是否帶有multiple_lock鎖標記，EMP_FLAG_MASK用來判斷是否為空槽。

warp內(nèi)的每個線程根據(jù)其帶有的threadIdx確定其應(yīng)該讀取的槽，threadIdx指示線程在warp內(nèi)的下標，線程應(yīng)該讀取的槽position與下標為hash的槽hash_pos的偏移正好與threadIdx.x相對應(yīng)（line 4）。雖然讀取的槽不同，但每個線程都需要對hash_pos中數(shù)據(jù)進行條件判斷，因此此時會使用shuffle指令將第一個線程讀取到的hash_pos中數(shù)據(jù)分發(fā)給其他線程（line 5）。

GLHT執(zhí)行查找操作時，會反復(fù)讀取hash_pos和領(lǐng)近槽的數(shù)據(jù)，并檢查hash_pos中的項是否帶有multiple_lock鎖標記（line 3～6）。

在檢查是否有等于查找鍵的槽數(shù)據(jù)時，首先根據(jù)判斷條件設(shè)置同名變量predict（line 8～13），然后用ballot和ffs指令并行地對所有warp內(nèi)線程持有的數(shù)據(jù)進行同步判斷（line 14）。

4.3.2 插入操作

插入操作使用了shuffle指令，主要用在兩方面：（1）在所有warp內(nèi)線程間同步hash_pos數(shù)據(jù)；（2）在某一個線程執(zhí)行atomicCAS操作后，將表達操作最終成功與否的變量廣播給其他線程以同步地推進所有warp內(nèi)線程的控制流判斷。

除了shuffle指令，插入操作還會在以下情況使用ballot和ffs指令組合：（1）在find_empty階段，找到最靠近hash_pos的不帶multiple_lock鎖標記的空槽；（2）在swap_value_into_empty階段，找到不帶任何鎖標記的swap槽。插入操作的ballot和ffs指令組合的具體使用方式與查找操作的方式（見偽代碼line 8～14）相似。

4.3.3 刪除操作

GLHT的刪除操作只使用到了與插入操作一樣的shuffle指令使用方式。

4.4 warp間完全并發(fā)：特殊的并發(fā)控制策略

之前的相關(guān)工作也使用了全局內(nèi)存配合CUDA原子操作，但全局內(nèi)存訪問速度要比共享內(nèi)存慢幾個數(shù)量級。為了提升性能，GLHT在此基礎(chǔ)上設(shè)計了特殊的并發(fā)控制策略（包括4.2節(jié)描述的鎖標記和暫時重復(fù)策略兩方面），保證了讀操作的無等待特性，在一定程度上彌補了全局內(nèi)存訪問慢的缺點。

在GLHT插入操作的swap_value_into_empty階段，需要置換target和swap的項。但是這個置換過程并非原子過程，且GLHT沒有結(jié)構(gòu)鎖，其他warp的讀操作很容易發(fā)生在置換過程的各個操作之間，很可能出現(xiàn)其他warp在讀取swap_head及其“從屬”槽時，讀取不到swap中有效鍵的情況。為此，GLHT設(shè)計了“暫時重復(fù)策略”，即先將swap中的項復(fù)制到target中，再將swap置為空。雖然造成了短暫的項重復(fù)，但保證了數(shù)據(jù)的正確性（即warp不會出現(xiàn)讀取不到正確存儲在表中的有效值）和讀取操作不需要等待的設(shè)計要求。

GLHT的查找操作不涉及任何原子操作，因此可以保證在有限的步驟內(nèi)完成，因此是無等待的，實際上所有的讀操作都是無等待的。除了hash_pos的multiple_lock，任何其他的寫操作標記都不會影響查找操作的進程，從而消除了讀操作與寫操作對于資源的互斥等待。這也正是GLHT將swap_lock和multiple_lock分開設(shè)計的原因，就是為了提高讀操作效率。無論是鍵-值對映射還是鍵集合操作，從統(tǒng)計經(jīng)驗上來說，應(yīng)用程序的讀操作數(shù)量相對寫操作會多一些，因此，提高讀操作效率對提高整體操作效率是非常有意義的。

需要強調(diào)的是，GLHT的設(shè)計方案只能保證GPU上數(shù)據(jù)結(jié)構(gòu)的無鎖并發(fā)安全性，CPU上無法實現(xiàn)相同的安全效果。這是因為warp內(nèi)的并行模式保證了多個位置的內(nèi)存讀操作是真正并行的，相當于在同一時間給多個位置的內(nèi)存狀態(tài)做了一個快照，后續(xù)所有對這些內(nèi)存的聯(lián)合判斷都相當于是在同一時間內(nèi)完成的，從而保證了操作的并發(fā)安全，而CPU無法做到這一點。

5 實驗評估

本實驗全部在Intel Xeon E5-2620服務(wù)器上執(zhí)行，該服務(wù)器擁有1個Socket，每個Socket有6個核，每個核有2個超線程。內(nèi)存為2×16 GB DDR3 SDRAM。高速緩存為32 KB L1數(shù)據(jù)緩存，32 KB L1指令緩存，256 KB L2緩存，15 360 KB L3緩存。操作系統(tǒng)為64位的Ubuntu 16.04.3。CPU代碼采用打開O3優(yōu)化的gcc-5.4.0編譯器編譯。GPU部分是在NVDIA GeForce GTX 1080上進行評估比較的，GDDR5X容量為8 GB。CUDA代碼采用CUDA 8.0編譯器（V8.0.61）編譯。

實驗評估分為兩方面：首先是靜態(tài)基準，以兩個操作階段（批量構(gòu)建和檢索）分步執(zhí)行的方式與其他GPU靜態(tài)哈希表（線性探測、平方探測和CUDPP的杜鵑哈希實現(xiàn)[6]）進行比較；其次是動態(tài)并發(fā)基準，以并發(fā)執(zhí)行隨機混合操作（插入、刪除和查找操作按比例混合）的方式與CPU跳步哈希表和Misra和Chaud-huri實現(xiàn)的完全并發(fā)且可動態(tài)更新的GPU無鎖鏈式哈希表[2]進行比較。

5.1 靜態(tài)基準

GPU靜態(tài)哈希表有兩個操作階段：（1）批量構(gòu)建階段，給定一個固定的負載因子（可以簡單地按照預(yù)先設(shè)計的內(nèi)存使用率來表示）和一個鍵-值對輸入數(shù)組，以批量的插入操作構(gòu)建整個數(shù)據(jù)結(jié)構(gòu)，若構(gòu)建階段發(fā)生插入失敗則需要從頭重建。（2）檢索階段，在批量構(gòu)建階段結(jié)束后，以鍵數(shù)組作為輸入，在數(shù)據(jù)結(jié)構(gòu)中執(zhí)行批量的查找操作，并將返回找到的對應(yīng)的值存儲在輸出數(shù)組中。

本實驗基準以吞吐量（操作總數(shù)量/執(zhí)行時間）作為衡量數(shù)據(jù)結(jié)構(gòu)性能的指標。所有數(shù)據(jù)結(jié)構(gòu)選取的槽數(shù)組都是大小一致的，并固定內(nèi)存使用率為0.8。各數(shù)據(jù)結(jié)構(gòu)的哈希函數(shù)也保持一致。操作總數(shù)作為橫坐標。GPU數(shù)據(jù)結(jié)構(gòu)的線程數(shù)量就等于操作總數(shù)量。在確定GPU數(shù)據(jù)結(jié)構(gòu)的線程數(shù)量后，需要決定每個線程塊的線程數(shù)量（線程塊數(shù)量=線程總數(shù)/每個線程塊的線程數(shù)量）。

圖9是各數(shù)據(jù)結(jié)構(gòu)的構(gòu)建速度比較。GLHT雖然比線性探測和平方探測靜態(tài)哈希表慢，但作為動態(tài)哈希表，它的速度基本上還是可以接受的。

Fig.9 Comparison on build speed圖9 構(gòu)建速度比較

預(yù)設(shè)所有檢索鍵都已存在于數(shù)據(jù)結(jié)構(gòu)。圖10是各數(shù)據(jù)結(jié)構(gòu)的檢索速度比較。與其他靜態(tài)哈希表相比，GLHT的速度仍然較為合理。

5.2 動態(tài)并發(fā)基準

Fig.10 Comparison on retrieve speed圖10 檢索速度比較

文獻[4]已提出了CPU跳步哈希表的并發(fā)版本，后續(xù)實驗中以CPU lock-based hopscotch表示。此外，為了與之前他人提出的GPU哈希表進行比較，本實驗基準選擇了Misra和Chaudhuri提供的GPU上的無鎖鏈式哈希表[2]作為參照。注意到文獻[2]的槽數(shù)組實際是鏈表結(jié)點的指針數(shù)組，操作過程中需要動態(tài)地為鏈表結(jié)點進行內(nèi)存分配。文獻[2]稱，為了確保性能評估可以集中在數(shù)據(jù)結(jié)構(gòu)本身可實現(xiàn)的原始吞吐量上而不受內(nèi)存分配開銷的任何干擾，在GPU內(nèi)核函數(shù)啟動之前從CPU預(yù)先分配了足夠數(shù)量的鏈表結(jié)點到GPU內(nèi)存中，以便并發(fā)操作過程中不從GPU調(diào)用動態(tài)內(nèi)存分配。本文把這個過程稱為“預(yù)先分配內(nèi)存”。這么做的原因是，在操作過程中從GPU調(diào)用動態(tài)內(nèi)存分配是非常耗時的事情。但GLHT不需要這樣的預(yù)分配過程和相關(guān)的耗時操作，因此更具有靈活性。若以文獻[2]不計算“預(yù)先分配內(nèi)存”的執(zhí)行時間與GLHT直接相比，GLHT的優(yōu)勢將不能體現(xiàn)，因此為了公平，本實驗將同時考慮文獻[2]的不計算“預(yù)先分配內(nèi)存”的情況（以GPU chained without allocation表示）和計算“預(yù)先分配內(nèi)存”的情況（以GPU chained with allocation表示）下的吞吐量。

本實驗基準以吞吐量（操作總數(shù)量/執(zhí)行時間）作為衡量數(shù)據(jù)結(jié)構(gòu)性能的指標。數(shù)據(jù)結(jié)構(gòu)的性能可能取決于不同操作的混合比例、鍵的取值范圍以及操作總數(shù)量。為評估不同的操作組合，將不同混合比例表示為三元組[x,y,z]，表示具有x%的插入操作、y%的刪除操作和z%的查找操作。本實驗選取了兩個操作組合，[20,20,60]和[40,40,20]。為評估鍵的取值范圍，在每個操作組合上設(shè)計4個不同的整數(shù)鍵范圍，[0,100]，[0,1 000]，[0,10 000]和[0,100 000]。操作總數(shù)固定為100 000。每個測試的操作序列都是根據(jù)混合比例和總數(shù)量預(yù)先生成的，操作鍵從被評估的鍵范圍中隨機生成。每個測試都需要在GPU上或CPU上評估3次，并且以中值作為其真實執(zhí)行時間。所有數(shù)據(jù)結(jié)構(gòu)選取的槽數(shù)組大小都是固定一致的，哈希函數(shù)也保持一致。線程數(shù)量對于CPU數(shù)據(jù)結(jié)構(gòu)的執(zhí)行性能來說，并不是越多越好。在本實驗環(huán)境下，為CPU數(shù)據(jù)結(jié)構(gòu)選擇了達到最佳性能的線程數(shù)16。而GPU數(shù)據(jù)結(jié)構(gòu)的線程數(shù)量是根據(jù)每次測試的操作總數(shù)量決定的，文獻[2]稱每個線程執(zhí)行一個操作時效果是最好的，于是GPU數(shù)據(jù)結(jié)構(gòu)的線程數(shù)量就等于操作總數(shù)量。文獻[2]選擇每個線程塊512個線程，而GLHT根據(jù)設(shè)計方案選擇每個線程塊32個線程。

操作組合[20,20,60]偏向讀操作。從圖11可以看出，雖然GPU chained without allocation具有明顯的性能優(yōu)勢，但是計算上預(yù)先分配內(nèi)存時間后的GPU chained with allocation恰是執(zhí)行時間最長的，實際上GLHT對GPU chained with allocation有200倍左右的性能提升。并且，隨著鍵范圍的增大，GPU chained without allocation對GLHT的性能優(yōu)勢也沒有那么明顯了，從2、3倍的優(yōu)勢降低到了1倍多。而GLHT對CPU lock-based hopscotch大概有4、5倍的性能優(yōu)勢。

Fig.11 Comparison on throughput of combination[20,20,60]圖11 組合[20,20,60]的吞吐量對比

操作組合[40,40,20]偏向?qū)懖僮?。依舊是GPU chained without allocation比較具優(yōu)勢，而GPU chained with allocation最差。但是從圖12可以看到，與操作組合[20,20,60]相比，GLHT的優(yōu)勢越來越明顯，GPU chained without allocation對GLHT僅有1倍多的性能優(yōu)勢，甚至在鍵范圍較大的情況下，存在GLHT性能超越GPU chained without allocation的現(xiàn)象；GLHT對GPU chained with allocation有200～400倍的性能比；另一方面，GLHT依舊具有對基于鎖的CPU跳步哈希表的優(yōu)勢，且優(yōu)勢擴大到了5～9倍。

Fig.12 Comparison on throughput of combination[40,40,20]圖12 組合[40,40,20]的吞吐量對比

從以上實驗數(shù)據(jù)可以明顯看出，無論是讀操作比重較大的情況還是寫操作比重較大的情況，本章實現(xiàn)的GLHT對CPU上的跳步哈希表具有絕對的性能優(yōu)勢（4～9倍）。

至于文獻[2]提供的GPU上的無鎖鏈式哈希表，雖然它也支持并發(fā)的插入、刪除和查找操作，但其實仍然不是完全動態(tài)的數(shù)據(jù)結(jié)構(gòu)。GPU內(nèi)核函數(shù)通常無法直接訪問CPU內(nèi)存，因此在處理CPU內(nèi)存之前必須將數(shù)據(jù)復(fù)制到GPU上，然后再寫回CPU。但是，將數(shù)據(jù)復(fù)制到GPU或從GPU復(fù)制數(shù)據(jù)需要付出非常昂貴的時間代價，文獻[2]正是采用了這種昂貴的方式為實驗中的所有插入操作都預(yù)先分配了結(jié)點資源（必須在編譯時知道具體分配計劃），并且不能在運行時動態(tài)分配新項和釋放已刪除項。這是GPU上鏈式哈希表的一個最大的限制。而GLHT就沒有這樣的限制，因此更具靈活性。

鏈式哈希表必須為每個插入結(jié)點分配相應(yīng)的內(nèi)存，但幸運的是開放尋址的哈希表可以避免大量的內(nèi)存分配，GLHT中作為結(jié)構(gòu)基礎(chǔ)的跳步哈希表正是開放尋找哈希表的一個典型。雖然表面上看GLHT比GPU chained without allocation性能差，但在真實生產(chǎn)環(huán)境中更關(guān)心的是程序的總體運行時間，也就是GPU chained with allocation的運行性能，因此可以毫不猶豫地說，GLHT更具有競爭優(yōu)勢，畢竟它相對GPU chained with allocation有200～400倍的性能比。退一步說，即使不考慮GPU chained with allocation，GLHT也已經(jīng)在寫操作比重較大的工作負載中超越了GPU chained without allocation。

6 結(jié)束語

跳步哈希表可以使用高效的GPU合并訪問完成讀取請求，相對其他哈希表，更適合用于GPU設(shè)計，本文提出并實現(xiàn)了一種GPU跳步哈希表GLHT，它是首個GPU完全并發(fā)且可動態(tài)更新的跳步哈希表。GLHT與之前的工作相比，具有以下兩個特點：（1）warp內(nèi)單個操作并行，采用warp協(xié)同工作共享策略，減少程序控制流中的分支與發(fā)散；（2）warp間多個操作并發(fā)，使用全局內(nèi)存配合CUDA原子操作以及特殊的并發(fā)控制策略設(shè)計，在實現(xiàn)完全并發(fā)和無鎖特性的同時保證了讀操作的無等待特性。GLHT與其他GPU靜態(tài)哈希表相比，具有可以接受的構(gòu)建和檢索速度；與現(xiàn)有的CPU跳步哈希相比，具有4～9倍的性能優(yōu)勢；比采取預(yù)先分配內(nèi)存的GPU無鎖鏈式哈希表[1]更加靈活，并且在寫操作較多的工作負載中獲得了更好的性能。

本文實現(xiàn)的GLHT中，為了模型設(shè)計和說明的簡便，直接以unsigned long long int作為數(shù)據(jù)結(jié)構(gòu)的項，未來可以將鍵值存儲部分改為指向鍵-值對的指針以提高使用性。另外，由于目前GPU原子操作的限制（例如atomicCAS操作只涉及整數(shù)數(shù)據(jù)類型），GLHT的設(shè)計模型仍顯粗糙，未來可以等GPU原子操作可以涉及結(jié)構(gòu)對象時，繼續(xù)豐富本模型。