【CUDA】Dynamic parallelismによるQuicksort

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処理時間を図ってみたのが以下,
f:id:yusuke_ujitoko:20160205144629p:plain
OS:Windows7professional
CPU:Intel Core i7 3.33GHz
GPU:GeForce780Ti

CPUに惨敗...
しかしサンプルを見てみると,もっと良い実装がある模様.