OpenCL奇偶排序程序异常求助：元素丢失、重复及错位问题

我需要实现基于奇偶排序（odd-even sort）的OpenCL程序，作为OpenCL新手完成代码后遇到异常：排序后原数组部分元素丢失、出现重复值且存在错位（原数组无重复元素）。以下是我的代码：

原Kernel代码 #

const char* kernelSource =
"__kernel void odd_even_sort(__global ulong * arr, const unsigned int n) {
"
"    unsigned int id = get_global_id(0);
"
"    for (unsigned int phase = 0; phase < n; phase++) {
"
"        unsigned int swap_idx;
"
"        if (phase % 2 == 0) {
"
"            swap_idx = id * 2;
"
"        }
"
"        else {
"
"            swap_idx = id * 2 + 1;
"
"        }
"
"        if (swap_idx + 1 < n) {
"
"            if (arr[swap_idx] > arr[swap_idx + 1]) {
"
"                ulong temp = arr[swap_idx];
"
"                arr[swap_idx] = arr[swap_idx + 1];
"
"                arr[swap_idx + 1] = temp;
"
"            }
"
"        }
"
"        barrier(CLK_GLOBAL_MEM_FENCE);
"
"    }
"
"}
";

原主机端运行代码 #

clSetKernelArg(kernel, 0, sizeof(cl_mem), (cl_ulong*)&A_mem);
clSetKernelArg(kernel, 1, sizeof(unsigned int), &nums);
size_t global_work_size = nums / 2;
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);

其中nums是数组长度且为2的幂。

问题原因分析 #

1. Barrier同步失效：Kernel内的barrier(CLK_GLOBAL_MEM_FENCE)仅能同步同一工作组内的工作项，若未指定局部工作大小，OpenCL会自动划分多个工作组，不同组的工作项无法被该Barrier同步。这会导致部分工作项提前进入下一个相位，并发操作全局内存时出现数据竞争，最终引发元素重复、丢失或错位。
2. Kernel参数类型错误：clSetKernelArg的第一个参数传递了(cl_ulong*)&A_mem，但A_mem是cl_mem类型，强制类型转换会导致未定义行为，可能破坏参数传递的正确性。
3. 全局工作大小不合理：奇数相位时，有效的工作项数量应为(nums-1)/2，但固定使用nums/2会产生无意义的空操作工作项，虽不直接引发错误，但浪费计算资源。

修复方案 #

方案1：拆分相位为独立Kernel调用（推荐，适合大数组） #

将奇偶排序的每一个相位拆分为单独的Kernel调用，在主机端同步，确保所有工作项完成当前相位后再进入下一个。

修改后的Kernel代码

// 偶数相位Kernel：处理索引0&1, 2&3,...
const char* even_phase_kernel =
"__kernel void even_phase(__global ulong * arr, const unsigned int n) {
"
"    unsigned int id = get_global_id(0);
"
"    unsigned int swap_idx = id * 2;
"
"    if (swap_idx + 1 < n) {
"
"        if (arr[swap_idx] > arr[swap_idx + 1]) {
"
"            ulong temp = arr[swap_idx];
"
"            arr[swap_idx] = arr[swap_idx + 1];
"
"            arr[swap_idx + 1] = temp;
"
"        }
"
"    }
"
"}
";

// 奇数相位Kernel：处理索引1&2, 3&4,...
const char* odd_phase_kernel =
"__kernel void odd_phase(__global ulong * arr, const unsigned int n) {
"
"    unsigned int id = get_global_id(0);
"
"    unsigned int swap_idx = id * 2 + 1;
"
"    if (swap_idx + 1 < n) {
"
"        if (arr[swap_idx] > arr[swap_idx + 1]) {
"
"            ulong temp = arr[swap_idx];
"
"            arr[swap_idx] = arr[swap_idx + 1];
"
"            arr[swap_idx + 1] = temp;
"
"        }
"
"    }
"
"}
";

修改后的主机端代码

// 创建两个Kernel对象
cl_kernel even_kernel = clCreateKernel(program, "even_phase", &ret);
cl_kernel odd_kernel = clCreateKernel(program, "odd_phase", &ret);

// 设置Kernel参数（修正类型错误）
ret = clSetKernelArg(even_kernel, 0, sizeof(cl_mem), &A_mem);
ret = clSetKernelArg(even_kernel, 1, sizeof(unsigned int), &nums);
ret = clSetKernelArg(odd_kernel, 0, sizeof(cl_mem), &A_mem);
ret = clSetKernelArg(odd_kernel, 1, sizeof(unsigned int), &nums);

// 循环执行每一个相位
for (unsigned int phase = 0; phase < nums; phase++) {
    size_t global_size;
    cl_kernel current_kernel;
    if (phase % 2 == 0) {
        current_kernel = even_kernel;
        global_size = nums / 2;
    } else {
        current_kernel = odd_kernel;
        global_size = (nums - 1) / 2;
    }
    // 执行Kernel
    ret = clEnqueueNDRangeKernel(command_queue, current_kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
    // 等待当前相位完成，确保所有工作项同步
    ret = clFinish(command_queue);
}

// 释放资源
clReleaseKernel(even_kernel);
clReleaseKernel(odd_kernel);

方案2：单Kernel内使用工作组同步（仅适合小数组） #

若数组长度较小（不超过设备最大工作组尺寸，通常为1024），可设置全局工作大小等于局部工作大小，确保所有工作项在同一工作组内，此时Kernel内的Barrier可有效同步。

修改后的主机端代码

ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &A_mem); // 修正类型错误
ret = clSetKernelArg(kernel, 1, sizeof(unsigned int), &nums);
size_t global_work_size = nums / 2;
size_t local_work_size = global_work_size; // 所有工作项在同一工作组
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

注意：此方案仅适用于小数组，若nums/2超过设备最大工作组尺寸，clEnqueueNDRangeKernel会返回错误。

关键修复点总结 #

• 修正clSetKernelArg的参数类型，直接传递&A_mem而非强制转换。
• 确保每一轮相位的所有工作项完成后再进入下一轮，优先选择拆分Kernel+主机端同步的方案。
• 针对不同相位设置合理的全局工作大小，避免无意义的空操作。

内容的提问来源于stack exchange，提问作者kylo_gg