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