http://blogs.nvidia.com/blog/2012/09/12/how-tesla-k20-speeds-up-quicksort-a-familiar-comp-sci-code/CC3.5以上で使えるDynamic parallelismは,デバイス内で再帰的にカーネルを呼び出すことできるので,
Quicksortなどが簡易に実装できる.
【CUDA】Dynamic Parallelism - 緑茶思考ブログ
たとえば,quicksortをそのまま単純に実装したのが↓.
#include <stdio.h> #define N 10 #define BLOCK_SIZE 1 #define GRID_SIZE 1 __device__ int pivot(int *a, int i, int j){ int k = i + 1; while(k<=j && a[k]==a[i]){ k++; } if(k>j){ return -1; } if(a[i] >= a[k]){ return i; } return k; } __device__ int partition(int *a, int i, int j, int x){ int left = i; int right = j; while(left <= right){ while(left <= j && a[left] < x){ left++; } while(right >= i && a[right] >= x){ right--; } if(left > right){ break; } int tmp = a[left]; a[left] = a[right]; a[right] = tmp; left++; right--; } return left; } __global__ void quickSort(int *a, int *sorted_a, int n, int i, int j){ if(threadIdx.x < n){ if(i==j){ return; } int p = pivot(a, i, j); if(p != -1){ int k = partition(a,i,j,a[p]); quickSort<<<GRID_SIZE, BLOCK_SIZE>>>(a, sorted_a, n, i, k-1); quickSort<<<GRID_SIZE, BLOCK_SIZE>>>(a, sorted_a, n, k, j); } } } int main(void){ int n = N; int nbytes = n*sizeof(int); int *a = 0; cudaMallocHost((void**)&a, nbytes); for(int i=0;i<n;i++){ a[i] = rand(); printf("a[%d] = %d\n",i,a[i]); } int *a_dev = 0; int *sorted_a_dev = 0; cudaMalloc((void**)&a_dev, nbytes); cudaMemcpy(a_dev, a, nbytes, cudaMemcpyHostToDevice); cudaMalloc((void**)&sorted_a_dev, nbytes); quickSort<<<GRID_SIZE, BLOCK_SIZE>>>(a_dev, sorted_a_dev, n, 0, n-1); cudaDeviceSynchronize(); cudaMemcpy(a, a_dev, nbytes, cudaMemcpyDeviceToHost); printf("\n"); for(int i=0;i<n;i++){ printf("a[%d] = %d\n",i,a[i]); } return 0; }
HOW TESLA K20 SPEEDS QUICKSORT, A FAMILIAR COMP-SCI CODE | The Official NVIDIA Blog
では,CUDAによるDynamic parallelismありなしでの計算時間が載っている.
CPU-GPU間のデータ転送が少ないため,Dynamic parallelismを使った時のほうが処理が速い.
また,CPUだけでquicksort実行した場合とDynamic parallelismを使ったGPU処理時間を図ってみたのが以下,
OS:Windows7professional
CPU:Intel Core i7 3.33GHz
GPU:GeForce780Ti
CPUに惨敗...
しかしサンプルを見てみると,もっと良い実装がある模様.