OpenCL死锁可能性

我正在使用全球原子来同步OpenCL中的工作组。

因此内核使用如下代码

... global volatile uint* counter;

if(get_local_id(0) == 0) {
    while(*counter != expected_value);
}
barrier(0);

等到 counter 成为 expected_value .

在另一个地方

if(get_local_id(0) == 0) atomic_inc(counter);

从理论上讲,如果所有工作组都在并发运行,那么该算法应该始终有效。但是,如果一个工作组在另一个工作组完全完成后才开始,那么内核可能会死锁。

在CPU和GPU(NVidia CUDA平台)上,它似乎总是工作的,有大量的工作组(超过8000个)。

对于算法来说,这似乎是最有效的实现。(它对2D缓冲区中的每一行进行前缀求和。)

OpenCL和/或NVidia的OpenCL实现是否保证了这一点?