張 娟，孫建伶

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

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

1 引言

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

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

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

GLHT的基礎(chǔ)數(shù)據(jù)結(jié)構(gòu)是跳步哈希表[4]。跳步哈希表的插入操作保持?jǐn)?shù)據(jù)的緊湊。當(dāng)發(fā)生數(shù)據(jù)沖突時(shí)，新數(shù)據(jù)插入到哈希槽（哈希槽即指鍵原始應(yīng)該被哈希到的槽）隨后的H個(gè)槽，這H個(gè)槽稱為當(dāng)前槽的鄰域，H是用戶設(shè)置的常數(shù)。每個(gè)槽關(guān)聯(lián)一個(gè)由H+1個(gè)bit組成的bitmap，指示當(dāng)前槽和后續(xù)H個(gè)槽中的項(xiàng)是否是最初哈希到當(dāng)前槽的項(xiàng)。若某個(gè)槽的項(xiàng)本來應(yīng)該哈希到前面的槽，則稱這個(gè)槽“從屬”于前面的那個(gè)槽。圖1是鍵v插入跳步哈希表的過程，白色表示空槽，灰色表示槽中有項(xiàng)，該哈希表的H為3。鍵v本應(yīng)哈希到槽6，但是發(fā)生了數(shù)據(jù)沖突。于是，首先通過線性探測(cè)找到距槽6最近的空槽13。如果兩個(gè)槽的距離小于等于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ù)與原始哈希槽的距離不會(huì)大于H，因此查找時(shí)只需檢查哈希槽及其鄰域中是否有目標(biāo)鍵，若無則可確定目標(biāo)鍵不存在，由此保證任何情況下的查找時(shí)間都是O(1)。

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

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

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

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

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

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

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

2 相關(guān)工作

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

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

最近，Ashkiani等人[3]設(shè)計(jì)了一種完全并發(fā)的GPU動(dòng)態(tài)無鎖鏈?zhǔn)焦１怼猄lab Hash。他們認(rèn)為，GFSL無論在插入、刪除還是查找操作中，都無法擊敗Slab Hash的性能峰值。

3 設(shè)計(jì)

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

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

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

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

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

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

如圖4，雖然GLHT讓warp大小的整個(gè)線程塊內(nèi)的線程協(xié)同地處理同一個(gè)操作任務(wù)，但不同線程塊之間仍然是操作獨(dú)立且完全并發(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)存的哪個(gè)層次。GPU的內(nèi)存結(jié)構(gòu)分為三個(gè)層次：可以被設(shè)備內(nèi)所有線程訪問的大的全局內(nèi)存；每個(gè)線程塊有著的更小但更快的共享內(nèi)存；線程塊中每個(gè)線程的本地寄存器。共享內(nèi)存很?。ㄍǔ?6 KB），并且它進(jìn)行了分區(qū)，因此來自不同塊的線程無法訪問另一個(gè)塊的共享內(nèi)存。GPU的全局內(nèi)存容量大，可供所有線程訪問。由于數(shù)以百萬計(jì)的線程可以執(zhí)行GPU內(nèi)核函數(shù)，但只有有限數(shù)量的SM存在，因此線程塊需要排隊(duì)等待SM。因此，除了內(nèi)核函數(shù)結(jié)束的時(shí)候，并沒有辦法可以全局地同步所有線程。為了實(shí)現(xiàn)warp間操作的完全并發(fā)，GLHT通過全局內(nèi)存實(shí)現(xiàn)各線程對(duì)所有數(shù)據(jù)狀態(tài)的共享。

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

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

4 實(shí)現(xiàn)

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

4.1 拆分操作階段

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

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

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階段將目標(biāo)鍵通過atomic-CAS放進(jìn)hash_pos，而刪除操作的update階段將target通過atomicCAS置為空，update階段的槽角色如圖6，注意target可能與hash_pos重合。

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

find_closer_empty階段的目標(biāo)是將找到的空槽向前移動(dòng)一次，find_closer_empty階段循環(huán)多個(gè)swap_value_into_empty階段，直到移動(dòng)成功。swap_value_into_empty階段每次對(duì)一塊置換區(qū)域操作，置換區(qū)域的第一個(gè)槽以角色swap_head表示，置換區(qū)域的最后一個(gè)槽即前面找到的那個(gè)空槽。從前到后在置換區(qū)域中尋找一個(gè)“從屬”于swap_head的槽，將這個(gè)槽以角色swap表示，并置換target和swap的項(xiàng)，置換完成則find_closer_empty階段也完成了；但若沒有找到swap，則find_closer_empty階段將角色swap_head向后推動(dòng)一個(gè)位置，并循環(huán)swap_value_into_empty階段。find_closer_empty階段最初將swap_head定在target前的第H個(gè)位置。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 鎖標(biāo)記

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

multiple_lock的含義如下：

（1）當(dāng)其標(biāo)記在非空槽時(shí)，表示該槽正處于插入或刪除操作的update階段的hash_pos角色。

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

swap_lock的含義如下：

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

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

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

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

Fig.8 Operations on lockflag during insert operation圖8 插入操作過程中對(duì)鎖標(biāo)記的操作

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

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

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

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

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

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鎖標(biāo)記，EMP_FLAG_MASK用來判斷是否為空槽。

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

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

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

4.3.2 插入操作

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

除了shuffle指令，插入操作還會(huì)在以下情況使用ballot和ffs指令組合：（1）在find_empty階段，找到最靠近hash_pos的不帶multiple_lock鎖標(biāo)記的空槽；（2）在swap_value_into_empty階段，找到不帶任何鎖標(biāo)記的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)存慢幾個(gè)數(shù)量級(jí)。為了提升性能，GLHT在此基礎(chǔ)上設(shè)計(jì)了特殊的并發(fā)控制策略（包括4.2節(jié)描述的鎖標(biāo)記和暫時(shí)重復(fù)策略兩方面），保證了讀操作的無等待特性，在一定程度上彌補(bǔ)了全局內(nèi)存訪問慢的缺點(diǎn)。

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

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

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

5 實(shí)驗(yàn)評(píng)估

本實(shí)驗(yàn)全部在Intel Xeon E5-2620服務(wù)器上執(zhí)行，該服務(wù)器擁有1個(gè)Socket，每個(gè)Socket有6個(gè)核，每個(gè)核有2個(gè)超線程。內(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上進(jìn)行評(píng)估比較的，GDDR5X容量為8 GB。CUDA代碼采用CUDA 8.0編譯器（V8.0.61）編譯。

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

5.1 靜態(tài)基準(zhǔn)

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

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

圖9是各數(shù)據(jù)結(jié)構(gòu)的構(gòu)建速度比較。GLHT雖然比線性探測(cè)和平方探測(cè)靜態(tài)哈希表慢，但作為動(dòng)態(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 動(dòng)態(tài)并發(fā)基準(zhǔn)

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

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

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

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

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

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

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

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

至于文獻(xiàn)[2]提供的GPU上的無鎖鏈?zhǔn)焦１?，雖然它也支持并發(fā)的插入、刪除和查找操作，但其實(shí)仍然不是完全動(dòng)態(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ù)需要付出非常昂貴的時(shí)間代價(jià)，文獻(xiàn)[2]正是采用了這種昂貴的方式為實(shí)驗(yàn)中的所有插入操作都預(yù)先分配了結(jié)點(diǎn)資源（必須在編譯時(shí)知道具體分配計(jì)劃），并且不能在運(yùn)行時(shí)動(dòng)態(tài)分配新項(xiàng)和釋放已刪除項(xiàng)。這是GPU上鏈?zhǔn)焦１淼囊粋€(gè)最大的限制。而GLHT就沒有這樣的限制，因此更具靈活性。

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

6 結(jié)束語(yǔ)

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

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