如何干净地退出 OpenCL 代码

Question

我正在寻找一种简单明了的方法来告诉我的主机，OpenCL 中的某些代码会导致错误，它应该放弃进一步的工作。我（认为）知道，try、catch 或 assert 在 OpenCL C 中不起作用。此外，内核必须定义为非返回函数，因此简单的错误代码返回也不在图像中。我唯一的想法是在主机和内核之间传递一个 cl_mem 对象，并在内核入队或启动之间检查它的值，这以某种方式强制执行了一种非常强大的序列化。有没有更好的主意，也许使用事件？

Answer 1

为了不引起序列化，使用异步解决方案。每个内核都有一个 cl_mem 对象，如果它想告诉主机端关闭，它会写入一个标记值。当主机稍后检测到这一点时，它会停止排队工作。因为这是异步的，一些额外的工作项可能会在写入和读取之间排队，您的系统将需要处理这种情况。您可能需要 cl_mem 个对象的环形缓冲区，因为您需要在主机上读取或映射它们以检查它们的内容，并且您不想阻止任何内核来执行此操作。

在 OpenCL 2.x 中，您或许可以使用管道进行此通信，但我没有使用过它们，因此无法提供有关此解决方案的更多详细信息。

Answer 2

如果你需要类似的东西

for(i 0 to N)
{
    do work (i)
    error ? break;
}

以并行方式，

int threadId=get_global_id(0);

// a broadcast read(for a new gpu) so no performance hit
mem_fence(CLK_GLOBAL_MEM_FENCE_READ)
if(error[0]==0) // or use atomic_add(&error[0],0) to read atomically(when total number of threads is low like thousands)
{
     do work (threadId);
     error? atomic_add(&error[0],errCode)
     mem_fence(CLK_GLOBAL_MEM_FENCE_WRITE)
}

所以至少你在线程组级别保存了周期，如果它们在原子错误写入之后启动，应该让后面的线程快速完成。原子操作很慢，但错误处理应该会降低它的重要性，对吧？此外，根据设备类型和驱动程序，在原子写入和适当的非原子读取之间可能至少需要数千个线程，因此对于一百万个线程，它可能是有效的，但对于一千个线程，您应该使用原子读取（原子添加值为零）因此每个线程都会在实际工作开始之前添加额外的周期，但至少它的延迟可能会被繁重的计算所隐藏。

如果您有多个设备需要相互通知错误，您应该使用 USE_HOST_PTR 错误缓冲区直接在主机内存上 read/write 错误代码，而不是使用设备内存。这可能比设备内存性能低，因为错误缓冲区不会被缓存并且会远离设备，因此可能是 5GB/s 的 pci-e 带宽瓶颈而不是 5TB/s 的设备内存（假设以单周期广播到所有内核, 对于最新的显卡)

如何干净地退出 OpenCL 代码

How to cleanly exit OpenCL code

opencl

pyopencl