如何干净地退出 OpenCL 代码答案

【问题标题】：How to cleanly exit OpenCL code如何干净地退出 OpenCL 代码
【发布时间】：2016-12-28 13:36:10
【问题描述】：

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

【问题讨论】：

标签： opencl pyopencl

【解决方案1】：

如果你需要类似的东西

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 作为错误缓冲区来直接在主机内存上读取/写入错误代码，而不是使用设备内存。这可能比设备内存的性能低，因为错误缓冲区不会被缓存并且会远离设备，因此可能是 5GB/s 的 pci-e 带宽瓶颈，而不是 5TB/s 的设备内存（假设广播到所有内核单周期, 对于最新的显卡）

【讨论】：

谢谢，这是一个很好的答案。是否保证所有线程都命中第二个（写入）mem_fence？如果不是，那对我来说似乎是令人讨厌的死锁的根源。
mem fence 仅用于简单读取。使用原子读取，您不需要。因为它们是原子的。您应该测试原子和非原子版本并在出现错误时进行基准测试（或者您故意犯错误），这样只会发生一个错误，然后不会出现其他错误，因此原子写入性能损失只发生一次。
您还可以将第二个内核加入队列以恢复任何“错误后计算”线程效果。
我理解你的两句话。但是，在发生错误的情况下，并非所有线程都有可能到达内存围栏。对我来说这已经足够好了，不过我还是想把它作为评论留给更多的读者。
这就是为什么我说只在线程组级别保存周期。我的意思是，如果在一个线程组中保存了一些周期，那么同一组中的其他线程的机会很高。如果不是，那么其他人也不会，但其他组可能已经为他们自己的线程捕获错误。它是关于在第一次写入后立即阅读的。可能有几个写入，但它们肯定都在进行中，并且受到 gpu 的限制。您还可以检查线程中每个 x 计算迭代的错误以更快地退出，但前提是您可以隐藏其延迟以不影响性能

【解决方案2】：

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

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

【讨论】：