CUDA加速排序算法

排序算法被广泛用在计算应用中。有很多种排序算法，像是枚举排序或者说秩排序，冒泡排序和归并排序。这些排序算法具有不同的(时间和空间)复杂度，因此对同一个数组也有不同的排序时间。对于很大的数组来说，所有的算法都需要很长的时间来完成。如果排序能用CUDA来加速，则会对很多计算应用产生很大帮助。
为了演示CUDA如何能加速不同种类的排序算法，我们将通过CUDA实现一个秩排序算法。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

#define arraySize 5
#define threadPerBlock 5

__global__ void addKernel(int *d_a, int *d_b)
{
	int count = 0;
	int tid = threadIdx.x;
	int ttid = blockIdx.x * threadPerBlock + tid;
	int val = d_a[ttid];
	__shared__ int cache[threadPerBlock];
	for (int i = tid; i < arraySize; i += threadPerBlock)
	{
		cache[tid] = d_a[i];
		__syncthreads();
		for (int j = 0; j < threadPerBlock; ++j)
			if (val > cache[j])
				count++;
		__syncthreads();
	}
	d_b[count] = val;
}

int main()
{
    int h_a[arraySize] = { 5, 9, 3, 4, 8 };
    int h_b[arraySize];
	int *d_a, *d_b;
	 
    cudaMalloc((void**)&d_b, arraySize * sizeof(int));
    cudaMalloc((void**)&d_a, arraySize * sizeof(int));
    
    // Copy input vector from host memory to GPU buffers.
    cudaMemcpy(d_a, h_a, arraySize * sizeof(int), cudaMemcpyHostToDevice);
    
    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<arraySize/threadPerBlock, threadPerBlock>>>(d_a, d_b);
    
    cudaDeviceSynchronize();
    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(h_b, d_b, arraySize * sizeof(int), cudaMemcpyDeviceToHost);
	printf("The Enumeration sorted Array is: \n");
	for (int i = 0; i < arraySize; i++) 
	{
		printf("%d\n", h_b[i]);
	}
    
    cudaFree(d_a);
    cudaFree(d_b);
    return 0;
}

该算法对于数组中的每个元素，通过统计小于它的数组中的其他元素的数量，从而确定该元素在结果数组中的位置。然后，我们根据位置将元素放入结果数组即可。对于(需要排序的）数组中的每个元素都重复进行一次该过程，则我们得到了一个排序后的数组。
内核函数使用两个数组作为参数,d_a是输入数组，而d_b则是结果输出数组。用count变量保存当前元素在排序后的结果数组中的位置。用tid变量保存块中的当前线程索引。而ttid变量则用来表示所有的块中的当前线程的唯一索引，或者说整个Grid中的当前线程的索引。使用共享内存来减少直接访问全局内存的时间。而内核中使用的共享内存中数组的元素数量等于块中的线程数，如同我们之前讨论过的那样。val变量保存了每个线程的当前元素。从全局内存中读取数据，填充到共享内存中。填充的这些数据用来和每个线程的当前val变量做比较。最终用count变量统计出来这些数据中小于当前val变量的数据的个数。因为共享内存中的元素数量有限，所以外层多次循环直到所有的数据都分阶段地填充到共享内存，并和val做比较，统计到count中。当最终的循环完成后，count变量中保存了最终排序数组中的位置，我们则在d_b数组按照这个位置保存当前元素。d_b就是最终的排序后的结果输出数组。