排序算法被广泛用在计算应用中。有很多种排序算法,像是枚举排序或者说秩排序,冒泡排序和归并排序。这些排序算法具有不同的(时间和空间)复杂度,因此对同一个数组也有不同的排序时间。对于很大的数组来说,所有的算法都需要很长的时间来完成。如果排序能用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就是最终的排序后的结果输出数组。