奇偶冒泡排序的CUDA实现

由于奇偶冒泡排序分奇下标和偶下标排序，这就确保了在一轮排序【奇/偶】过程中，隔一个处理一个，同一轮的排序中被处理的各元素是互不影响的，所以可以并行处理。。。事实上，奇偶冒泡排序在串行CPU上本没一点优势,他揍是为并行而生.

[cpp]  view plain copy

1. #include <cuda.h>  
2. #include <time.h>  
3. #include <stdio.h>  
4. #include <stdlib.h>  
5. #include <math.h>  
6. #include <pz_cpy>  
7. #include "cuPrintf.cu"  
8. bool InitGPUSet()  
9. {  
10.     cudaDeviceProp tCard;  
11.     int num = 0;  
12.     if(cudaSuccess == cudaGetDeviceCount(&num))  
13.     {  
14.         for(int i = 0; i < num; ++ i)  
15.         {  
16.             cudaSetDevice(i);  
17.             cudaGetDeviceProperties(&tCard, i);  
18.             printf("GPU: %s/n", tCard.name);  
19.         }  
20.     }  
21.     else  return false;  
22.     return true;  
23. }  
24. bool InitCuPrint()  
25. {  
26.     cudaError_t err = cudaPrintfInit();  
27.     return 0 == strcmp("no error", cudaGetErrorString(err));  
28. }  
29. __global__ void even_sort(int* ary, int size)  
30. {  
31.     int tid = threadIdx.x;//线程从0开始编号  
32.     if(1 == (tid + 1 ) % 2)//第奇数个轮回  
33.     {  
34.         cuPrintf("before swap %d/n", tid);  
35.         for(int i = 0; i < size; ++ i)  cuPrintf("%d ", ary[i]);  
36.         cuPrintf("/n");  
37.         if(ary[tid] > ary[tid + 1] && tid + 1 < size)  
38.         {  
39.             int tp = ary[tid];  
40.             ary[tid] = ary[tid + 1];  
41.             ary[tid + 1] = tp;  
42.             cuPrintf("after swap %d/n", tid);  
43.             for(int i = 0; i < size; ++ i)  cuPrintf("%d ", ary[i]);  
44.             cuPrintf("/n");  
45.         }  
46.     }  
47.     __syncthreads();  
48.     //cuPrintf("block: %d/t", bid);  
49.     //for(int i = 0; i < size; ++ i)  cuPrintf("%d ", ary[i]);  
50.     //cuPrintf("/n");  
51. }  
52. __global__ void odd_sort(int* ary, int size)  
53. {  
54.     int tid = threadIdx.x;//线程从0开始编号  
55.     if(0 == (tid + 1 ) % 2)//第奇数个轮回  
56.     {  
57.         /*cuPrintf("before swap %d/n", tid); 
58.         for(int i = 0; i < size; ++ i)  cuPrintf("%d ", ary[i]); 
59.         cuPrintf("/n");*/  
60.         if(ary[tid] > ary[tid + 1] && tid + 1 < size)  
61.         {  
62.             int tp = ary[tid];  
63.             ary[tid] = ary[tid + 1];  
64.             ary[tid + 1] = tp;  
65. /* 
66.             cuPrintf("after swap %d/n", tid); 
67.             for(int i = 0; i < size; ++ i)  cuPrintf("%d ", ary[i]); 
68.             cuPrintf("/n"); 
69.             */  
70.         }  
71.     }  
72.     //__syncthreads();  
73. }  
74. int main(void)  
75. {  
76.     const int SIZE = 27;  
77.     const int BLOCK_SIZE = SIZE;  
78.     const int THREAD_SIZE = SIZE;  
79.     if(!InitGPUSet())  puts("GPU failed");  
80.     else  if(!InitCuPrint())  puts("cuPrintf failed");  
81.     else  
82.     {  
83.         int ary[SIZE];// = {6, 5, 4, 3, 2, 1};  
84.         srand(time(0));  
85.         for(int i = 0; i < SIZE; ++ i)  ary[i] = rand() % 100;  
86.         for(int i = 0; i < SIZE; ++ i)  printf("%d ", ary[i]);  
87.         puts("");  
88.         int* gary;  
89.         cudaMalloc((void**) &gary, SIZE * sizeof(int));  
90.         cudaMemcpy(gary, ary, SIZE * sizeof(int), cudaMemcpyHostToDevice);  
91.         for(int i = 0; i < BLOCK_SIZE; ++ i)  
92.         {  
93.             even_sort<<<1, THREAD_SIZE, 0>>>(gary, SIZE);  
94.             //cudaPrintfDisplay(stdout, false);  
95.             odd_sort<<<1, THREAD_SIZE, 0>>>(gary, SIZE);  
96.         }  
97.         //cudaPrintfEnd();  
98.         cudaMemcpy(ary, gary, SIZE * sizeof(int), cudaMemcpyDeviceToHost);  
99.         for(int i = 0; i < SIZE; ++ i)  printf("%d ", ary[i]);  
100.         puts("");  
101.         cudaFree(gary);  
102.     }  
103.     return 0;  
104. }