OpenCL 2.0 异构计算 [第三版] (中文)
  • Introduction
  • 序言
  • 第1章 简介异构计算
    • 1.1 关于异构计算
    • 1.2 本书目的
    • 1.3 并行思想
    • 1.4 并发和并行编程模型
    • 1.5 线程和共享内存
    • 1.6 消息通讯机制
    • 1.7 并行性的粒度
    • 1.8 使用OpenCL进行异构计算
    • 1.9 本书结构
  • 第2章 设备架构
    • 2.1 介绍
    • 2.2 硬件的权衡
    • 2.3 架构设计空间
    • 2.4 本章总结
  • 第3章 介绍OpenCL
    • 3.1 简介OpenCL
    • 3.2 OpenCL平台模型
    • 3.3 OpenCL执行模型
    • 3.4 内核和OpenCL编程模型
    • 3.5 OpenCL内存模型
    • 3.6 OpenCL运行时(例子)
    • 3.7 OpenCL C++ Wapper向量加法
    • 3.8 CUDA编程者使用OpenCL的注意事项
  • 第4章 OpenCL案例
    • 4.1 OpenCL实例
    • 4.2 直方图
    • 4.3 图像旋转
    • 4.4 图像卷积
    • 4.5 生产者-消费者
    • 4.6 基本功能函数
    • 4.7 本章总结
  • 第5章 OpenCL运行时和并发模型
    • 5.1 命令和排队模型
    • 5.2 多命令队列
    • 5.3 内核执行域:工作项、工作组和NDRange
    • 5.4 原生和内置内核
    • 5.5 设备端排队
    • 5.6 本章总结
  • 第6章 OpenCL主机端内存模型
    • 6.1 内存对象
    • 6.2 内存管理
    • 6.3 共享虚拟内存
    • 6.4 本章总结
  • 第7章 OpenCL设备端内存模型
    • 7.1 同步和交互
    • 7.2 全局内存
    • 7.3 常量内存
    • 7.4 局部内存
    • 7.5 私有内存
    • 7.6 统一地址空间
    • 7.7 内存序
    • 7.8 本章总结
  • 第8章 异构系统下解析OpenCL
    • 8.1 AMD FX-8350 CPU
    • 8.2 AMD RADEON R9 290X CPU
    • 8.3 OpenCL内存性能的考量
    • 8.4 本章总结
  • 第9章 案例分析:图像聚类
    • 9.1 图像聚类简介
    • 9.2 直方图的特性——CPU实现
    • 9.3 OpenCL实现
    • 9.4 性能分析
    • 9.5 本章总结
  • 第10章 OpenCL的分析和调试
    • 10.1 设置本章的原因
    • 10.2 使用事件分析OpenCL代码
    • 10.3 AMD CodeXL
    • 10.4 如何使用AMD CodeXL
    • 10.5 使用CodeXL分析内核
    • 10.6 使用CodeXL调试OpenCL内核
    • 10.7 使用printf调试
    • 10.8 本章总结
  • 第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度
    • 11.1 简要介绍现状
    • 11.2 简单介绍C++ AMP
    • 11.3 编译器的目标 —— OpenCL 2.0
    • 11.4 C++ AMP与OpenCL对比
    • 11.5 C++ AMP的编译流
    • 11.6 编译之后的C++ AMP代码
    • 11.7 OpenCL 2.0提出共享虚拟内存的原因
    • 11.8 编译器怎样支持C++ AMP的线程块划分
    • 11.9 地址空间的推断
    • 11.10 优化数据搬运
    • 11.11 完整例子:二项式
    • 11.12 初步结果
    • 11.13 本章总结
  • 第12章 WebCL:使用OpenCL加速Web应用
    • 12.1 介绍WebCL
    • 12.2 如何使用WebCL编程
    • 12.3 同步机制
    • 12.4 WebCL的交互性
    • 12.5 应用实例
    • 12.6 增强安全性
    • 12.7 服务器端使用WebCL
    • 12.8 WebCL的状态和特性
  • 第13章 其他高级语言中OpenCL的使用
    • 13.1 本章简介
    • 13.2 越过C和C++
    • 13.3 Haskell中使用OpenCL
    • 13.4 本章总结
Powered by GitBook
On this page
  • 7.7.1 原子访问
  • 7.7.2 栅栏

Was this helpful?

  1. 第7章 OpenCL设备端内存模型

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支持原子操作。那么就来介绍一下内存序和内存域,这里我们简单的回顾一下原子操作。

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

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

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

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

C atomic_fetch_<key>_explicit(volatile A *object, M operand, memory_order order)

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

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

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

bool
atomic_compare_exchange_strong_explicit(
  volatile A *object,
  C *expected,
  C desired,
  memory_order success,
  memory_order failure,
  memory_scope scope)

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

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

#define ATOMIC_VAR_INIT(C value)

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

global atomic_int sync = ATOMIC_VAR_INIT(0);

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

local atomic_int sync;
if (get_local_id(0) == 0){
  atomic_init(&sync, 0);
}
work_group_barrier(CLK_LOCAL_MEM_FENCE);

7.7.2 栅栏

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

void
atomic_work_item_fence(
  cl_mem_fence_flags flags,
  memory_order order,
  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] 多线程/多工作项访问同一变量会产生数据竞争。

Previous7.6 统一地址空间Next7.8 本章总结

Last updated 6 years ago

Was this helpful?