OpenCL 同步函数

通过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并行程序的读者,这类错误是比较常见的。

赞(0)
未经允许不得转载:极客笔记 » OpenCL 同步函数
分享到: 更多 (0)

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址