将Barrier移至OpenCL内核if-else分支内的优化是否存在问题？

警告（伪代码）

假设我们有如下内核：

def kernel(array):
   a = get_global_id(0)
   if a > 1:
      array[0] = 10
   barrier(LOCAL_MEM_FENCE)

将代码优化为如下形式是否错误？

def kernel(array):
   a = get_global_id(0)
   if a > 1:
      array[0] = 10
      barrier(LOCAL_MEM_FENCE)
   else:
      barrier(LOCAL_MEM_FENCE)

尽管该优化看似合理，却会导致错误结果与数据竞争。针对Intel x86架构下的OpenCL Barrier实现，错误原因如下：

• OpenCL Barrier的核心规则：同一工作组内的所有线程必须到达同一个同步点调用barrier，才能继续执行后续逻辑。如果线程在不同代码路径调用barrier，会破坏同步语义，引发执行异常。
• 数据竞争的触发逻辑：原始代码中，所有线程先完成各自的条件执行（a>1的线程修改array[0]，其余线程无操作），再统一在barrier处等待，确保所有线程对array[0]的修改操作全部完成后，才会推进后续流程。而优化后的代码中，a>1的线程修改array[0]后立即调用barrier，a≤1的线程直接调用barrier，这会导致：
  • 部分线程还在修改array[0]时，另一部分线程可能已经通过barrier进入后续逻辑，直接读取未完成修改的array[0]，引发数据竞争。
  • OpenCL的barrier实现依赖工作组内所有线程到达同一同步点，不同代码路径的barrier会被识别为不同的同步点，可能导致线程永久等待（死锁）或同步逻辑完全失效，彻底破坏工作组的执行一致性。
• x86架构的局限性：x86架构的内存模型虽相对宽松，但无法弥补OpenCL同步语义的错误。不管底层架构如何，都必须遵守“同一工作组所有线程到达同一barrier点”的规则，否则依然会出现数据可见性问题。

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