當(dāng)前位置：首頁 > 编程资源 > 编程问答 >内容正文

编程问答

CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现

發(fā)布時間：2025/3/21 编程问答 46 豆豆

生活随笔收集整理的這篇文章主要介紹了 CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现小編覺得挺不錯的,現(xiàn)在分享給大家,幫大家做個參考.

在第五講中我們學(xué)習(xí)了GPU三個重要的基礎(chǔ)并行算法： Reduce, Scan 和 Histogram，分析了其作用與串并行實(shí)現(xiàn)方法。在第六講中，本文以冒泡排序 Bubble Sort、歸并排序 Merge Sort 和排序網(wǎng)絡(luò)中的雙調(diào)排序 Bitonic Sort?為例，講解如何從數(shù)據(jù)結(jié)構(gòu)課上學(xué)的串行并行排序方法轉(zhuǎn)換到并行排序，并附GPU實(shí)現(xiàn)代碼。

在并行方法中，我們將考慮到并行方法需要注意的特點(diǎn)進(jìn)行設(shè)計，希望大家在讀完本文后對GPU上并行算法的設(shè)計有一些粗淺的認(rèn)識。需注意的特點(diǎn)如下：?
1. 充分發(fā)揮硬件能力（盡量不要有空閑且一直處于等待狀態(tài)的SM）?
2. 限制branch divergence（見CUDA系列學(xué)習(xí)（二））?
3. 盡量保證內(nèi)存集中訪問（即防止不命中）

( 而我們在數(shù)據(jù)結(jié)構(gòu)課上學(xué)習(xí)的sort算法往往不注意這幾點(diǎn)。)

CUDA系列學(xué)習(xí)目錄：

CUDA系列學(xué)習(xí)（一）An Introduction to GPU and CUDA

CUDA系列學(xué)習(xí)（二）CUDA memory & variables - different memory and variable types

CUDA系列學(xué)習(xí)（三）GPU設(shè)計與結(jié)構(gòu)QA & coding練習(xí)

CUDA系列學(xué)習(xí)（四）Parallel Task類型與 Memory Allocation

CUDA系列學(xué)習(xí)（五）GPU基礎(chǔ)算法: Reduce, Scan, Histogram

I. Bubble Sort

冒泡排序，相信大家再熟悉不過了。經(jīng)典冒泡排序通過n輪有序冒泡(n為待排序的數(shù)組長度)得到有序序列，其空間復(fù)雜度O(1), 時間復(fù)雜度O(n^2)。

那么如何將冒泡排序算法改成并行算法呢？這里就需要解除一些依賴關(guān)系，比如是否能解除n輪冒泡間的串行依賴 & 是否能解除每一輪冒泡內(nèi)部的串行依賴，使得同樣的n^2次冒泡操作可以通過并行，降低step complexity。?
1996年， J Kornerup針對這些問題提出了odd-even sort算法，并在論文中證明了其排序正確性。?

I.1 從Bubble Sort到Odd-even Sort

先來看一下odd-even sort的排序方法:?

?
圖1.1

上圖為odd-even sort的基本方法。?
奇數(shù)步中， array中奇數(shù)項(xiàng)array[i]與右邊的item(array[i + 1])比較；?
偶數(shù)步中， array中奇數(shù)項(xiàng)array[i]與左邊的item(array[i - 1]) 比較；

這樣，同一個step中的各個相鄰比較就可以并行化了。

PS：對于array中有偶數(shù)個元素的數(shù)組也是一樣：?

?
圖1.2

I.2 Odd-even Sort復(fù)雜度

在odd-even sort的算法下，原本O(n^2)的總比較次數(shù)不變，但是由于并行，時間復(fù)雜度降到O(n), 即

step complexity = O(n)?
work complexity = O(n^2)

code詳見 <?Bubble sort and its variants?>

II. Merge Sort

看過odd-even sort后，我們來看如何將歸并排序并行化。數(shù)據(jù)結(jié)構(gòu)課上我們都學(xué)過經(jīng)典歸并排序：基于divide & conquer 思想，每次將一個array分拆成兩部分，分別排序后合并兩個有序序列。可以通過 T(n)=2T(n/2)+n 得到，其complexity = O(nlogn)。和I.1節(jié)類似，我們看看merge sort中的哪些步是可以并行的。

這里可以將基于merge sort的大規(guī)模數(shù)據(jù)排序分為三部分。經(jīng)過divide步之后，數(shù)據(jù)分布如圖所示：?

?
圖2.1

最下端的為?大量-短序列?的合并；?
中間一塊為中等數(shù)量-中等長度序列?的合并；?
最上端的為少量-長序列?的合并；

我們分這三部分進(jìn)行并行化。之后大家會明白為啥要這么分~

II.1 Step 1: Huge number of small tasks

這一部分，每一部分序列合并的代價很小，而有眾多這樣的任務(wù)。所以我們采取給每個merge一個thread去執(zhí)行， thread 內(nèi)部就用串行merge的方法。?

II.2 Step 2: Mid number of mid task

在這一階段，有中等數(shù)量的task，每個task的工作量也有一定增長。所以我們采取給每個merge一個SM去執(zhí)行，每個block上運(yùn)行的多個thread并行merge的方法。和step1中的主要區(qū)別就是block內(nèi)部merge從串行改成了并行。那么怎樣去做呢？

如下圖所示，假如現(xiàn)在有兩個長為4個元素的array，想對其進(jìn)行排序，將merge排序結(jié)果index 0 - 7 寫入數(shù)據(jù)下方方塊。

?
圖2.2

做法：?
對于array中的每個數(shù)字，看兩個位置：?
1. 自己所在序列的位置：看它前面有幾個元素?
2. 另一個序列的位置：另一個序列中有多少個元素比它小（采用二分搜索）

這樣做來，第一步O(1)可得到，第二步二分查找O（logn）可得到；整個merge過程用shared memory存結(jié)果。

II.3 Step 3: Small number of huge task

第三個環(huán)節(jié)中，也就是merge的頂端（最后一部分），每個merge任務(wù)的元素很多，但是merge任務(wù)數(shù)很少。這種情況下，如果采用Step2的方法，最壞的情況就是只有一個很大的task在跑，此時只有1個SM在忙，其他空閑SM卻沒法利用，所以這里我們嘗試將一個任務(wù)分給多個SM執(zhí)行。

方法：如圖2.3所示，將每個序列以256個元素為單位分段，得到兩個待merge序列In1和In2。然后，對這些端節(jié)點(diǎn)排序，如EABFCGDH。如step2中的方法，我們計算每個段節(jié)點(diǎn)在另一個短序列（長256）中的位置，然后只對中間那些不確定的部分進(jìn)行merge排序，每個merge分配一個SM。

如E~A的部分，?
1. 計算出E在In1中的位置posE1, A在In2中的位置posA2?
2. merge In1中 posE1~A 和 In2中 E~posA2的元素

?
圖2.3

II.4 Merge sort in GPU

以上面step 1的merge sort為例，其gpu代碼中kernel函數(shù)如下：

其中temp為排好序的序列，每次排序兩個大小為sortedsize的block，為temp賦值2 * sortedsize個元素。

所以實(shí)際上，sortedsize就是一個sorted block的大小。

<code class="hljs objectivec has-numbering" style="display: block; padding: 0px; color: inherit; box-sizing: border-box; font-family: 'Source Code Pro', monospace;font-size:undefined; white-space: pre; border-radius: 0px; word-wrap: normal; background: transparent;">__global__ void mergeBlocks(int *a, int *temp, int sortedsize) {int id = blockIdx.x;int index1 = id * 2 * sortedsize;int endIndex1 = index1 + sortedsize;int index2 = endIndex1;int endIndex2 = index2 + sortedsize;int targetIndex = id * 2 * sortedsize;int done = 0;while (!done){if ((index1 == endIndex1) && (index2 < endIndex2))temp[targetIndex++] = a[index2++];else if ((index2 == endIndex2) && (index1 < endIndex1))temp[targetIndex++] = a[index1++];else if (a[index1] < a[index2])temp[targetIndex++] = a[index1++];elsetemp[targetIndex++] = a[index2++];if ((index1 == endIndex1) && (index2 == endIndex2))done = 1;} } </code><ul class="pre-numbering" style="box-sizing: border-box; position: absolute; width: 50px; top: 0px; left: 0px; margin: 0px; padding: 6px 0px 40px; border-right-width: 1px; border-right-style: solid; border-right-color: rgb(221, 221, 221); list-style: none; text-align: right; background-color: rgb(238, 238, 238);"><li style="box-sizing: border-box; padding: 0px 5px;">1</li><li style="box-sizing: border-box; padding: 0px 5px;">2</li><li style="box-sizing: border-box; padding: 0px 5px;">3</li><li style="box-sizing: border-box; padding: 0px 5px;">4</li><li style="box-sizing: border-box; padding: 0px 5px;">5</li><li style="box-sizing: border-box; padding: 0px 5px;">6</li><li style="box-sizing: border-box; padding: 0px 5px;">7</li><li style="box-sizing: border-box; padding: 0px 5px;">8</li><li style="box-sizing: border-box; padding: 0px 5px;">9</li><li style="box-sizing: border-box; padding: 0px 5px;">10</li><li style="box-sizing: border-box; padding: 0px 5px;">11</li><li style="box-sizing: border-box; padding: 0px 5px;">12</li><li style="box-sizing: border-box; padding: 0px 5px;">13</li><li style="box-sizing: border-box; padding: 0px 5px;">14</li><li style="box-sizing: border-box; padding: 0px 5px;">15</li><li style="box-sizing: border-box; padding: 0px 5px;">16</li><li style="box-sizing: border-box; padding: 0px 5px;">17</li><li style="box-sizing: border-box; padding: 0px 5px;">18</li><li style="box-sizing: border-box; padding: 0px 5px;">19</li><li style="box-sizing: border-box; padding: 0px 5px;">20</li><li style="box-sizing: border-box; padding: 0px 5px;">21</li><li style="box-sizing: border-box; padding: 0px 5px;">22</li><li style="box-sizing: border-box; padding: 0px 5px;">23</li><li style="box-sizing: border-box; padding: 0px 5px;">24</li><li style="box-sizing: border-box; padding: 0px 5px;">25</li></ul>

主函數(shù)中，定義block大小并調(diào)用kernel function：

<code class="hljs cs has-numbering" style="display: block; padding: 0px; color: inherit; box-sizing: border-box; font-family: 'Source Code Pro', monospace;font-size:undefined; white-space: pre; border-radius: 0px; word-wrap: normal; background: transparent;"> int blocks = BLOCKS/2;int sortedsize = THREADS;while (blocks > 0){mergeBlocks<<<blocks,1>>>(dev_a, dev_temp, sortedsize);cudaMemcpy(dev_a, dev_temp, N*sizeof(int), cudaMemcpyDeviceToDevice);blocks /= 2;sortedsize *= 2;}cudaMemcpy(a, dev_a, N*sizeof(int), cudaMemcpyDeviceToHost); </code><ul class="pre-numbering" style="box-sizing: border-box; position: absolute; width: 50px; top: 0px; left: 0px; margin: 0px; padding: 6px 0px 40px; border-right-width: 1px; border-right-style: solid; border-right-color: rgb(221, 221, 221); list-style: none; text-align: right; background-color: rgb(238, 238, 238);"><li style="box-sizing: border-box; padding: 0px 5px;">1</li><li style="box-sizing: border-box; padding: 0px 5px;">2</li><li style="box-sizing: border-box; padding: 0px 5px;">3</li><li style="box-sizing: border-box; padding: 0px 5px;">4</li><li style="box-sizing: border-box; padding: 0px 5px;">5</li><li style="box-sizing: border-box; padding: 0px 5px;">6</li><li style="box-sizing: border-box; padding: 0px 5px;">7</li><li style="box-sizing: border-box; padding: 0px 5px;">8</li><li style="box-sizing: border-box; padding: 0px 5px;">9</li><li style="box-sizing: border-box; padding: 0px 5px;">10</li><li style="box-sizing: border-box; padding: 0px 5px;">11</li></ul>

III. Bitonic Sort

III.1 Bitonic Sequence 雙調(diào)序列

不同于以上兩種排序方法，現(xiàn)在我們要接觸的雙調(diào)排序是排序網(wǎng)絡(luò)方法中的一種。想起當(dāng)年在浙大面試某導(dǎo)師的實(shí)驗(yàn)室時就是讓實(shí)現(xiàn)的雙調(diào)排序，并不斷優(yōu)化，不過當(dāng)時土得一坨，就沒聽說過這個算法。。。最后寫出個多線程就結(jié)束了，后來也沒再整理。現(xiàn)在我們來看看bitonic sort是個什么鬼。

雙調(diào)排序是排序網(wǎng)絡(luò)中最快的方法之一。所謂的排序網(wǎng)絡(luò)是data-independent的排序，即網(wǎng)絡(luò)比較順序與數(shù)據(jù)無關(guān)的排序方法，所以特別適合硬件做并行化。

在了解雙調(diào)排序算法之前，我們先來看看什么是雙調(diào)序列。雙調(diào)序列是一個先單調(diào)遞增后單調(diào)遞減或者先單調(diào)遞減后單調(diào)遞增的序列。

III.2 雙調(diào)排序算法

假如我們現(xiàn)在拿到了雙調(diào)序列，怎樣對它按照從小到大進(jìn)行排序呢？形象一點(diǎn)來看，我們將一個雙調(diào)序列切成兩半，每一段的單調(diào)性統(tǒng)一，然后如下圖圖3.1所示，將兩段疊放起來，進(jìn)行兩兩比較，這樣一定能夠在左右兩段分別得到一個雙調(diào)序列（想想為什么得到的是兩個雙調(diào)序列），且左邊的雙調(diào)序列中元素全部小于右側(cè)得到的雙調(diào)序列的所有元素。迭代這個過程，每次都能將序列二分成兩個子雙調(diào)序列，直到這個子雙調(diào)序列的長度為2，也就變成了一個單調(diào)子序列，這個過程排序后原先的長雙調(diào)序列就變?yōu)橛行蛄?。整個過程如下圖圖3.2所示。

?
圖3.1

?
圖3.2

III.3 任意序列生成雙調(diào)序列

好，III.2中講了怎樣對雙調(diào)序列進(jìn)行排序，那問題來了，怎樣從任意序列生成雙調(diào)序列呢？這里可以看看本文最后的參考文獻(xiàn)3，寫得很詳細(xì)。這個過程叫Bitonic merge, 實(shí)際上也是divide and conquer的思路。和III.2中的思路正相反，我們可以將兩個相鄰的，單調(diào)性相反的單調(diào)序列看作一個雙調(diào)序列，?每次將這兩個相鄰的，單調(diào)性相反的單調(diào)序列merge生成一個新的雙調(diào)序列，然后排序（同III.2）。這樣只要每次兩個相鄰長度為n的序列的單調(diào)性相反，就可以通過連接得到一個長度為2n的雙調(diào)序列。 n開始為1，每次翻倍，直到等于數(shù)組長度，就只需要一遍單方向（單調(diào)性）排序了。

?
圖3.3

以16個元素的array為例，?
1. 相鄰兩個元素合并形成8個單調(diào)性相反的單調(diào)序列，?
2. 兩兩序列合并，形成4個雙調(diào)序列，分別按相反單調(diào)性排序?
3. 4個長度為4的相反單調(diào)性單調(diào)序列，相鄰兩個合并，生成兩個長度為8的雙調(diào)序列，分別排序?
4. 2個長度為8的相反單調(diào)性單調(diào)序列，相鄰兩個合并，生成1個長度為16的雙調(diào)序列，排序

總算講完了，那么腫么實(shí)現(xiàn)呢？我們看這個過程需要控制哪些地方？如上圖所示，我們可以將len=16的array的雙調(diào)排序分成4部分，每一部分結(jié)束都會形成若干長度為?i?的單調(diào)序列。在每一部分中，用?j?表示比較的間隔，如下圖所示每一時刻的i和j。

?
圖3.4?

III.4 雙調(diào)排序的并行實(shí)現(xiàn)

本著“talk is cheap, show me the code”的優(yōu)良作風(fēng)，拿粗來雙調(diào)排序的GPU實(shí)現(xiàn)代碼share如下：

<code class="hljs cpp has-numbering" style="display: block; padding: 0px; color: inherit; box-sizing: border-box; font-family: 'Source Code Pro', monospace;font-size:undefined; white-space: pre; border-radius: 0px; word-wrap: normal; background: transparent;">/** Author: Rachel* <zhangruiqing01@baidu.com>** File: bitonic_sort.cu* Create Date: 2015-08-05 17:10:44**/#include<iostream> #include<stdio.h> #include<stdlib.h> #include"gputimer.h" #include<time.h> #define NThreads 8 #define NBlocks 4#define Num NThreads*NBlocksusing namespace Gadgetron;__device__ void swap(int &a, int &b){int t = a;a = b;b = t; }__global__ void bitonic_sort(int* arr){extern __shared__ int shared_arr[];const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;//const unsigned int tid = threadIdx.x;shared_arr[tid] = arr[tid];__syncthreads();//for(int i=2; i<=blociDim.x; i<<=1){for(unsigned int i=2; i<=Num; i<<=1){for(unsigned int j=i>>1; j>0; j>>=1){unsigned int tid_comp = tid ^ j;if(tid_comp > tid){if((tid & i)==0){ //ascendingif(shared_arr[tid]>shared_arr[tid_comp]){swap(shared_arr[tid],shared_arr[tid_comp]);}}else{ //desendingif(shared_arr[tid]<shared_arr[tid_comp]){swap(shared_arr[tid],shared_arr[tid_comp]);}}}__syncthreads();}}arr[tid] = shared_arr[tid]; }int main(int argc, char* argv[]) {GPUTimer timer;int* arr= (int*) malloc(Num*sizeof(int));//init array valuetime_t t;srand((unsigned)time(&t));for(int i=0;i<Num;i++){arr[i] = rand() % 1000; }//init device variableint* ptr;cudaMalloc((void**)&ptr,Num*sizeof(int));cudaMemcpy(ptr,arr,Num*sizeof(int),cudaMemcpyHostToDevice);for(int i=0;i<Num;i++){printf("%d\t",arr[i]);}printf("\n");dim3 blocks(NBlocks,1);dim3 threads(NThreads,1);timer.start();bitonic_sort<<<blocks,threads,Num*sizeof(int)>>>(ptr);//bitonic_sort<<<1,Num,Num*sizeof(int)>>>(ptr);timer.stop();cudaMemcpy(arr,ptr,Num*sizeof(int),cudaMemcpyDeviceToHost);for(int i=0;i<Num;i++){printf("%d\t",arr[i]);}printf("\n");cudaFree(ptr);return 0; } </code><ul class="pre-numbering" style="box-sizing: border-box; position: absolute; width: 50px; top: 0px; left: 0px; margin: 0px; padding: 6px 0px 40px; border-right-width: 1px; border-right-style: solid; border-right-color: rgb(221, 221, 221); list-style: none; text-align: right; background-color: rgb(238, 238, 238);"><li style="box-sizing: border-box; padding: 0px 5px;">1</li><li style="box-sizing: border-box; padding: 0px 5px;">2</li><li style="box-sizing: border-box; padding: 0px 5px;">3</li><li style="box-sizing: border-box; padding: 0px 5px;">4</li><li style="box-sizing: border-box; padding: 0px 5px;">5</li><li style="box-sizing: border-box; padding: 0px 5px;">6</li><li style="box-sizing: border-box; padding: 0px 5px;">7</li><li style="box-sizing: border-box; padding: 0px 5px;">8</li><li style="box-sizing: border-box; padding: 0px 5px;">9</li><li style="box-sizing: border-box; padding: 0px 5px;">10</li><li style="box-sizing: border-box; padding: 0px 5px;">11</li><li style="box-sizing: border-box; padding: 0px 5px;">12</li><li style="box-sizing: border-box; padding: 0px 5px;">13</li><li style="box-sizing: border-box; padding: 0px 5px;">14</li><li style="box-sizing: border-box; padding: 0px 5px;">15</li><li style="box-sizing: border-box; padding: 0px 5px;">16</li><li style="box-sizing: border-box; padding: 0px 5px;">17</li><li style="box-sizing: border-box; padding: 0px 5px;">18</li><li style="box-sizing: border-box; padding: 0px 5px;">19</li><li style="box-sizing: border-box; padding: 0px 5px;">20</li><li style="box-sizing: border-box; padding: 0px 5px;">21</li><li style="box-sizing: border-box; padding: 0px 5px;">22</li><li style="box-sizing: border-box; padding: 0px 5px;">23</li><li style="box-sizing: border-box; padding: 0px 5px;">24</li><li style="box-sizing: border-box; padding: 0px 5px;">25</li><li style="box-sizing: border-box; padding: 0px 5px;">26</li><li style="box-sizing: border-box; padding: 0px 5px;">27</li><li style="box-sizing: border-box; padding: 0px 5px;">28</li><li style="box-sizing: border-box; padding: 0px 5px;">29</li><li style="box-sizing: border-box; padding: 0px 5px;">30</li><li style="box-sizing: border-box; padding: 0px 5px;">31</li><li style="box-sizing: border-box; padding: 0px 5px;">32</li><li style="box-sizing: border-box; padding: 0px 5px;">33</li><li style="box-sizing: border-box; padding: 0px 5px;">34</li><li style="box-sizing: border-box; padding: 0px 5px;">35</li><li style="box-sizing: border-box; padding: 0px 5px;">36</li><li style="box-sizing: border-box; padding: 0px 5px;">37</li><li style="box-sizing: border-box; padding: 0px 5px;">38</li><li style="box-sizing: border-box; padding: 0px 5px;">39</li><li style="box-sizing: border-box; padding: 0px 5px;">40</li><li style="box-sizing: border-box; padding: 0px 5px;">41</li><li style="box-sizing: border-box; padding: 0px 5px;">42</li><li style="box-sizing: border-box; padding: 0px 5px;">43</li><li style="box-sizing: border-box; padding: 0px 5px;">44</li><li style="box-sizing: border-box; padding: 0px 5px;">45</li><li style="box-sizing: border-box; padding: 0px 5px;">46</li><li style="box-sizing: border-box; padding: 0px 5px;">47</li><li style="box-sizing: border-box; padding: 0px 5px;">48</li><li style="box-sizing: border-box; padding: 0px 5px;">49</li><li style="box-sizing: border-box; padding: 0px 5px;">50</li><li style="box-sizing: border-box; padding: 0px 5px;">51</li><li style="box-sizing: border-box; padding: 0px 5px;">52</li><li style="box-sizing: border-box; padding: 0px 5px;">53</li><li style="box-sizing: border-box; padding: 0px 5px;">54</li><li style="box-sizing: border-box; padding: 0px 5px;">55</li><li style="box-sizing: border-box; padding: 0px 5px;">56</li><li style="box-sizing: border-box; padding: 0px 5px;">57</li><li style="box-sizing: border-box; padding: 0px 5px;">58</li><li style="box-sizing: border-box; padding: 0px 5px;">59</li><li style="box-sizing: border-box; padding: 0px 5px;">60</li><li style="box-sizing: border-box; padding: 0px 5px;">61</li><li style="box-sizing: border-box; padding: 0px 5px;">62</li><li style="box-sizing: border-box; padding: 0px 5px;">63</li><li style="box-sizing: border-box; padding: 0px 5px;">64</li><li style="box-sizing: border-box; padding: 0px 5px;">65</li><li style="box-sizing: border-box; padding: 0px 5px;">66</li><li style="box-sizing: border-box; padding: 0px 5px;">67</li><li style="box-sizing: border-box; padding: 0px 5px;">68</li><li style="box-sizing: border-box; padding: 0px 5px;">69</li><li style="box-sizing: border-box; padding: 0px 5px;">70</li><li style="box-sizing: border-box; padding: 0px 5px;">71</li><li style="box-sizing: border-box; padding: 0px 5px;">72</li><li style="box-sizing: border-box; padding: 0px 5px;">73</li><li style="box-sizing: border-box; padding: 0px 5px;">74</li><li style="box-sizing: border-box; padding: 0px 5px;">75</li><li style="box-sizing: border-box; padding: 0px 5px;">76</li><li style="box-sizing: border-box; padding: 0px 5px;">77</li><li style="box-sizing: border-box; padding: 0px 5px;">78</li><li style="box-sizing: border-box; padding: 0px 5px;">79</li><li style="box-sizing: border-box; padding: 0px 5px;">80</li><li style="box-sizing: border-box; padding: 0px 5px;">81</li><li style="box-sizing: border-box; padding: 0px 5px;">82</li><li style="box-sizing: border-box; padding: 0px 5px;">83</li><li style="box-sizing: border-box; padding: 0px 5px;">84</li><li style="box-sizing: border-box; padding: 0px 5px;">85</li><li style="box-sizing: border-box; padding: 0px 5px;">86</li><li style="box-sizing: border-box; padding: 0px 5px;">87</li><li style="box-sizing: border-box; padding: 0px 5px;">88</li><li style="box-sizing: border-box; padding: 0px 5px;">89</li><li style="box-sizing: border-box; padding: 0px 5px;">90</li><li style="box-sizing: border-box; padding: 0px 5px;">91</li><li style="box-sizing: border-box; padding: 0px 5px;">92</li><li style="box-sizing: border-box; padding: 0px 5px;">93</li><li style="box-sizing: border-box; padding: 0px 5px;">94</li><li style="box-sizing: border-box; padding: 0px 5px;">95</li><li style="box-sizing: border-box; padding: 0px 5px;">96</li><li style="box-sizing: border-box; padding: 0px 5px;">97</li><li style="box-sizing: border-box; padding: 0px 5px;">98</li><li style="box-sizing: border-box; padding: 0px 5px;">99</li></ul>

code中，?
tid^j用于單方向判斷，防止同一元素比較兩次；?
tid & i == 0 用于判斷這個部分應(yīng)該是單增還是單減，因?yàn)榉较蛟诿總€長為i的單調(diào)序列中是一致的，所以選用i判斷單調(diào)方向。

參考文獻(xiàn)：?
1.?Bubble sort and its variants?
2.?nvidia的mergesort實(shí)現(xiàn)?
3.?我用過的淺顯易懂的Bitonic sort文檔

from:?http://blog.csdn.net/abcjennifer/article/details/47110991

總結(jié)

以上是生活随笔為你收集整理的CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。

上一篇： CUDA系列学习（五）GPU基础算法:
下一篇： Learning to Rank简介