[12] 使用 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;
}

count 变量保存当前元素在排序后的结果数组中的位置
使用共享内存来减少直接访问全局内存的时间
val变量保存了每个每个线程的当前元素。
从全局内存中读取数据，填充到共享内存，填充的这些数据用来和每个线程的当前val变量做比较
最终用count变量统计出这些数据中小于当前val 变量的数据的个数，因为共享内存中的元素数量有限，所以外层多次循环知道所有的数据都分阶段地填充到共享内存，并和val做比较，统计到count中
当最终地循环完成后，count变量中保存了最终排序数组地位置，我们则在d_b数组中按照这个位置保存当前元素，d_b就是最终地排序后地结果输出数组
main（）函数代码如下：

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;
}