极客笔记通过OpenCL执行模型我们知道，对于一个内核函数，会有多个工作组参与计算。每个工作组中会有多个工作项参与计算。对于这些工作组和工作项是否有同步机制呢？
在OpenCL中定义了一个相对宽松的同步机制。在这个同步机制中，多个工作组之间没办法同步；而在同一个工作组内的工作项，可以通过内置函数实现同步；而不同工作组内的工作项是没办法通过内建的API进行同步。
一个工作组内的工作项可以通过如下函数实现同步：
void work_group_barrier(cl_mem_fence_flags flags)
void work_group_barrier(cl_mem_fence_flags flags,
memory_scope scope)
对于如上函数，在一个工作组内，在计算单元上执行内核的所有工作项在越过这个函数继续执行之前，必须执行这个函数,也可以看做这是所有工作项的一个同步点。执行内核的工作组中所有工作项都必须执行到这个函数。
如果这个函数在一个条件语句中，且工作组内任何一个工作项进入了这个条件语句并执行这个函数，那么工作组内所有工作项都必须执行此函数。
如果这个函数在一个循环中，对于每一次循环迭代，工作组内任何工作项能够越过这个函数继续执行之前，所有工作项都必须执行这个函数。
参数scope指定工作组内工作项对存储空间的访问是否对工作组内的工作项、设备或所有SVM设备可见的。如果没有这个scope，则默认为memory_scope_work_group。
参数flags指定了何种存储空间在参数scope指定的恰当存储范围内变得可见，可取的值为：
- CLK_LOCAL_MEM_FENCE：局部存储器访问对工作组内所有工作项是可见的。
-
CLK_GLOBAL_MEM_FENCE：全局存储器访问对工作组内所有工作项是可见的。
-
CLK_IMAGE_MEM_FENCE:图像存储器访问对工作组内所有工作项是可见的。
下面列举一些使用work_group_barrier不正确的用法。
kernel void readlocal(global int ＊g, local int ＊shared)
{
int id = get_global_id(0);
……
if(id 〈 3)
{
work_group_barrier(CLK_LOCAL_MEM_FENCE)
……
}
……
}
在上述例子中，假定每一个工作组中有16个工作项。在这种情况下，不是所有的工作项都会执行work_group_barrier（）。对于这种情况，结果是未知的，可能会造成硬件中的一个死锁。
kernel void MemcpyCheck(global int ＊a.global int ＊out)
{
int temp;
int id = get_global_id(0);
a[id] = a[id] + id;
barrier(CLK_GLOBAL_MEM_FENCE);
temp = (a[id] + a[id + 1] + a[id + 2]);
out[id] = temp;
}
在上述代码中，假定全局工作项大小为8，有两个工作组。对于id=3的工作项计算temp值时，会用到另一个工作组中a[4]和a[5]的值。而工作组之间无法实现存储器一致性，所以对于上述情况，结果是可能某次计算结果正确，下次计算结果又是错误的。对于刚开始编写OpenCL并行程序的读者，这类错误是比较常见的。