【OpenCL基础 · 二 】OpenCL架构

---

前言

通过【OpenCL基础 · 一】因源，我们了解了OpenCL的起源和应用场景。在异构并行平台上，OpenCL可以协助开发者方便地使用计算资源，使得异构平台编程变得容易。

那么如何使用这个便捷的开发工具就是开发人员需要学习的了，我们需要从宏观到细节去理解OpenCL是如何运作的，才能进行实际的代码开发工作。这篇博文将介绍OpenCL的整体框架，展示它是如何利用异构平台进行控制、资源调度及运算的。包括OpenCL平台模型、执行模型和存储器模型。

参考：《OpenCL异构并行计算——原理、机制与优化实践》

---

一、OpenCL平台模型

平台模型描述了OpenCL是如何看待底层异构硬件平台的，下图是一个非常直观的展示：

这个图很核心哦~

一个异构平台包括多个、多类型处理器，OpenCL这个小可爱主观上将它们分为主机【Host】和与其相连的一个或多个OpenCL设备【Compute Device】，主机的功能是控制，OpenCL设备的主要功能则是计算。因此主机一般是包含X86或ARM处理器的计算平台，OpenCL设备可以是CPU、GPU、DSP、FPGA等处理器。

每个OpenCL设备包含一个或多个计算单元【Compute Unit，CU】，每个计算单元又由一个或多个处理单元【Processing Element，PE】组成，PE是执行计算的最小单元，不再向下细分。

在实际算法执行过程中，PE就是一个计算核，开发者可以在指定PE上运行代码。列举一个常见的使用场景，我们可以让大量的PE同时运行相同的代码段，不过它们的输入值不同，以此实现SIMT形式的并行运算效果。若将OpenCL设备（以下简称为设备）看作多核处理器的话，那么每个PE就像是一个核。在执行模型中会有详细介绍。

二、OpenCL执行模型

执行模型描述了OpenCL是如何协调主机和设备进行计算的。

实现功能，必然要依靠代码。和平台模型对应，OpenCL 的代码区分为主机端程序和设备端的内核（kernel）程序。主机端程序运行在主机处理器上，它将内核程序和数据从主机提交至设备，设备接收到程序与数据后就在PE上进行计算。计算完成后，主机会从设备端取走运算结果。

OpenCL设备一般不具有IO口，因此数据需要从主机端获取，并且计算结果也要返还给主机。

从这样的工作流中，我们不难发现，内核程序会是一些计算量大、逻辑较简单的函数，而主机程序则是主要的控制程序，掌握整个流程和数据传输。OpenCL提供了丰富的API供主机端使用，开发者可以方便地按照规定流程进行调用，实现一个OpenCL工程。

可以看出内核程序是并行运算的核心，OpenCL定义了以下三种内核：

• OpenCL内核： 用OpenCL C语言编写，并用OpenCL C编译器编译的函数。所有OpenCL都支持该类型内核的执行。
• 原生内核： OpenCL之外创建的函数，比如主机端源码中定义的函数，在OpenCL中通过一个函数指针来访问。执行原生内核是一个可选功能。
• 内建内核： 被绑定至特定设备，不需要源码编译成程序对象的函数，这是OpenCL的扩展功能。常见用法是针对公开固定函数硬件，将它们关联到一个特定设备上，一些处理器厂家会针对自家产品提供一些实用的内建内核。

执行模型的框架如上所述，其具体实施就要依赖三个关键要素：上下文、命令队列和内核。下面对它们逐一介绍。

1.上下文

上下文【context】是主机为了内核程序的顺利运行而创建出来的一个执行环境，它包含以下内容：

• 设备【Device】：OpenCL平台包含的一个或多个设备，即指定可以运行内核程序的设备。
• 内核对象【Kernel】：要在设备上运行的内核函数。
• 程序对象【Program】：实现内核程序的源码和目标二进制码。
• 存储器对象【MemObject】：对主机和设备可见的存储对象，内核执行时操作这些对象的实例。

在实际开发中，主机端会通过API先创建一个context，将要使用的设备编号会作为参数传入，用于指示使用哪些设备完成之后的内核计算。之后主机在调用相应的API创建内核对象、程序对象等对象时，会把已经建好的context作为参数传进去，表示这些对象都建在这个传入的context中，以此形成对应的关系。

上下文和设备的对应关系：
1、上下文要求其中所关联的设备全部来自同一个平台；对于不同平台的设备，必须创建不同的上下文进行管理。
2、在一个平台中，一台设备可以同时关联到多个上下文中。
3、主机可以使用多个上下文来管理多个设备。

2.命令队列

在OpenCL架构中，主机通过命令队列与设备进行交互。主机或运行在设备上的内核可以提交命令给命令队列，命令会在队列中等待，直到调度到设备上被执行。一个命令队列在一个上下文中被创建，并关联到一个设备上。

命令可以分为以下三类：

• 【内核入队命令】将一个内核程序入队至关联了相同设备的命令队列中，等待设备执行。比如对应的API有：

cl_int clEnqueueNDRangeKernel (
cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
......

• 【存储器入队命令】与存储空间设置相关的命令入队。比如主机与设备数据传输的命令、内存空间映射的命令、内存释放命令等。比如对应的API有：

cl_int clEnqueueSVMMemcpy (
cl_command_queue command_queue, cl_bool blocking_copy, void *dst_ptr, const void *src_ptr, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)；
cl_int clEnqueueMigrateMemObjects ( 
cl_command_queue command_queue, cl_uint num_mem_objects, const cl_mem *mem_objects, cl_mem_migration_flags flags, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
cl_int clEnqueueWriteBufferRect (
cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, const size_t *buffer_origin, const size_t *host_origin, const size_t *region, size_t buffer_row_pitch, size_t buffer_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
......

• 【同步命令】对命令队列中需要执行的命令施加执行顺序约束。

命令是异步方式执行，即主机或设备内核提交了命令后不用等命令执行完成，可以直接继续工作。

命令队列中的命令执行顺序包括：

• 【按序执行】按先入先出的顺序执行命令。
• 【乱序执行】命令以任意顺序执行，通过显式的同步点或显式事件依赖项来约束顺序。

一般情况下是按序执行，所有OpenCL平台均支持。

3.内核的执行——NDRange

异构平台上可能有多个计算设备，每个设备又有数量众多的核，如何将计算任务分配到各个核上是一个需要考虑的管理问题。OpenCL采用编号的方式来管理任务（即内核程序）实例，每个实例就是在一个处理单元（还记得吗，是平台模型中的PE哦==）上运行的。

主机提交一个内核到设备上运行时，OpenCL会创建一个N维整数索引空间，N可以取1、2或3。这个N维的网格就称为NDRange。
内核程序在设备上执行时，为了利用多核并行性，OpenCL会创建出多个内核实例，一个实例被分配到一个PE上同时运行，它们执行的计算是相同的。我们把实例称为一个工作项（work-item），在NDRange中一个实例就对应网格中的一个点，这个点是用坐标标识的。比如下图是一个二维NDRange，其中最小的一个方块就表示一个工作项，涂黑的那个工作项对应的全局坐标【全局ID】 值就是（6,5）。一般坐标从（0,0）开始。

图中的G_X和G_Y值描述了全局空间的大小，也表明了分配给当前内核程序使用的PE数量。

为了更多的管理需求，OpenCL将多个工作项组织为一个工作组（work-group），工作组横跨了全局索引空间，提供了对索引空间的粗粒度分解。上图中的NDRange就被划分成9个工作组，每个工作组包含了4x4个工作项。工作组也有自己的索引ID，上图正中间的工作组的ID就是（1,1）；同时工作项会被定义一个组内ID（局部ID），来表示工作项在其所属工作组中的位置。
通常情况下，我们使用工作项的全局ID就足够了。

通过上述内容，我们可以进行合理推测，N取1、2、3时分别适配的计算类型为，一维数组计算、二维矩阵计算和三维矩阵计算（三维的情况在卷积神经网络中就很常见了）。

三、OpenCL存储器模型

1.存储器区域

存储是运算过程中必不可少的部分，用于存储运算数据和代码。OpenCL异构平台由主机端和设备端构成，存储器区域包含了主机与设备的内存， 根据功能定义以下几种不同的存储器区域：

• 【主机内存】主机可以直接使用的内存，通过OpenCL API或共享虚拟存储器接口，实现存储器对象在主机和设备间的传输。
• 【全局存储器】这个存储器区域允许上下文中任何设备中所有工作组的所有工作项读写。
• 【常量存储器】属于全局存储器中的一块区域，在内核实例执行器件其保存的数据保持不变，对于工作项而言该存储器是只读的，主机负责对该存储器对象进行分配和初始化。
• 【局部存储器】该存储器区域对于工作组局部可见，它可以用来分配由该工作组中所有工作项共享的变量。
• 【私有存储器】该存储器区域是一个工作项的私有区域，其中的变量其他工作项不可见。

以上除了主机内存外，剩下的几种存储器模型都是属于OpenCL设备的。

OpenCL的内存模型如下图所示，其中全局存储器和常量存储器可以在一个上下文内的一个或多个设备间共享，OpenCL设备可能包含缓存来支持对这两个存储器的高效访问。一个OpenCL设备关联自己的局部存储器和私有存储器。

2.存储器对象

全局存储器中的数据内容通过存储器对象来表示，一个存储器对象就是对全局存储器区域的一个引用。存储器对象可以分为三种类型：

• 【缓冲 buffer】内核可用的一个连续存储区域，可以将内建数据类型、矢量数据类型等映射到缓冲区，内核通过指针来访问缓冲区。
• 【图像 image】用于存储标准格式的图像。它是一个不透明的数据结构，使用OpenCL API来管理。
• 【管道 pipe】管道存储器是数据项有序的队列，同一时刻仅有一个内核实例（即工作项）可以向一个管道内写数据，同一时刻，也仅有一个内核实例可以从一个管道中取数据。

3.主机与设备的数据交互

在计算过程中，主机和设备必然要进行数据交互，OpenCL提供三种交互方式：

• 【读/写/填充】显示的数据传输方式，需要主机将对应的命令入队，实现数据读写操作。
• 【映射和解映射】映射是指允许主机将一个存储器区域映射到主机内存的一个区域上，主机就可以直接访问到。在实际操作中，主机先将映射命令入队，实现存储器区域的映射，再对该区域进行读写；完成后再入队一个解映射命令，使得内核可以安全地读写缓冲。
• 【拷贝】将存储器对象在两个缓冲间拷贝，这两个缓冲可以停留在主机或设备上。

OpenCL 2.0提供了一种新的主机-设备交互方式，称为共享虚拟存储器（Shared Virtual Memory，SVM），有以下三种类型：

• 【粗粒度SVM】共享发生在OpenCL缓冲存储器对象 区域的粒度；
• 【细粒度SVM】共享发生在OpenCL缓冲存储器对象里独立地以字节加载/存储的粒度；
• 【细粒度系统SVM】共享发生在主机内存内任何地方独立地以字节加载/存储的粒度。

简单来说，粗粒度对应区域，细粒度对应字节，系统对应主机内存。

---

总结

本文对OpenCL的整体架构进行了介绍，实际构建项目时的具体步骤均依托于这些模型进行，因此先在脑海中建立起大框架有助于后续的开发学习。
我也是初学者，内容可能存在理解有误的地方，欢迎大家指正~