OpenCL实现的归并排序（二）

本文给出一个归并排序算法的例子。本例子支持任意数量元素的数组。代码是从DeepSeek上得到的，只是把C语言的OpenCL库换成了C++版的。这个版本是上个版本的优化版，因为上个版本的一个GPU线程处理一对有序子数组的合并，所以在反复合并几次后线程数会越来越少导致加速效果不好。从代码上可以看出上一版的越循环global size越小。而这个版本则利用有序子数组的合并的特性优化了一下，一个GPU线程只处理一个元素的选择。这样global size可以一直保持不变，从而得到较大的加速效果。

归并排序是通过两个有序子数组之间不断地合并实现的。从1个元素的两个子数组开始合并，然后是2个元素的两个子数组合并，然后是4个元素的两个子数组合并……直到最后把整个数组合并成有序数组。下述代码也是这样实现的，代码中segmentSize就表示子数组的大小，对应的在核函数中索引从left_start到left_end是第一个子数组，从right_start到right_end是第二个子数组，把这两个子数组合并。本算法运行环境是VS2019、OpenCL3，CPU是Intel i5，显卡是英伟达的4060。

本例不需要头文件，下面是CPP文件：

const string strKernel = R"(
__kernel void merge_sort_scalar(
    __global const float* restrict input,
    __global float* restrict output,
    const int segment_size,      // 当前每个有序段的大小（以元素为单位）
    const int total_size         // 总元素数
) {
    int gid = get_global_id(0);  // 全局线程ID，每个线程处理一个输出位置
    
    if (gid >= total_size) return;
    
    // 1. 确定这个输出位置属于哪一对要合并的段
    int segment_pair = gid / (segment_size * 2); /* 这里是除法 */
    int left_start = segment_pair * segment_size * 2;
    int left_end = min(left_start + segment_size, total_size);
    int right_start = left_end;
    int right_end = min(right_start + segment_size, total_size);
    
    // 在这个段对内的相对位置
    int k = gid - left_start;
    
    // 2. 二分搜索：在左段中查找应该有多少个元素在位置k之前
    // 我们要找到最小的i，使得left[i]最接近right[k-i-1]
    // 如果没有这样的i，则i = left_end - left_start
    
    int left_len = left_end - left_start;
    int right_len = right_end - right_start;
    
    int low = 0;
    int high = left_len;  // 搜索范围[0, left_len]
    
    while (low < high) {
        int mid = (low + high) / 2;
        
        float left_val;
        if (mid < left_len) {
            left_val = input[left_start + mid];
        } else {
            left_val = INFINITY;  // 超出左段，设为无穷大
        }
        
        int right_pos = k - mid - 1;
        float right_val;
        if (right_pos >= 0 && right_pos < right_len) {
            right_val = input[right_start + right_pos];
        } else if (right_pos < 0) {
            right_val = -INFINITY;  // 超出右段，设为负无穷大
        } else {
            right_val = INFINITY;   // 超出右段，设为正无穷大
        }
        
        if (left_val <= right_val) {
            low = mid + 1;  // left_val <= right_val，所以mid还可以更大
        } else {
            high = mid;     // left_val > right_val，所以mid应该更小
        }
    }
    
    // 3. 确定最终输出值
    int left_idx = low;           // 左段中应该出现在k之前的元素数量
    int right_idx = k - left_idx; // 右段中应该出现在k之前的元素数量
    
    float result;
    
    // 有几种情况：
    if (left_idx < left_len) {
        // 左段有候选元素
        float left_candidate = input[left_start + left_idx];
        
        if (right_idx < right_len) {
            // 右段也有候选元素，取较小的那个
            float right_candidate = input[right_start + right_idx];
            result = (left_candidate <= right_candidate) ? left_candidate : right_candidate;
        } else {
            // 右段没有候选元素了，取左段
            result = left_candidate;
        }
    } else {
        // 左段没有候选元素了，取右段
        result = input[right_start + right_idx];
    }
    
    output[gid] = result;
})";

int main()
{
    cl::Program program(strKernel);
    try
    {
        program.build("-cl-std=CL2.0");
    }
    catch (...)
    {
        cl_int buildErr = CL_SUCCESS;
        auto buildInfo = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&buildErr);
        for (auto& pair : buildInfo)
        {
            std::cerr << pair.second << std::endl << std::endl;
        }
        return 0;
    }
    auto kernel = cl::KernelFunctor<cl::Buffer, cl::Buffer, int, int>(program, "merge_sort_scalar");

    /* 准备初始数据 */
    const int DATA_SIZE = 518;
    vector<float> number(DATA_SIZE, 100);
    theRNG().fill(number, RNG::UNIFORM, 0, 100); /* OpenCV的函数随机数初始化 */

    cl::Buffer bufferInput(number.begin(), number.end(), false);
    cl::Buffer bufferOutput(CL_MEM_READ_WRITE, DATA_SIZE * sizeof(float));

    int segmentSize = 1;
    bool useBufferAsInput = true;

    while (segmentSize < DATA_SIZE)
    {
        cl::Buffer inputBuffer = useBufferAsInput ? bufferInput : bufferOutput;
        cl::Buffer outputBuffer = useBufferAsInput ? bufferOutput : bufferInput;

        kernel(cl::EnqueueArgs(cl::NDRange(DATA_SIZE), cl::NDRange()), inputBuffer, outputBuffer, segmentSize, DATA_SIZE);

        segmentSize *= 2;
        useBufferAsInput = !useBufferAsInput;
    }
    vector<float> after(DATA_SIZE);
    cl::Buffer finalBuffer = useBufferAsInput ? bufferInput : bufferOutput;
    cl::copy(finalBuffer, after.begin(), after.end());

    for (auto item : after)
    {
        cout << item << " ";
    }

    return 0;
}

下面是程序运行的结果截图。