同OpenCL上下文多GPU间如何用clEnqueueMigrateMemObjects传缓冲区

箭头标注流程的核心逻辑 #

你标注的「保存缓冲区1处理流程关联的最后一个事件」是跨设备内存同步的核心，OpenCL所有入队API都是异步执行，API返回仅代表命令加入队列，不代表GPU已经执行完成：

• 源设备完成对共享缓冲区的写操作后，必须捕获该操作对应的完成事件，才能标记共享缓冲区的最新数据已经就绪
• 目标设备发起缓冲区迁移操作时，需要将源设备的完成事件作为依赖项传入，OpenCL调度器会自动阻塞迁移操作，直到源设备所有前置写操作全部完成，从根本上避免读写数据竞争。

示例代码问题修正 #

你当前编写的代码存在4个核心问题，会直接导致传输失败或数据错误：

• 内核参数传参错误：clSetKernelArg传入缓冲区参数时，需要传入cl_mem类型缓冲区对象的指针，而非主机端数组地址，你代码中传入&dev0_arr、&common_arr的写法不符合接口要求
• 迁移标志使用错误：CL_MIGRATE_MEM_OBJECT_HOST会强制将缓冲区迁回主机内存，跨GPU传输不需要绕主机中转，传入标志值0即可，驱动会自动选择最优路径（支持P2P的硬件会直接走GPU间直连通道，不需要经过主机内存拷贝，延迟更低）
• 缺失事件依赖链：当前代码所有入队操作未配置事件依赖，会出现源设备还未完成缓冲区写入，目标设备就已经开始读取的竞态问题，得到的数据是未更新的脏数据
• 共享缓冲区重复初始化：同一上下文下的cl_mem对象是全局可访问的，不需要在两个设备的命令队列中分别对common_buf执行初始化写入，单次初始化即可。

修正后的单时间步双向传输核心代码如下：

cl_event dev0_calc_done, migrate_to_dev1_done, dev1_calc_done, migrate_to_dev0_done;

// 设备0执行本侧计算，更新共享边界缓冲区
cl_int status = clEnqueueNDRangeKernel(cmdq_dev0, kernel0, 3, NULL, global_size, local_size,
                                       0, NULL, &dev0_calc_done);
// 将共享缓冲区迁移到设备1，迁移命令入队到目标设备（设备1）的队列，依赖设备0计算完成事件
status |= clEnqueueMigrateMemObjects(cmdq_dev1, 1, &common_buf, 0,
                                     1, &dev0_calc_done, &migrate_to_dev1_done);
// 设备1执行本侧计算，依赖缓冲区迁移完成事件
status |= clEnqueueNDRangeKernel(cmdq_dev1, kernel1, 3, NULL, global_size, local_size,
                                 1, &migrate_to_dev1_done, &dev1_calc_done);

// 双向传输：将设备1更新后的共享缓冲区迁回设备0，迁移命令入队到目标设备（设备0）的队列，依赖设备1计算完成事件
status |= clEnqueueMigrateMemObjects(cmdq_dev0, 1, &common_buf, 0,
                                     1, &dev1_calc_done, &migrate_to_dev0_done);
// 设备0启动下一时间步计算，依赖缓冲区迁回完成事件
status |= clEnqueueNDRangeKernel(cmdq_dev0, kernel0, 3, NULL, global_size, local_size,
                                 1, &migrate_to_dev0_done, NULL);

// 释放事件对象避免内存泄漏
clReleaseEvent(dev0_calc_done);
clReleaseEvent(migrate_to_dev1_done);
clReleaseEvent(dev1_calc_done);
clReleaseEvent(migrate_to_dev0_done);

双向传输的命令队列选择规则 #

clEnqueueMigrateMemObjects的命令队列选择逻辑非常明确：

• 迁移命令永远入队到缓冲区要迁往的目标设备所绑定的命令队列
• 从设备0向设备1传common_buf时，目标设备是设备1，迁移命令入队到cmdq_dev1
• 从设备1向设备0传common_buf时，目标设备是设备0，迁移命令入队到cmdq_dev0

注意：禁止两个设备同时对同一个共享缓冲区执行写操作，必须严格遵循「源设备写完成→缓冲区迁移→目标设备读写」的顺序，否则会出现未定义的计算错误。

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