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)
  3. 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] 多线程/多工作项访问同一变量会产生数据竞争。