一種基于GPU的枚舉排序算法及其并行化

谷國太，孫陸鵬，張紅艷，肖漢

(1.河南省新聞出版學校，河南 鄭州 450044；2.鄭州師范學院 信息科學與技術(shù)學院，河南 鄭州 450044)

0 引 言

在數(shù)據(jù)處理的過程中，排序是其中非常重要的一個環(huán)節(jié)，而且排序在各種系統(tǒng)或應用軟件中都有著非常廣泛的應用，在處理數(shù)據(jù)序列的各種應用中也發(fā)揮著重要的作用[1-2]。枚舉排序是通過把目標元素與其他元素進行比較，最終確定其在整個元素序列中的位置。隨著社會信息化的快速發(fā)展，無論在科學研究、工業(yè)生產(chǎn)、娛樂，還是在社會民生、環(huán)境與商業(yè)等領(lǐng)域，數(shù)據(jù)量都呈現(xiàn)爆炸式增長，這就對計算速度提出了更高的要求，提高枚舉排序的運算速度有著非常重要的意義[3-4]。

傳統(tǒng)的枚舉并行排序算法是在CPU上運行，雖然多核CPU可以提高運算速度，但是受到空間、電力、冷卻等因素限制，多核系統(tǒng)目前面臨著核數(shù)超過16個以后性能無法隨內(nèi)核數(shù)線性擴展以及并行軟件限制等問題[5]。在基于圖形處理器(graphic processing unit,GPU)的并行計算中，GPU解決了上述多核CPU在運算上的不足，同時降低了成本和功耗，為解決越來越復雜的圖形和數(shù)據(jù)計算問題提供了新的方案?？梢灶A見，在今后的發(fā)展中，GPU處理器的地位越來越重要，并將會成為提高運算速度的重要突破口[6-8]。

國內(nèi)外學者對枚舉算法進行了大量的研究。He G H等[9]提出一種軟輸入軟輸出度量優(yōu)先MIMO檢測的算法和結(jié)構(gòu)。在該算法中，混合枚舉策略避免了全部枚舉和排序，大大降低了計算復雜度；J.Davila等[10]提出一種算法，該算法解決了稀疏枚舉排序問題，比以往的算法性能更好；孫斌[11]利用MPI實現(xiàn)了并行枚舉算法，將數(shù)據(jù)序列分成若干個小的子序列，分派給多個處理器同時進行處理排序，之后再將每個處理器的排序結(jié)果送往主節(jié)點進行整合并完成最終排序，實現(xiàn)了多進程并行計算，大大降低了運算時間；S.Rajasekaran[12]提出一種合并排序(LMM)算法，證明在超立方體上可以得到比Nassimi和Sahni的經(jīng)典算法更簡單的稀疏枚舉排序；王勇超[13]利用High Performance Linpack基準測試方法對集群系統(tǒng)的運算速度和性能進行了研究和測試，進而深入探討高性能計算的枚舉排序算法；H.M.Alnuweiri[14]提出一種新的VLSI最優(yōu)排序器設(shè)計方法，它將輪換排序和枚舉排序相結(jié)合，對N個數(shù)進行排序，降低了算法時間復雜度；A.E.Perez等[15]公開了一種符號枚舉排序過程的方法及使用該方法的裝置，允許準確和有效地計算度量以及確定低復雜度的星座符號的最佳搜索順序；羅明山[16]在OpenMP編程模型上實現(xiàn)了枚舉并行排序，但是，多核處理器需要內(nèi)核之間進行頻繁的信息交換，實現(xiàn)核與核之間的同步需要大量的指令處理時間，這就需要在降低單個內(nèi)核最高頻率的基礎(chǔ)上增加內(nèi)核的數(shù)量。核數(shù)與主頻相互牽制，并不能很好地解決問題。

目前國內(nèi)外利用CUDA進行枚舉排序并行算法的研究較少，大部分研究基于CPU傳統(tǒng)的MPI或者OpenMP并行計算模型，雖然可以進行并行計算，但是加速效果不理想。為此，本文將GPU并行計算和枚舉排序算法相結(jié)合，可以大大提高算法的運算能力和排序效率，著重研究如何將枚舉排序算法在主從式異構(gòu)計算模型下，利用GPU并行完成每個數(shù)據(jù)與其他數(shù)據(jù)大小比較以及統(tǒng)計記錄值的計算，從而驗證該結(jié)構(gòu)下枚舉排序并行算法對于大規(guī)模數(shù)據(jù)集數(shù)據(jù)處理的有效性。

1 CUDA體系結(jié)構(gòu)

GPU由數(shù)量眾多的邏輯運算部件ALU構(gòu)成，這些邏輯運算部件又被劃分為若干個流多處理器，而每個流多處理器中又包含若干個流處理器[17]。流多處理器是GPU結(jié)構(gòu)中執(zhí)行和調(diào)度的單元，包含了各自獨立的控制單元，可以獨立控制指令的分發(fā)和執(zhí)行，流處理器是GPU最基本的計算單元[18]。

CUDA編程模型采用CPU與GPU協(xié)同工作、各司其職的工作模式，如圖1所示。CPU的主要任務(wù)是邏輯分析運算和串行運算，GPU的任務(wù)主要是并行計算，目的是把并行計算任務(wù)細化和線程化。CPU與GPU都有自己相應的存儲空間，且各自獨立。kernel函數(shù)不是一個完整的任務(wù)，而是整個任務(wù)中可以被并行執(zhí)行的步驟[19-20]。

圖1 CUDA編程模型

2 枚舉排序算法描述

枚舉排序是一種簡單排序算法。算法主要思想為：對每一個待排序的元素統(tǒng)計小于它的所有元素的個數(shù)，從而得到該元素最終處于序列中的位置。假定待排序的n個數(shù)存在a[1],…，a[n]中[21]，首先將a[1]與a[2],…，a[n]比較，記錄比其小的數(shù)的個數(shù)，令其為k，a[1] ，存入有序的數(shù)組b[1],…，b[n]的b[k+1]位置上；其次將a[2]與a[1],a[3],…，a[n]比較，記錄比其小的數(shù)的個數(shù)，依此類推[22]。

算法：枚舉排序串行算法

輸入：無序數(shù)組a[1],…,a[n]

輸出：有序數(shù)組b[1],…，b[n]

Begin

for i=1 to n do

(1)k=1

(2)for j=1 to n do

if a[i]>a[j] then

k=k+1

end if

end for

(3)b[k]=a[i]

end for

End

串行的枚舉排序算法描述如下。

步驟一隨機產(chǎn)生一個數(shù)組序列a，包含n個待排序數(shù)據(jù)。

步驟二對于待排序數(shù)組元素a[i]，從j=1開始依次比較a[i]和a[j],并記錄比a[i]小的數(shù)的個數(shù)，記錄在數(shù)組b[k+1]位置上[23]。

步驟三若i小于數(shù)組的個數(shù)n，則循環(huán)執(zhí)行上述步驟二到步驟三，直至待排序數(shù)組a中的數(shù)據(jù)全部排序,否則執(zhí)行步驟四。

步驟四完成排序工作[24]。

串行秩排序算法中比較操作共n*(n-1)次，時間復雜度為O(n2)。

3 GPU加速的枚舉排序并行算法的分析與設(shè)計

3.1 枚舉排序并行算法設(shè)計

在算法并行特征分析的基礎(chǔ)上可以看出，每個待排序數(shù)的排序都是獨立排序。結(jié)合GPU高效并行的特點，假設(shè)對一個長為n的輸入序列使用n個線程進行排序，只需分配每個線程負責完成對其中一個待排序元素在有序序列中的定位和存儲，使整個待排序序列的排序并行實現(xiàn)，接著將所有存儲的信息集中到主機中即可。CUDA枚舉排序并行算法描述如下。

算法：枚舉排序并行算法

輸入：無序數(shù)組a[1],…，a[n]

輸出：有序數(shù)組a[1],…，b[n]

Begin

(1)host send a[1],…,a[n]to device

(2)for all Tiwhere 1≤i≤n par-do

(2.1)k=1

(2.2)for j=1 to n do

if(a[i]≥a[j])and(i>j)then

k=k+1

end if

end for

b[k]=a[k]

end for

(3)device send b[1],…，b[n] to host

End

在枚舉排序并行算法中使用了n個線程，由于每個線程定位一個元素，所以步驟(2)的時間復雜度為O(1)；步驟(3)中設(shè)備完成的數(shù)組元素重定位操作的時間復雜度為O(n)；所以總的時間復雜度為O(n)。由此可見，枚舉排序并行算法時間復雜度O(n)

3.2 枚舉排序算法并行化方案

(1)為待排序序列a和有序序列b分配設(shè)備存儲器空間。

cudaMalloc((void**)& a,(n+5)*sizeof(double))

cudaMalloc((void**)& b,(n+5)*sizeof (double))

(2)把主機端的數(shù)據(jù)傳遞到設(shè)備端。

cudaMemcpy(a,A,(n+5)*sizdof (double),cudaMemcpyHostToDevice)

(3)定義kernel配置。

線程塊分配的線程數(shù)：

dim3 threads(256)

網(wǎng)格分配的線程塊數(shù)：

dim3 blocks((n+threads.x-1)/threads.x)

通過實驗尋找最合適的線程分配數(shù)，如表1所示。綜合比較可以看出，當線程為128或256時，運算速度較快。

(4)發(fā)射kernel進行并行計算。

enume<<>>(a,b,n)

(5)將已排序數(shù)據(jù)從設(shè)備端傳輸?shù)街鳈C端進行輸出。

cudaMemcpy(B,b,(n+5)*sizeof(double),cudaMemcpyDeviceToHost)

表1 不同線程數(shù)量運行時間對比

4 實驗與分析

4.1 實驗運算平臺

對枚舉排序并行算法的計算效率進行測試，測試平臺如下。

CPU為Intel Core i5，主頻為2.50 GHz，內(nèi)存為4 GB。GPU為NVIDIA GTX 580，GPU型號為GF119，屬于Fermi構(gòu)架。內(nèi)置512個處理器核心，最高處理頻率1 544 MHz。操作系統(tǒng)為Microsoft Window7 64位，仿真實驗軟件為MATLAB2013b，GPU的應用程序編程接口為CUDA7.0，開發(fā)環(huán)境為 Microsoft Visual Studio 2015。

4.2 實驗結(jié)果和性能分析

4.2.1 實驗數(shù)據(jù)

步驟一排序數(shù)據(jù)規(guī)模為100～40 000時，對基于CPU的枚舉排序串行算法的運算時間計時。

步驟二排序數(shù)據(jù)規(guī)模為100～40 000時，對基于GPU的枚舉排序并行算法的運算時間計時。

步驟三基于CPU的枚舉排序串行實現(xiàn)和基于GPU的枚舉排序并行實現(xiàn)運算時間如表2所示。

4.2.2 加速性能分析

由圖2可以看出，當數(shù)據(jù)規(guī)模達到20 000時，基于GPU的枚舉排序相比基于CPU的枚舉排序才呈現(xiàn)加速效果，當數(shù)組的數(shù)值達到40 000時，基于GPU的枚舉排序時間是基于CPU的枚舉排序時間的1/4左右。

在枚舉并行排序的過程中，如果排序數(shù)據(jù)規(guī)模比較小，CPU與GPU相比，在運算速度上有很大的優(yōu)勢。這是因為在GPU計算的過程中，數(shù)據(jù)在內(nèi)存和全局存儲器之間來回傳輸，數(shù)據(jù)交互帶來的時間延遲是不能被忽略的，消耗了很多時間。

表2 枚舉排序串并行性能對比

圖2 不同數(shù)據(jù)規(guī)模的枚舉排序GPU加速效率趨勢圖

通過實驗數(shù)據(jù)可以看出，如果數(shù)據(jù)規(guī)模太小，使用CUDA并不能帶來處理速度提升。因為，對于一些數(shù)據(jù)規(guī)模非常小的系統(tǒng)而言，串行算法運算時間非常短，甚至能夠在ms級內(nèi)完成。但是在使用基于CUDA計算時，由于GPU系統(tǒng)存在啟動時間問題，在GPU上的執(zhí)行時間相對就較長，而隨著數(shù)據(jù)規(guī)模的不斷增大，GPU運算時間增加的非常慢，而CPU運算時間則呈現(xiàn)迅速增長趨勢，GPU強大的浮點運算能力就顯現(xiàn)出來了。

4.2.3 系統(tǒng)性能瓶頸分析

在基于GPU的枚舉排序并行算法中需要對存儲器進行讀寫操作。對于n個待排序數(shù)據(jù)，需要讀取n次存儲器，寫入n次存儲器。設(shè)待排序數(shù)據(jù)集合數(shù)據(jù)為20 000，每個數(shù)據(jù)值分配存儲空間大小是4字節(jié)。所以，存儲器存取數(shù)據(jù)總量約為1.6 GB。除以kernel實際執(zhí)行時間0.013 s，得到的帶寬數(shù)值約為123.08 GB/s，這已經(jīng)接近NVIDIA GTX 580顯示存儲器的192.38 GB/s帶寬了。因而，分析可得，基于GPU的枚舉排序并行算法的效率受限于全局存儲器帶寬。

5 結(jié) 論

(1)提出利用GPU并行運算提升大規(guī)模數(shù)據(jù)處理速度的方法是可行，且效果顯著。

(2)通過枚舉排序并行算法，在不同數(shù)據(jù)規(guī)模下，尋找每個線程塊分配的最佳線程數(shù)，開展枚舉排序并行算法分析，確定設(shè)計方案。

(3)通過枚舉排序串并行算法實驗分析發(fā)現(xiàn)，當數(shù)據(jù)規(guī)模達到40 000時，并行枚舉排序算法時間是串行算法運行時間的1/4左右，隨著數(shù)據(jù)規(guī)模的不斷增大，加速比會越來越大。

(4)本文所提枚舉排序并行算法是一重并行，將來研究中會繼續(xù)優(yōu)化并行算法，使枚舉排序算法充分利用GPU的高效并行性。同時，目前主流的異構(gòu)眾核協(xié)處理器除了GPU以外，還有Intel的基于MIC架構(gòu)眾核協(xié)處理器Xeon Phi，輔助計算的能力也非常強大。下一步計劃完成基于CPU+MIC的枚舉排序并行算法。