第8章 Events - 事件对象测试

Events 测试套件验证 OpenCL 事件对象的正确性。事件是 OpenCL 中用于跟踪命令执行状态、实现同步和依赖关系管理的核心机制，对于构建高效的异步执行流程至关重要。

8.1 事件对象概述

8.1.1 什么是事件对象

OpenCL 事件对象（cl_event）用于：

• 跟踪命令状态：监控命令的执行进度（排队、提交、运行、完成）
• 建立依赖关系：通过等待列表确保命令按正确顺序执行
• 同步操作：协调主机和设备之间的执行流程
• 性能分析：记录命令执行的时间戳

事件生命周期：

排队(QUEUED) → 提交(SUBMITTED) → 运行(RUNNING) → 完成(COMPLETE)

8.1.2 事件状态

状态	宏定义	说明
已排队	CL_QUEUED	命令已进入队列
已提交	CL_SUBMITTED	命令已提交到设备
正在运行	CL_RUNNING	命令正在设备上执行
已完成	CL_COMPLETE	命令执行完成
错误	负值	命令执行失败

8.1.3 测试列表

test_definition test_list[] = {
    // 基本事件状态测试
    ADD_TEST(event_get_execute_status),
    ADD_TEST(event_get_write_array_status),
    ADD_TEST(event_get_read_array_status),
    ADD_TEST(event_get_info),
    
    // 事件等待与同步
    ADD_TEST(event_wait_for_execute),
    ADD_TEST(event_wait_for_array),
    ADD_TEST(event_flush),
    ADD_TEST(event_finish_execute),
    ADD_TEST(event_finish_array),
    ADD_TEST(event_release_before_done),
    
    // 标记与屏障
    ADD_TEST(event_enqueue_marker),
    ADD_TEST(event_enqueue_marker_with_event_list),
    ADD_TEST(event_enqueue_barrier_with_event_list),
    
    // 乱序队列测试
    ADD_TEST(out_of_order_event_waitlist_single_queue),
    ADD_TEST(out_of_order_event_waitlist_multi_queue),
    ADD_TEST(out_of_order_event_waitlist_multi_queue_multi_device),
    
    // 等待列表测试
    ADD_TEST(waitlists),
    
    // 用户事件
    ADD_TEST(userevents),
    ADD_TEST(userevents_multithreaded),
    
    // 事件回调
    ADD_TEST(callbacks),
    ADD_TEST(callbacks_simultaneous),
};

---

8.2 事件创建与查询

8.2.1 自动创建事件

大多数 OpenCL 命令都可以返回事件对象：

cl_event event;
cl_int error;

// 内核执行返回事件
error = clEnqueueNDRangeKernel(
    queue, kernel, 1, NULL, &global_size, &local_size,
    0, NULL,      // 无等待事件
    &event);      // 返回事件
test_error(error, "Failed to enqueue kernel");

// Buffer 读取返回事件
error = clEnqueueReadBuffer(
    queue, buffer, CL_FALSE,
    0, size, host_ptr,
    0, NULL,
    &event);
test_error(error, "Failed to enqueue read");

// 记得释放事件
clReleaseEvent(event);

8.2.2 查询事件状态

cl_int status;
cl_int error;

// 查询命令执行状态
error = clGetEventInfo(
    event,
    CL_EVENT_COMMAND_EXECUTION_STATUS,
    sizeof(status),
    &status,
    NULL);
test_error(error, "Failed to get event status");

// 检查状态
switch (status)
{
    case CL_QUEUED:
        log_info("Command queued\n");
        break;
    case CL_SUBMITTED:
        log_info("Command submitted\n");
        break;
    case CL_RUNNING:
        log_info("Command running\n");
        break;
    case CL_COMPLETE:
        log_info("Command complete\n");
        break;
    default:
        log_error("Command failed with error %d\n", status);
        return -1;
}

8.2.3 查询事件信息

// 查询命令类型
cl_command_type cmd_type;
error = clGetEventInfo(event, CL_EVENT_COMMAND_TYPE,
                      sizeof(cmd_type), &cmd_type, NULL);

// 查询关联的命令队列
cl_command_queue event_queue;
error = clGetEventInfo(event, CL_EVENT_COMMAND_QUEUE,
                      sizeof(event_queue), &event_queue, NULL);

// 查询关联的上下文
cl_context event_context;
error = clGetEventInfo(event, CL_EVENT_CONTEXT,
                      sizeof(event_context), &event_context, NULL);

// 查询引用计数
cl_uint ref_count;
error = clGetEventInfo(event, CL_EVENT_REFERENCE_COUNT,
                      sizeof(ref_count), &ref_count, NULL);
log_info("Event reference count: %u\n", ref_count);

---

8.3 事件等待与同步

8.3.1 clWaitForEvents - 等待事件完成

// 等待单个事件
cl_event event;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                              &global_size, NULL,
                              0, NULL, &event);

error = clWaitForEvents(1, &event);
test_error(error, "Failed to wait for event");

// 现在命令肯定已经完成
log_info("Kernel execution completed\n");
clReleaseEvent(event);

等待多个事件：

cl_event events[3];

// 启动三个操作
clEnqueueWriteBuffer(queue, buffer1, CL_FALSE, 0, size,
                    data1, 0, NULL, &events[0]);
clEnqueueWriteBuffer(queue, buffer2, CL_FALSE, 0, size,
                    data2, 0, NULL, &events[1]);
clEnqueueWriteBuffer(queue, buffer3, CL_FALSE, 0, size,
                    data3, 0, NULL, &events[2]);

// 等待所有操作完成
error = clWaitForEvents(3, events);
test_error(error, "Failed to wait for events");

log_info("All write operations completed\n");

// 释放事件
for (int i = 0; i < 3; i++)
    clReleaseEvent(events[i]);

8.3.2 clFlush - 刷新命令队列

clFlush 确保已入队的命令提交到设备，但不等待完成：

// 入队多个命令
clEnqueueNDRangeKernel(queue, kernel1, 1, NULL, &size, NULL,
                      0, NULL, NULL);
clEnqueueNDRangeKernel(queue, kernel2, 1, NULL, &size, NULL,
                      0, NULL, NULL);
clEnqueueNDRangeKernel(queue, kernel3, 1, NULL, &size, NULL,
                      0, NULL, NULL);

// 立即刷新，确保命令提交
error = clFlush(queue);
test_error(error, "Failed to flush queue");

// 命令已提交，但可能还在执行
log_info("Commands flushed to device\n");

// 继续做其他工作...

8.3.3 clFinish - 完成所有命令

clFinish 阻塞直到队列中的所有命令完成：

// 入队多个命令
clEnqueueWriteBuffer(queue, buffer, CL_FALSE, 0, size,
                    data, 0, NULL, NULL);
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      0, NULL, NULL);
clEnqueueReadBuffer(queue, buffer, CL_FALSE, 0, size,
                   result, 0, NULL, NULL);

// 等待所有命令完成
error = clFinish(queue);
test_error(error, "Failed to finish queue");

// 现在所有操作都已完成，result 中有有效数据
log_info("All operations completed\n");

clFlush vs clFinish：

函数	行为	使用场景
clFlush	提交命令但不等待	确保命令及时提交
clFinish	等待所有命令完成	需要结果时同步

---

8.4 事件依赖与等待列表

8.4.1 建立命令依赖关系

通过等待列表（wait list）建立命令间的依赖：

cl_event write_event, kernel_event, read_event;

// 第1步：写入数据到 buffer
error = clEnqueueWriteBuffer(
    queue, buffer, CL_FALSE,
    0, size, input_data,
    0, NULL,           // 无依赖
    &write_event);     // 返回事件

// 第2步：执行内核（依赖于写入完成）
error = clEnqueueNDRangeKernel(
    queue, kernel, 1, NULL, &global_size, NULL,
    1, &write_event,   // 等待写入完成
    &kernel_event);    // 返回事件

// 第3步：读取结果（依赖于内核完成）
error = clEnqueueReadBuffer(
    queue, buffer, CL_FALSE,
    0, size, output_data,
    1, &kernel_event,  // 等待内核完成
    &read_event);      // 返回事件

// 等待最后一个操作完成
clWaitForEvents(1, &read_event);

// 释放事件
clReleaseEvent(write_event);
clReleaseEvent(kernel_event);
clReleaseEvent(read_event);

8.4.2 多依赖示例

一个命令可以依赖多个事件：

cl_event write_events[3];
cl_event kernel_event;

// 并行写入三个 buffer
clEnqueueWriteBuffer(queue, buffer1, CL_FALSE, 0, size,
                    data1, 0, NULL, &write_events[0]);
clEnqueueWriteBuffer(queue, buffer2, CL_FALSE, 0, size,
                    data2, 0, NULL, &write_events[1]);
clEnqueueWriteBuffer(queue, buffer3, CL_FALSE, 0, size,
                    data3, 0, NULL, &write_events[2]);

// 内核等待所有写入完成
error = clEnqueueNDRangeKernel(
    queue, kernel, 1, NULL, &global_size, NULL,
    3, write_events,   // 等待3个写入事件
    &kernel_event);

clWaitForEvents(1, &kernel_event);

// 释放事件
for (int i = 0; i < 3; i++)
    clReleaseEvent(write_events[i]);
clReleaseEvent(kernel_event);

8.4.3 跨队列依赖

事件可以在不同队列之间传递：

cl_command_queue queue1, queue2;
cl_event event_from_queue1;

// 在 queue1 中执行操作
clEnqueueNDRangeKernel(queue1, kernel1, 1, NULL, &size, NULL,
                      0, NULL, &event_from_queue1);

// 在 queue2 中等待 queue1 的操作
clEnqueueNDRangeKernel(queue2, kernel2, 1, NULL, &size, NULL,
                      1, &event_from_queue1,  // 跨队列依赖
                      NULL);

clFinish(queue2);
clReleaseEvent(event_from_queue1);

---

8.5 用户事件

8.5.1 创建用户事件

用户事件允许主机端控制命令的执行：

// 创建用户事件
cl_event user_event = clCreateUserEvent(context, &error);
test_error(error, "Failed to create user event");

// 入队依赖于用户事件的命令
cl_event kernel_event;
error = clEnqueueNDRangeKernel(
    queue, kernel, 1, NULL, &global_size, NULL,
    1, &user_event,    // 等待用户事件
    &kernel_event);

log_info("Kernel queued, waiting for user event\n");

// 做一些主机端工作...
// ...

// 触发用户事件，允许内核执行
error = clSetUserEventStatus(user_event, CL_COMPLETE);
test_error(error, "Failed to set user event status");

log_info("User event triggered, kernel will execute\n");

// 等待内核完成
clWaitForEvents(1, &kernel_event);

clReleaseEvent(user_event);
clReleaseEvent(kernel_event);

8.5.2 用户事件作为门控

用户事件可以作为"门"来控制一批命令的执行：

cl_event gate_event = clCreateUserEvent(context, &error);

// 入队多个依赖门事件的命令
cl_event events[5];
for (int i = 0; i < 5; i++)
{
    clEnqueueNDRangeKernel(queue, kernels[i], 1, NULL,
                          &size, NULL,
                          1, &gate_event,
                          &events[i]);
}

log_info("All kernels queued, waiting for gate\n");

// 等待某个外部条件...
wait_for_external_condition();

// 打开门，所有内核开始执行
clSetUserEventStatus(gate_event, CL_COMPLETE);
log_info("Gate opened, kernels executing\n");

// 等待所有内核完成
clWaitForEvents(5, events);

clReleaseEvent(gate_event);
for (int i = 0; i < 5; i++)
    clReleaseEvent(events[i]);

8.5.3 用户事件错误处理

用户事件可以设置为错误状态：

cl_event user_event = clCreateUserEvent(context, &error);

// 入队依赖命令
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      1, &user_event, NULL);

// 如果发生错误，设置用户事件为错误状态
if (some_error_condition)
{
    // 使用负值表示错误
    clSetUserEventStatus(user_event, -1);
    log_error("User event set to error state\n");
}
else
{
    // 正常完成
    clSetUserEventStatus(user_event, CL_COMPLETE);
}

clReleaseEvent(user_event);

---

8.6 事件回调函数

8.6.1 注册事件回调

事件回调允许在事件达到特定状态时执行自定义代码：

// 回调函数签名
void CL_CALLBACK event_callback(
    cl_event event,
    cl_int event_command_status,
    void *user_data)
{
    log_info("Event callback triggered with status %d\n",
             event_command_status);
    
    // 处理用户数据
    int *counter = (int*)user_data;
    (*counter)++;
}

// 测试函数
int test_event_callbacks()
{
    cl_event event;
    int callback_counter = 0;
    
    // 执行命令
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                          0, NULL, &event);
    
    // 注册回调（当事件完成时调用）
    error = clSetEventCallback(
        event,
        CL_COMPLETE,           // 监听完成状态
        event_callback,        // 回调函数
        &callback_counter);    // 用户数据
    test_error(error, "Failed to set event callback");
    
    // 等待事件完成
    clWaitForEvents(1, &event);
    
    // 回调可能已执行，但不保证立即执行
    // 可能需要短暂等待
    usleep(100000);  // 等待 100ms
    
    if (callback_counter == 0)
    {
        log_error("Callback was not triggered\n");
        return -1;
    }
    
    log_info("Callback was triggered %d times\n", callback_counter);
    clReleaseEvent(event);
    return 0;
}

8.6.2 多状态回调

可以为同一事件的不同状态注册多个回调：

typedef struct {
    cl_int expected_status;
    int callback_index;
} CallbackData;

void CL_CALLBACK status_callback(
    cl_event event,
    cl_int status,
    void *user_data)
{
    CallbackData *data = (CallbackData*)user_data;
    log_info("Callback %d: Expected %d, Got %d\n",
             data->callback_index,
             data->expected_status,
             status);
}

// 为不同状态注册回调
CallbackData data[3];
data[0] = (CallbackData){ CL_SUBMITTED, 0 };
data[1] = (CallbackData){ CL_RUNNING, 1 };
data[2] = (CallbackData){ CL_COMPLETE, 2 };

cl_event event;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      0, NULL, &event);

// 注册三个回调
clSetEventCallback(event, CL_SUBMITTED, status_callback, &data[0]);
clSetEventCallback(event, CL_RUNNING, status_callback, &data[1]);
clSetEventCallback(event, CL_COMPLETE, status_callback, &data[2]);

clWaitForEvents(1, &event);
clReleaseEvent(event);

8.6.3 回调注意事项

重要限制：

• 回调在独立线程中执行
• 回调中不应调用阻塞的 OpenCL 函数（如 clWaitForEvents、clFinish）
• 回调执行时间应尽量短
• 回调可能在任何时刻触发，即使在 clWaitForEvents 之后

// 危险：在回调中调用阻塞函数
void CL_CALLBACK bad_callback(cl_event event, cl_int status, void *data)
{
    // 不要这样做！可能导致死锁
    clFinish(some_queue);  // 危险！
}

// 安全：使用标志通知
bool g_event_completed = false;

void CL_CALLBACK safe_callback(cl_event event, cl_int status, void *data)
{
    // 只设置标志，不调用阻塞函数
    g_event_completed = true;
}

---

8.7 标记与屏障

8.7.1 clEnqueueMarker - 插入标记

标记用于获取之前所有命令的完成事件：

// OpenCL 1.1 及之前
cl_event marker_event;
error = clEnqueueMarker(queue, &marker_event);
test_error(error, "Failed to enqueue marker");

// marker_event 会在之前所有命令完成后变为 COMPLETE
clWaitForEvents(1, &marker_event);
log_info("All previous commands completed\n");
clReleaseEvent(marker_event);

8.7.2 clEnqueueMarkerWithWaitList - 带等待列表的标记

OpenCL 1.2+ 支持带等待列表的标记：

#ifdef CL_VERSION_1_2
cl_event events[3];

// 执行多个操作
clEnqueueWriteBuffer(queue, buf1, CL_FALSE, 0, size,
                    data1, 0, NULL, &events[0]);
clEnqueueWriteBuffer(queue, buf2, CL_FALSE, 0, size,
                    data2, 0, NULL, &events[1]);
clEnqueueWriteBuffer(queue, buf3, CL_FALSE, 0, size,
                    data3, 0, NULL, &events[2]);

// 创建标记，等待特定事件
cl_event marker;
error = clEnqueueMarkerWithWaitList(
    queue,
    3, events,    // 等待这3个事件
    &marker);
test_error(error, "Failed to enqueue marker");

// marker 完成意味着这3个写入都完成了
clWaitForEvents(1, &marker);

for (int i = 0; i < 3; i++)
    clReleaseEvent(events[i]);
clReleaseEvent(marker);
#endif

8.7.3 clEnqueueBarrier - 插入屏障

屏障确保之前的所有命令完成后才执行后续命令：

// OpenCL 1.1
// 执行一批命令
clEnqueueWriteBuffer(queue, buffer, CL_FALSE, 0, size, data, 0, NULL, NULL);
clEnqueueNDRangeKernel(queue, kernel1, 1, NULL, &size, NULL, 0, NULL, NULL);

// 插入屏障
error = clEnqueueBarrier(queue);
test_error(error, "Failed to enqueue barrier");

// 屏障后的命令必须等待之前所有命令完成
clEnqueueNDRangeKernel(queue, kernel2, 1, NULL, &size, NULL, 0, NULL, NULL);
clEnqueueReadBuffer(queue, buffer, CL_FALSE, 0, size, result, 0, NULL, NULL);

clFinish(queue);

8.7.4 clEnqueueBarrierWithWaitList - 带等待列表的屏障

OpenCL 1.2+ 支持更灵活的屏障：

#ifdef CL_VERSION_1_2
cl_event events[2];

clEnqueueNDRangeKernel(queue, kernel1, 1, NULL, &size, NULL,
                      0, NULL, &events[0]);
clEnqueueNDRangeKernel(queue, kernel2, 1, NULL, &size, NULL,
                      0, NULL, &events[1]);

// 插入屏障，等待特定事件
cl_event barrier;
error = clEnqueueBarrierWithWaitList(
    queue,
    2, events,    // 等待这2个内核
    &barrier);
test_error(error, "Failed to enqueue barrier");

// barrier 后的命令等待上述2个内核完成
clEnqueueReadBuffer(queue, buffer, CL_FALSE, 0, size,
                   result, 1, &barrier, NULL);

clReleaseEvent(events[0]);
clReleaseEvent(events[1]);
clReleaseEvent(barrier);
#endif

---

8.8 乱序队列中的事件

8.8.1 乱序队列概述

默认情况下，命令队列按顺序执行命令。乱序队列允许命令以任意顺序执行：

// 创建乱序队列
cl_command_queue out_of_order_queue = clCreateCommandQueue(
    context, device,
    CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,  // 启用乱序执行
    &error);
test_error(error, "Failed to create out-of-order queue");

8.8.2 乱序队列中的依赖管理

在乱序队列中，必须显式管理依赖：

cl_event write_event, kernel_event, read_event;

// 写入数据
clEnqueueWriteBuffer(out_of_order_queue, buffer, CL_FALSE,
                    0, size, data, 0, NULL, &write_event);

// 执行内核（显式依赖写入）
clEnqueueNDRangeKernel(out_of_order_queue, kernel, 1, NULL,
                      &size, NULL,
                      1, &write_event,      // 必须显式指定
                      &kernel_event);

// 读取结果（显式依赖内核）
clEnqueueReadBuffer(out_of_order_queue, buffer, CL_FALSE,
                   0, size, result,
                   1, &kernel_event,        // 必须显式指定
                   &read_event);

// 等待最后操作
clWaitForEvents(1, &read_event);

clReleaseEvent(write_event);
clReleaseEvent(kernel_event);
clReleaseEvent(read_event);

8.8.3 乱序队列的优势

允许独立命令并行执行：

cl_event events[4];

// 这4个内核相互独立，可以并行执行
clEnqueueNDRangeKernel(out_of_order_queue, kernel1, 1, NULL,
                      &size, NULL, 0, NULL, &events[0]);
clEnqueueNDRangeKernel(out_of_order_queue, kernel2, 1, NULL,
                      &size, NULL, 0, NULL, &events[1]);
clEnqueueNDRangeKernel(out_of_order_queue, kernel3, 1, NULL,
                      &size, NULL, 0, NULL, &events[2]);
clEnqueueNDRangeKernel(out_of_order_queue, kernel4, 1, NULL,
                      &size, NULL, 0, NULL, &events[3]);

// 设备可以并行执行这些内核
clWaitForEvents(4, events);

for (int i = 0; i < 4; i++)
    clReleaseEvent(events[i]);

---

8.9 事件测试最佳实践

8.9.1 正确管理事件生命周期

// 好的做法：始终释放事件
cl_event event;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      0, NULL, &event);
clWaitForEvents(1, &event);
clReleaseEvent(event);  // 必须释放

// 不好的做法：忘记释放事件
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      0, NULL, &event);
// 忘记 clReleaseEvent(event) - 内存泄漏！

8.9.2 使用事件包装器

// C++ 包装器自动管理生命周期
clEventWrapper event;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                              &size, NULL, 0, NULL, &event);
// 析构时自动释放

8.9.3 避免不必要的同步

// 不高效：频繁同步
for (int i = 0; i < 100; i++)
{
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                          0, NULL, NULL);
    clFinish(queue);  // 每次都同步 - 慢！
}

// 高效：批量执行后同步
for (int i = 0; i < 100; i++)
{
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                          0, NULL, NULL);
}
clFinish(queue);  // 只同步一次

8.9.4 合理使用阻塞与非阻塞

// 阻塞读取（简单但效率低）
clEnqueueReadBuffer(queue, buffer, CL_TRUE,  // 阻塞
                   0, size, result, 0, NULL, NULL);
// 函数返回时数据已就绪

// 非阻塞读取（更高效）
cl_event read_event;
clEnqueueReadBuffer(queue, buffer, CL_FALSE,  // 非阻塞
                   0, size, result, 0, NULL, &read_event);
// 可以继续做其他工作...
do_other_work();
// 需要数据时再等待
clWaitForEvents(1, &read_event);
clReleaseEvent(read_event);

8.9.5 检查事件错误

cl_event event;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL,
                      0, NULL, &event);

// 等待事件
cl_int wait_status = clWaitForEvents(1, &event);

// 检查等待是否成功
if (wait_status != CL_SUCCESS)
{
    log_error("Wait failed: %d\n", wait_status);
}

// 检查命令执行状态
cl_int exec_status;
clGetEventInfo(event, CL_EVENT_COMMAND_EXECUTION_STATUS,
              sizeof(exec_status), &exec_status, NULL);

if (exec_status < 0)
{
    log_error("Command failed with error: %d\n", exec_status);
}

clReleaseEvent(event);

---

本章小结

第8章详细介绍了 Events 测试套件的核心内容：

• 事件基础：事件状态、生命周期、创建与查询
• 等待与同步：clWaitForEvents、clFlush、clFinish
• 依赖管理：等待列表、跨队列依赖、多依赖关系
• 用户事件：主机端门控、错误处理
• 事件回调：回调注册、多状态监听、使用限制
• 标记与屏障：clEnqueueMarker、clEnqueueBarrier
• 乱序队列：并行执行、显式依赖管理
• 最佳实践：生命周期管理、避免过度同步、错误检查

事件是 OpenCL 异步执行模型的核心，正确使用事件可以构建高效的并行执行流程，充分发挥设备性能。

下一章将介绍 Profiling（性能分析）测试，详细讲解如何使用 OpenCL 事件进行性能测量和分析。