OpenCL 内核是异步执行的吗？答案

【问题标题】：Are OpenCL kernels executed asynchronously?OpenCL 内核是异步执行的吗？
【发布时间】：2020-03-15 16:21:03
【问题描述】：

对于 CUDA，我知道它们是在向默认流（空流）发出启动命令后异步执行的，那么在 OpenCL 中呢？示例代码如下：

cl_context context;
cl_device_id device_id;
cl_int err;
...
cl_kernel kernel1;
cl_kernel kernel2;
cl_command_queue Q = clCreateCommandQueue(context, device_id, 0, &err);
...
size_t global_w_offset[3] = {0,0,0};
size_t global_w_size[3] = {16,16,1};
size_t local_w_size[3] = {16,16,1};
err = clEnqueueNDRangeKernel(Q, kernel1, 3, global_w_offset, global_w_size, 
                             local_w_size, 0, nullptr, nullptr);
err = clEnqueueNDRangeKernel(Q, kernel2, 3, global_w_offset, global_w_size, 
                             local_w_size, 0, nullptr, nullptr);
clFinish(Q);

kernel1 和 kernel2 在命令入队后会异步执行吗？（即执行重叠）

更新
根据OpenCL Reference，似乎在clCreateCommandQueue 中设置properties 为CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 可以满足我的需要。但是，out_of_order 是否意味着异步？

【问题讨论】：

标签： c++ asynchronous opencl gpgpu

【解决方案1】：

在您当前的代码中：

err = clEnqueueNDRangeKernel(Q, kernel1, 3, global_w_offset, global_w_size, 
                             local_w_size, 0, nullptr, nullptr);
err = clEnqueueNDRangeKernel(Q, kernel2, 3, global_w_offset, global_w_size, 
                             local_w_size, 0, nullptr, nullptr);

kernel1先完成，然后kernel2被执行

使用


clCreateCommandQueue(context, device_id, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);

您可以同时执行多个不同的内核，但不能保证这样做。

但请注意，并非所有 OpenCL 实现都支持 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE。这也意味着您无法保证kernel1 将在kernel2 之前完成执行。如果kernel1 输出的任何对象需要作为kernel2 的输入，则可能会失败。

还可以创建多个命令队列并将命令排入队列，它们存在的原因是因为您希望解决的问题可能涉及主机中的一些（如果不是全部）异构设备。它们可以表示不共享数据的独立计算流，或者表示每个后续任务都依赖于前一个任务（通常是共享数据）的依赖计算流。但是，如果没有共享数据，这些命令队列将在设备上执行而不同步。如果数据是共享的，那么程序员需要使用 OpenCL 规范中的同步命令来保证数据的同步。

【讨论】：

【解决方案2】：

out_of_order 是否意味着异步

“乱序”队列意味着内核可能按照与队列不同的顺序执行（如果它们的事件/数据依赖关系允许的话）。它们也可能同时执行，但不一定。

另外，异步执行意味着除了执行重叠之外的其他东西（称为并行执行或并发）。异步执行意味着设备上的内核代码独立于主机代码执行——这在 OpenCL 中总是如此。

获得并发（执行重叠）的简单方法是在同一设备上使用 >1 个队列。这甚至适用于没有乱序队列功能的实现。它不保证执行重叠（因为 OpenCL 可以在比 CUDA 更多的设备上使用，并且在某些设备上您只是不能一次执行 >1 个内核），但根据我对大多数 GPU 的经验，您至少应该得到一些重叠。不过，您需要注意内核在单独队列中使用的缓冲区。

【讨论】：