OpenCL中如何等待设备端队列完成后再从主机端读取结果

问题代码 #

// 主机端调用主内核
err = clEnqueueNDRangeKernel(queue, main_kernel, 1, NULL, globalSize_g, 0, 0, NULL, NULL);

// 设备端的主内核实现如下：

__kernel void main_kernel(...)
{
    for (int i = 0; i <= 1000; ...)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_NO_WAIT, ...);
        // 多个子内核提交
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_NO_WAIT, ...);
    }
}

// 主机端从缓冲区读取结果
clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, ...);

问题描述 #

主内核在设备端循环提交多个子内核，且子内核需要按顺序执行（前一个完成后再执行下一个）。但主内核循环结束时间远早于所有子内核执行完毕，导致主机端读取缓冲区时，部分子内核的结果还未生成，数据不完整。

尝试过的无效方案：

• 读取前调用clFinish()等待队列，未解决问题
• 等待主内核完成的事件，仅能跟踪主内核结束，无法覆盖子内核
• 考虑过设备端同步函数，但未找到对应OpenCL规范接口；全局变量轮询不合理

解决方案 #

核心思路 #

设备端通过enqueue_kernel提交的子内核属于对应主机端命令队列的异步任务，主内核的完成事件仅标记自身执行结束，不包含子内核。需通过事件链+用户事件将设备端子内核的完成状态传递到主机端，确保主机端仅在所有子内核执行完毕后读取数据。

具体实现步骤 #

1. 主机端准备用户事件

创建一个用户事件，用于接收设备端所有子内核完成的信号：

cl_event user_event;
cl_int err = clCreateUserEvent(context, &err);

2. 修改设备端主内核

在主内核中维护子内核的事件链，确保子内核按顺序执行，最后等待所有子内核完成后标记用户事件：

__kernel void main_kernel(__global cl_event *user_event, ...)
{
    event_t prev_event = NULL;
    event_t last_sub_event = NULL;

    for (int i = 0; i <= 1000; i++)
    {
        event_t curr_sub_event;
        // 提交子内核，传入前一个子内核的事件作为依赖，保证顺序执行
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_NO_WAIT,
                       // 此处传入子内核的ndrange参数
                       prev_event, &curr_sub_event,
                       // 子内核及参数
                       sub_kernel, ...);
        
        prev_event = curr_sub_event;
        last_sub_event = curr_sub_event;
    }

    // 等待最后一个子内核执行完成
    enqueue_wait_events(get_default_queue(), 1, &last_sub_event);
    // 标记用户事件为完成，通知主机端所有子内核已结束
    set_user_event_status(*user_event, CL_COMPLETE);
}

3. 主机端关联事件并等待

将用户事件作为参数传入主内核，提交主内核后等待用户事件完成，再执行读取操作：

// 将用户事件设置为主内核参数（需替换为对应参数索引）
err = clSetKernelArg(main_kernel, USER_EVENT_ARG_INDEX, sizeof(cl_event), &user_event);

// 提交主内核
err = clEnqueueNDRangeKernel(queue, main_kernel, 1, NULL, globalSize_g, 0, 0, NULL, NULL);

// 等待用户事件完成（即所有子内核执行完毕）
err = clEnqueueWaitForEvents(queue, 1, &user_event);

// 读取缓冲区数据
err = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, ...);

// 释放用户事件资源
clReleaseEvent(user_event);

额外注意事项 #

• 确保使用OpenCL 2.0及以上版本，enqueue_kernel和set_user_event_status是OpenCL 2.0引入的特性
• 如果你的设备不支持OpenCL 2.0，需改为主机端循环提交子内核，通过主机端事件链控制依赖，避免设备端提交子内核的场景
• 若调用clFinish(queue)无效，检查队列是否启用了CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE，这种情况下必须通过显式事件依赖实现同步

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