7.7 内存序

对于任何编程语言来说,内存序对于内存模型来说十分重要,需要用一定的顺序来保证线程得到的是期望的结果。当我们使用多线程和共享数据时,内存一致性模型能帮助保证线程得到的是正确的结果。OpenCL需要提供可移植化的高度并行代码,那么内存模型在正式发布的标准文档中就尤为重要。

之前我们提到过,执行内核中的所有工作项都可以访问全局内存上的数据。另外,在同一工作组的工作项可以共享局部内存。直到现在,我们在处理内存时,更多的是使用OpenCL的“松散型”一致模型。对于全局内存,我们没有使用更加复杂的内存模型,并且默认让不同工作组中的工作项更新不同位置的全局内存。关于更新内存对象,我们不能认为和该对象有关的命令状态为CL_COMPLETE时,内存对象更新完成。在实践中,简单的内存模型覆盖了绝大多数的内核。第4章中,我们看到这种内存一致性模型也能支持直方图和卷积相关的应用。

近几年中,C/C++和Java都在支持“获取-释放”操作,为了就是在不使用锁的情况下,进行线程同步。这些操作有助于并行应用中核心代码的处理。OpenCL 2.0在基于C11标准的基础上,也支持“获取-释放”操作。另外,OpenCL的开发者可以将这种问题解决方式扩展到其他类型的应用上,使得支持OpenCL的高级语言可以更加容易进行线程同步。

对于编程者来说,顺序一致模型是最为直观的内存模型。如果系统由顺序一致模型实现,那么各个处理器上的内存操作将会按照程序的执行顺序进行,并且所有处理器上的操作顺序一致。不过,对于顺序一致性模型很难进行优化,因为对程序的正确性没有影响(比如,由编译器重拍指令顺序或在处理器上使用一个存储内存块)。因此,松散的内存模型需要和顺序一致性模型输出一样的结果才算正确。松散序一致性模型需要硬件和软件遵循某些规则,才能得到正确的结果。对于编程者来说,需要花点时间告诉硬件,数据在什么时候才能对其他线程可见。

不过有时同步操作,会成为程序性能的瓶颈。因此,OpenCL提供对应不同类型选项,供同步操作使用(供编程者指定),每种选项的粒度都有不同的粒度和范围。这些选项称为内存序(memory order)和内存域(memory scope)。

OpenCL提供三种不同程度的一致性顺序(从弱到强):松散、获取-释放和顺序。这些选项则由内存序选项指定:

  • 松散(memory_order_relaxed):这种内存序不会对内存序有任何的约束——编译器可以自由的对操作进行重排,包括后续的加载和存储操作。不过该方式可能会带来一些副作用,可能会造成结果错误。2.0之前的OpenCL标准中,原子操作就包含在松散的内存序中。因为缺少限制,所以编程者可能使用松散序获得最好的性能。
  • 获取(memory_order_acquire):获取操作和加载操作成对出现。当为同步操作指定该选项时,任何共享内存需要被其他执行单元(例如,其他工作项,或主机端线程)“释放”后才能进行存储。编译器需要将加载和存储操作移到同步操作之后。
  • 释放(memory_order_release):与获取操作不同,释放操作会和存储操作成对出现。当为同步操作指定释放序时,其会影响同步点之前的存储操作,使其操作对其他线程可见,并且在同步点之前的所有加载操作,必须在达到同步点前全部完成。编译器会将加载和同步操作移至同步点之前。
  • 获取-释放(memory_order_acq_rel):该内存序具有获取和释放的属性:其会在获取到其他执行单元的内存时,释放自己所获取的内存。这个选项通常用于“读改写”操作。
  • 顺序(memory_order_seq_cst):顺序一致性的内存序不存在数据数据竞争[1]。该内存序中,加载和存储操作的执行顺序和程序的执行顺序一致,这样加载和存储操作也就是简单的交错与不同的执行单元中。该选项要比memory_order_acq_rel更加严格,因为最后程序可以说是在串行执行。

当对全局内存进行同步时,指定内存序带来的性能开销,可能要超过计算时的开销。试想一个系统中具有多个设备,共享一个上下文,并且包含一个细粒度的SVM内存。当某个工作项使用释放型同步操作,那么就需要对所有设备上的工作项进行同步——如果不考虑算法的正确性,这将带来很大的性能开销。因此,对于很多操作来说,内存序参数会伴随一个内存域,其限制了指定执行单元可见操作的范围。

可以作为内存域指定的选项如下:

  • 工作项(memory_scope_work_item):指定内存序要应用到每个工作项中。这里需要对图像对象进行行操作。
  • 工作组(memory_scope_work_group):指定的内存序应用于工作组中的每个工作项。这个操作与栅栏操作相比,相当于一个轻量级的同步。
  • 设备(memory_scope_device):指定内存序用于某一个执行设备。
  • 所有设备(memory_scope_all_svm_devices):指定内存序应用于所有设备上的所有工作项,以及主机端(对细粒度SVM使用原子操作)。

与访问全局内存不同,访问局部内存不需要指定内存域(实际上指定了也会忽略)——局部原子操作通常具有默认内存域memory_scope_work_group。因为局部内存的访问只在同一工作组中存在,所以在外部设置memory_scope_device和memory_scope_all_svm_devices对于局部内存没有任何意义。

7.7.1 原子访问

本章开始时,我们说到OpenCL 2.0支持原子操作。那么就来介绍一下内存序和内存域,这里我们简单的回顾一下原子操作。

回想一下我们介绍过的原子操作,比如:加载、存储和“预取后修改”。我们展示一下“预取后修改”操作的函数声明:

  1. C atomic_fetch_<key>(volatile A *object, M operand)

这里的key可以替换成add、min或max。object参数为一种原子类型的变量的指针,operand代表操作数。返回值C,其类型是非原子的A类型。返回值时object地址中存储的值,这个返回值是没有进行操作前的数值。

上面的描述中,可以认为C/C++和OpenCL 2.0利用原子操作对内存序进行控制。因此,所有的原子操作都具有传入内存序和内存域的原子操作。例如,“预取后修改”函数具有以下函数声明:

  1. C atomic_fetch_<key>_explicit(volatile A *object, M operand, memory_order order)
  2. C atomic_fetch_<key>_explicit(volatile A *object, M operand, memory_order order, memory_scope scope)

通过这样的设计,就可以将线程间同步的任务交给原子操作来完成了。使用原子操作来做的同步的原因:其设置的标识可以让其他线程知道,应该在什么时候对某个区域的内存进行访问。因此,当一个线程想要看到其他线程所修改的内存时,通过对原子操作进行设置一些标志,然后等待共享数据释放。其他线程需要读取对应标志,在条件满足的情况下,将最后更新的数据拷贝到共享内存中。

另外,OpenCL除了支持加载、存储和“预取后修改”类型的原子操作之外,还支持交换、“比较后交换”和“测试后设置”类型的原子操作。这里列出一个“比较后交换”函数的声明:

  1. bool
  2. atomic_compare_exchange_strong_explicit(
  3. volatile A *object,
  4. C *expected,
  5. C desired,
  6. memory_order success,
  7. memory_order failure,
  8. memory_scope scope)

与之前看到函数声明不一样atomic_compare_exchange_strong_explicit()具有两个内存序参数——success和failure。这两个参数指定的是当比较操作成功和没成功时所使用到的内存序。编程者可以使用这种操作,来控制没有必要同步操作。比如,编程者将memory_order_relaxed传入failure,就是想在条件不成功的时候,不让工作项等待交换完成。

我们之前一直在讨论原子操作如何使用,并没有讨论如何对原子操作进行初始化。OpenCL C有两种方式对原子操作的操作域进行初始化。在程序范围内声明一个原子变量,可以使用ATOMIC_VAR_INIT()宏,该宏的声明如下所示:

  1. #define ATOMIC_VAR_INIT(C value)

这种方式初始化的原子对象是在程序域内进行声明,且分配在去全局地址空间内。例如:

  1. global atomic_int sync = ATOMIC_VAR_INIT(0);

原子变量在内核端需要使用非原子函数atomic_init()进行声明和初始化。注意,因为atomic_init()是非原子函数,但是也不能被多个工作项同时调用。也就是,初始化需要串行且同步的进行,例如下面代码所示:

  1. local atomic_int sync;
  2. if (get_local_id(0) == 0){
  3. atomic_init(&sync, 0);
  4. }
  5. work_group_barrier(CLK_LOCAL_MEM_FENCE);

7.7.2 栅栏

栅栏同步操作与内存的位置无关。虽然,实践中我们使用栅栏对工作组进行同步,但是我们从来没有说过栅栏操作如何使用内存序。在OpenCL C中,栅栏操作可以由atomic_work_item_fence()函数执行,其声明如下:

  1. void
  2. atomic_work_item_fence(
  3. cl_mem_fence_flags flags,
  4. memory_order order,
  5. memory_scope scope)

flags参数可以传入CLK_GLOBAL_MEM_FENCE, CLK_LOCAL_MEM_FENCE和CLK_IMAGE_MEM_FENCE,或将这几个参数使用“位或”(OR)的方式共同传入。共同传入的方式,与单独传入的效果是一样的。

很多系统上图像对象还是限制在非通用显示硬件上。OpenCL标准当然也注意到了这点,所以可向atomic_work_item_fence()传入CLK_IMAGE_MEM_FENCE,来保证图像对象在写之后才可读——即使对同一个工作项。如果有多个工作项要进行同步,然后可以读取同一工作组中前一工作项所写入图像中的数据,最后需要使用CLK_IMAGE_MEM_FENCE作为参数传入work_group_barrier()。另一种特别的方式,可以使用工作项栅栏对局部和全局内存的访问顺序进行统一控制。

之前我们介绍过,使用工作组栅栏和内存栅栏作为同步操作。理论上,这还是两种栅栏操作——就像出入栅栏一样。入栏就是指定标志和作用域释放栅栏。同样的,出栏也需要指定对应的标志和作用域。


[1] 多线程/多工作项访问同一变量会产生数据竞争。