GPU版基數排序（radix sort）

基數排序演算法的思想很簡單，但如何提高並行性，使得能用gpu進行高效的排序計算呢？本文介紹適合於gpu計算版本的排序演算法。

為了說明gpu版的排序演算法，我們先複習一下最初版本的排序演算法：

a picture is worth than a thousand words.

因此我們上圖說話。

基數排序是從最低位看起，把這一位數字值相同的按照掃瞄順序放入同乙個桶裡面，值小的桶在前面。當所有數字都掃瞄完，再使用高一位，迴圈上述步驟，直至達到所有數字所有的最高位數，最後輸出的就是排序後的答案。

input:unsigned int inputvals[n]

output:unsigned int outputvals[n]

for(i=0; i
< numbins; i+= numbits)
1. 計算inputval中第i位值為0到numbins-1的各有多少個，其值分別放入binhistogram[numbins]中
2. 對binhistogram[numbins]進行scan計算，其值分別放入binscan[numbins]中，即
binscan[k]=binhistogram[0]+···+binhistogram[k-1]
3. 遍歷inputval陣列，計算第j個量時，計算其值所屬的桶bin的值
outputvals[binscan[bin]] = vals_src[j];
binscan[bin]++;   
4. swap(inputvals,outputvals)

gpu版本的排序演算法與cpu在前兩步都是類似的。

for(i=0; i
< numbins; i+= numbits)
1. 同cpu
2. 同cpu
3. 決定每個數值的相對位置，比如：[001
1001]
-> [010
1232]
4. 根據第二步和第三步計算的結果，決定每個每個數字的絕對位置，通過map運算從相應的位置取值

終於到了我最想提醒大家的部分了。

1. 如何決定每個數值的相對位置？

決定每個數值的相對位置其實使用的是compact方法，以input為[0 0 1 1 0 0 1]為例子，計算如下：

其中scan使用hill/steele演算法。

2. 但用以上的方法能算出真正的相對位置麼？

不能。

what？不能你還寫這麼久？

其實也不是完全不能，當排序的數目小於block中最大threads數目時（比如我設定的是2048），是可以正確排序的。為什麼呢？因為我們使用的scan方法中需要同步。

先複習hill/steele演算法，以加法為例，對[1, 2, 3, 4, 5, 6, 7, 8]進行scan計算

平行計算中最怕的就是沒能同步啦，那麼再上hill/steele演算法的**。

// scan
__global__ void scan_cal(unsigned
int *d_cdf, const
unsigned
int *d_bins, const size_t size)
// exclusive scan
d_cdf[myid]= d_bins[myid]-current_value;
__syncthreads();
}

發現問題了麼？在圖中每一行計算完後，需要等整個一行同步後才能進行下一次的計算。而__syncthreads();函式是block-level的同步，即在不同的block中的計算並不會同步後再進行，因而一些block沒有算完的時候，其他block已經取了之前的值接著下面的計算了。因此如果只有乙個block，即排序的數目小於block中最大threads數目時，還是可以正確排序的。

3. 那，怎麼辦呢？

這裡提供乙個可以排序220的方法。我們對之前的gpu演算法進行小的修改：

blocksize是乙個block裡面有多少個執行緒（blockdim.x×blockdim.y×blockdim.z），對於cuda計算能力1.x的，blocksize最大值為1024；1.x+的值為2048。

for(i=0; i
< numbins; i+= numbits)
1. 計算inputval的每blocksize個sub-array中，第i位值為0到numbins-1的各有多少個，其值分別放入binhistogram[numbins×griddim]中
2. 對binhistogram[numbins×griddim]進行scan計算，其值分別放入binscan[numbins×griddim]中，即
binscan[k]=binhistogram[0]+···+binhistogram[k-1]
3. block內決定每個數值的相對位置，比如：[001
1001]
-> [010
1232]
4. （不變）根據第二步和第三步計算的結果，決定每個每個數字的絕對位置，通過map運算從相應的位置取值

即分組compact。還是以[0 0 1 1 0 0 1]為例，此時我們設定blocksize=4，numbins=2（0和1），griddim=ceil(7/blocksize)=2：

binhistogram=[ 2 2 | 2 1]

第二個2代表第二組中值為0的從位址2開始

第三個4代表第一組中值為1的從位址4開始

第四個6代表第二組中值為1的從位址6開始

計算值為0的相對位置

計算針對0的predicate: [true true false false | true true false]

其中true=1, false=0: [1 1 0 0 | 1 1 0]

exclusive scan: [0 1 2 2 | 0 1 2]

對於predicate為true的值，留下上一步的結果[0 1 nan nan | 0 1 nan]

計算值為1的相對位置

計算針對1的predicate:[false false true true | false false true]

其中true=1, false=0: [0 0 1 1 | 0 0 1]

exclusive scan: [0 0 0 1 | 0 0 0]

對於predicate為true的值，留下上一步的結果[0 1 0 1 | 0 1 1]

絕對位置=相對位置+binscan[bin的位置]

絕對位置=[ (0+0) (1+0) (0+4) (1+4) | (0+2) (1+2) (1+6) ]=[0 1 4 5 | 2 3 7]

__global__ void specific_address(unsigned
int *d_out_value, const
unsigned
int *const d_in_value,
const
unsigned
int *d_cdfs, const
unsigned
int *d_relative_pos,
const size_t numbins, const
unsigned
int pos, const size_t size)

之前提到的，上述演算法對最多220數字進行排序（對於計算能力1.x以上的顯示卡可以為222）。因為還是scan同步那裡進行了限制，因而還有修改的空間。因為我所做的工作計算量小於220因此就沒接著繼續研究了，有興趣的同學歡迎分享對更高資料集排序實現的idea，我們一起學習~

GPU版基數排序（radix sort）

排序基數排序

排序基數排序

桶排序基數排序計數基數排序 Java

GPU版基數排序（radix sort）

排序 基數排序

排序 基數排序

桶排序 基數排序 計數基數排序 Java

相關推薦

排序基數排序

排序基數排序

桶排序基數排序計數基數排序 Java