OpenCL的并发和执行模型

《OpenCL异构计算》 

第五章 OpenCL的并发与执行模型 读书笔记

Outline

• kernel, work_item, workgroup 概念, 执行域的概念 
• OpenCL 同步： kernel, fence 和 barrier 
• 队列与全局同步 
• 主机端内存模型 
• 设备端内存模型

kernel, work_item, workgroup 概念, 执行域的概念

基本概念
kernel 实例（线程）= work_item
NDRange 定义 并行执行空间。
一个 work_item 定义 one sliver of a large parallel execution space。
一个 work_item 运行在一个 PE (Processing Element) 上。
一个 PE 可处理多个 work_item （在一个 kernel 执行过程中）。
work-group 是 固定大小数目的 work_item 的集合。
一个 work-group 运行在一个 CU (Compute Unit) 上。

同步问题
问题1 work-item 之间是否定义同步？如何同步？
答：在运行中，work_item 之间相互独立，OpenCL 未定义 （任意两个）work_item 之间的同步。这样做，是为了保证 OpenCL 执行模型的可扩展性。允许一个 work-group 内部的本地同步。

通信问题
问题1 work-group 内部是否允许通信？
答：一定程度允许。但被限制，以提高可扩展性。Local Store？？？

重要接口函数

• uint get_work_dim(): Returns the number of dimensions in use in the dispatch.
• uint get_global_size(uint dimension): Returns the global number of work items in the dimension requested.
• uint get_global_id(uint dimension): Returns the index of the current work item in the global space and in the dimension requested.
• uint get_local_size(uint dimension): Returns the size of workgroups in this dispatch in the requested dimension.
• uint get_local_id(uint dimension): Returns the index of the current work item as an offset from the beginning of the current workgroup.
• uint get_num_groups(uint dimension): Returns the number of workgroups in the specified dimension of the dispatch. This is get_global_size divided by get_local_size.
• uint get_group_id(uint dimension): Returns the index of the current work group. That is, the global index of the first work-item in the workgroup, dividing by the workgroup size.

OpenCL 同步： kernel, fence 和 barrier

OpenCL 同步的执行方式要点如下：

• 一个 work-item 的读操作和另一个 work-item 的写操作没有任何顺序保证。
• OpenCL 采用宽松的同步模型和内存一致性模型。
• 为了获得松弛一致性，The solution is that OpenCL explicitly defines synchronization points where the programmer knows with certainty what the state of some part of the system is and can rely on that information to obtain expectations of behavior.
• 在 GPU 上，停止一个 wavefront 的运行不会释放它所占用的资源。因此，在 GPU 上，如果采用操作系统那种可以随时停止或者调度线程执行的方式，则可能出现这样一种状况：一个已经调度在设备上的 wavefront 需要得到另外一个还没有调度到设备上的 wavefront 释放信号量才能开始运行。为避免这种情况，全局同步只定义在 kernel 执行的边界处。
• 不同 work-group 的两个 work-item，并没有定义方法来确保它们的执行顺序。
• 数据共享（通信），主要是实现了同一个work-group 的 work-item 之间，通过 local memory 来进行数据共享。
• OpenCL 定义了一个 work-group 内部的同步操作。通过 barrier 函数。 barrier(CLK_LOCAL_MEM_FENCE);
• 在 kernel 代码中， 如果某个分支调用了 barrier 来同步，则必须保证同一个 work-group 中的所有 work-item 都会执行该分支。如果有 work-item 未执行包含该 barrier 调用的分支，则 barrier 的行为是不确定的。在很多设备上，会导致死锁。

任务队列与全局同步

• OpenCL内存一致性
• 事件
• 事件回调
• 命令 barrier 与 mareker

队列
问题1 能够加入 OpenCL 队列的命令有哪些？

• kernel 执行命令
• 内存操作命令
• 同步命令

问题2 入队的方式是？

答：kernel 执行命令和同步命令都是异步入队。

全局同步

特点： 一个命令的完成只能在同步点得到保证。

问题1 有哪些同步点？

• 调用 clFinish
• 等待一个特定事件的完成
• 执行一个阻塞访存操作

内存一致性

问题1 同步命令的作用是？

答：保证在同步点上， work 的同步 和 内存一致的视图。这意味着命令之间的一致性，也即通信的正确性在一个有序队列的各命令间或者在产生事件和等待该事件的命令之间是有保证的。

问题2 主机 host 与 设备 device 之间的一致性如何保证？

答：只能通过阻塞操作达到目的。问题1中一致性保证是在运行时保持内存对象的一致性，而对于主机端 API 来说是不可见的。

问题3 设备之间的一致性如何保证？

答：内存对象是与上下文相关而不是与设备相关，所以在数据被共享和发生相应的事件时，OpenCL 运行时要保证这些对象在不同设备之间的一致性。

Data is moved from one device to another such that if a kernel is to be executed on a second device, any results generated on the first will be available when necessary. The completion of an event on the first data structure is the guarantee that the data is OK to move and no separate buffer copy operation is

needed. 这2句话如何理解？

事件

问题1 构建命令队列的参数问题

cl_command_queue clCreateCommandQueue(	cl_context context,
	cl_device_id device,
	cl_command_queue_properties properties,
	cl_int *errcode_ret)

cl_command_queue_properties 设置可以是

Command-Queue Properties	Description
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE	Determines whether the commands queued in the command-queue are executed in-order or out-of-order. If set, the commands in the command-queue are executed out-of-order. Otherwise, commands are executed in-order.
CL_QUEUE_PROFILING_ENABLE	Enable or disable profiling of commands in the command-queue. If set, the profiling of commands is enabled. Otherwise profiling of commands is disabled. See clGetEventProfilingInfo for more information.

直接使用事件是与主机同步的第三种方法。OpenCL 通过构建一个任务图来指明各事件之间的依赖关系。也可以调用 clWaitForEvents 来指定一个 wait 事件。

问题2  如何管理多个设备上的命令队列
答：可以将多个队列映射到相同的设备上，实现不同命令之间的重叠执行或者 命令与 host-device 通信操作之间的重叠执行。在一个有多个设备的系统中，每个设备都需要自己的命令队列。在关联了多个设备的同一个上下文中，可为每个设备创建自己的命令队列。而且可以使用 事件 来实现命令间的同步。如果多个设备各自有各自自己的上下文，则不能使用事件进行同步。此时，只能同步依靠 clFinish 实现，或是 显式的在对象间拷贝数据。

问题3 OpenCL异构编程框架下有多少中执行模型？

• 两个或多个流水执行

• 多个独立并行

问题4 除同步外，事件还有哪些用途？

答：可用于查询 OpenCL 命令执行相关的错误状态或者程序剖分数据。通过 getInfo 函数实现。

问题5 用户事件是什么？如何使用？

答：用户自定义的时间。通过 clCreateUserEvent 创建。用于将依赖于其他任意任务的命令入队。用户事件可以像其它事件一样传给入队函数，但其执行状态是用户显式的给出的。

事件回调

问题1 事件回调有啥用？

答：同样是定义回调函数，在事件到达指定状态时触发。事件回调可 用于将新命令入队 或 用于调用主机端函数。回调函数应该是轻量级的。clSetEventCallback 函数的调用位置很重要。必须位于产生其所需的有效事件的函数调用之后。

问题2 事件回调的使用有啥讲究？

答：必须非常小心。

• 为同一个事件多个执行状态注册的回调函数，在命令状态改变时，不保证这些回调函数按照正确的顺序调用执行。
• 回调函数是线程安全的，会被异步调用。
• 当调用开销大的系统调用或阻塞式的 OpenCL API （如 clFinish）时，回调函数的行为是未定义的。

这种使用事件处理的机制的好处是应用级行为和同步操作可在不同厂商提供的 CPU 和 GPU 上用统一的方式进行处理。这样就将设备相关的调优局限于 计算核 中。

问题3 事件回调之外，还有啥替代方案？

答：可使用 native kernel。将普通 C 函数入队。

命令 barrier 与 marker

问题1 不同命令队列的同步除事件外，还有什么方法？

答：还有一种与 work-group 内部同步类似的方法，就是使用 barrier。marker 与 barrier 类似，但是不阻塞执行。当有事件在等待命令队列中的一些命令完成时，marker 非常有用，但是不会影响到其他命令的执行。还有一个同步原语是 waitForEvents，其行为与 marker 相反。不等待所有任务完成，只等待特定任务完成。

主机端内存模型

• buffer 对象
• image 对象

OpenCL 内存对象定义在上下文，而不是设备上。通常 buffer 上的操作是没有必要涉及到具体硬件设备的。保证数据在正确的时间到达正确的位置是运行时的工作。

buffer 对象
从 CPU 的意义来看，buffer 对象是一维数组，类似于 C 程序中 malloc 开辟的内存。是连续的，RMA。OpenCL 还定义了 sub-buffer 对象上下文，允许把一个单独的 buffer 切分成许多可以重叠的较小的 buffer，而这些 sub-buffer 的读写复制与 buffer 的操作是相同的。sub-buffer 的重叠，以及他与父buffer 的合并可能会造成别名问题。

imgae 对象
对图像数据进行了优化。体现在：

• GPU 的 cache 层次结构以及数据流结构都是为了优化图像数据类型的访问而设计的。
• GPU 驱动程序优化了数据布局以提供对硬件的高效访存的支持，特别是使用二维访存模式时。
• image 访存支持复杂的数据类型转换以允许数据以广泛的压缩格式进行存储。

不像 buffer，image 不同内存对象之间的位置关系对开发人员来说是不可见的。image 数据结构不仅对开发人员不可见，对 kernel 代码亦完全不可见。只能通过专门的访问函数来存取。

image 数据格式由通道序列和通道类型组成。支持运行时系统和硬件优化。image 对象在设备端不能通过指针进行直接访问，也不能在同一个 kernel 中进行读操作 和 写操作。消除了锯齿数据的可能性，可以进行安全缓存。

Z序映射方式增加图像数据局部性

设备端内存模型

• 设备端宽松的内存一致性模型
• 全局内存
• 本地内存
• 常量内存
• 私有内存

设备端宽松的内存一致性模型

• work-item 内部内存操作有可预测的顺序：即任意两个对同一地址的读和写是不会被硬件或者编译器重新排序的。
• 在同一个 work-group 的不同 work-item 之间，只在 barrier 操作处保证其内存一致性。
• 在 work-group 之间，在 kernel 执行完成前，不保证内存一致性。

为了使位于同一个或者不同 work-group 中的 work-item 之间开展某种程度的通信，OpenCL 定义了一组 fence 操作。但是即使使用这些 fence 操作也不能保证 work-item 的执行顺序。

• read_mem_fence( cl_mem_fence_flags flags )
• write_mem_fence( cl_mem_fence_flags flags )
• mem_fence( cl_mem_fence_flags flags )

保证不同 work-item 之间内存操作正确的另一个方法是 原子操作。可以在不影响其他 work-item 的情况下，保证数据读或读写以及其他操作的正确性。（浮点数原子操作设计复杂）OpenCL 目前仅定义了 整数原子操作，如 atomic_add 等。