opencl入门介绍

淡淡的烟草味﹌ 2022-07-31 16:25 656阅读 0赞

根据《OpenCL Overview》与《OpenCL Technical Overview》整理编写，由于第一次接触OpenCL，定会有部分的专业词汇翻译不当，敬请指正。

# **1.《OpenCL Overview》整理笔记** #

## **这是一个异构的世界** ##

**　　OpenCL**让程序员写**单一的可移植程序**，在一个**异构平台**上**使用所有的资源**。

## OpenCL使用 ##

1. 定义平台； 2. 在平台上执行代码； 3. 在内存中移动数据； 4. 编写(和编译)程序。

## OpenCL平台模型 ##

一个**主机(Host)** \+ 一个或更多的**计算设备**

- 每个**计算设备**由一个或多个**计算单元**组成

- 每个**计算单元**进一步分成一个或多个**处理元件**

![SouthEast][]

## OpenCL执行模型 ##

OpenCL应用程序运行在一个主机(host)上，主机提交工作到计算设备。

-**工作项(Work item)**：OpenCL上的基本工作单元

-**Kernel**：工作项目的代码，基本上是个C函数

-**程序(Program)**：收集kernel和其它函数(类似于动态库)

-**上下文(Context)**：工作项目执行的环境，包括设备和它们的内存与命令队列

应用程序队列kernel执行例子

- 顺序排队，队列到设备

- 顺序或乱序执行

![SouthEast 1][]

## OpenCL内存模型 ##

**私有内存(Private Memory)**

- 每个工作项

**局部内存(Local Memory)**

- 工作组内共享

**全局/常量内存(Global/Constant Memory)**

- 对所有工作组可见

**主机内存(Host Memory)**

- 在CPU上

**内存管理**是**显式**的，你必须搬移数据从**host -> global -> local ...**，然后返回。

![SouthEast 2][]

## kernel编程：OpenCL C语言 ##

**ISO C99的一个子集**

- 但是除了一些C99的特性，比如标准C99头文件，函数指针，递归，变长数组，和位域

**ISO C99的一个超集**

- 工作项和工作组

- 向量类型

- 同步

- 地址空间限定符(Address space qualifiers)

**包括一大组的内建函数**

- 图像处理

- 工作项的处理

- 专门的数学函数，等等

## kernel编程：数据类型 ##

**标量数据类型**

- char , uchar, short, ushort, int, uint, long, ulong, float

- bool, intptr\_t, ptrdiff\_t, size\_t, uintptr\_t, void, half (storage)

**　　图像类型**

- image2d\_t, image3d\_t, sampler\_t

**　　向量数据类型**

- Vector lengths 2, 4, 8, & 16 (char2, ushort4, int8, float16, double2, …)  
　　　　- Endian safe  
　　　　- 向量长度对齐  
　　　　- 向量操作  
　　　　- 内建函数

![SouthEast 3][]

## 建立程序对象 ##

**程序对象封装**

- 上下文

- 程序源代码/二进制文件

- 目标设备和构建选项的列表

**建立过程... 创建一个程序对象**

- clCreateProgramWithSource()  
　　　　- clCreateProgramWithBinary()

![SouthEast 4][]

## **OpenCL的同步：队列与事件(Queues & Events)** ##

事件可用于同步队列之间的kernel执行

例子：两个队列，两个设备

![SouthEast 5][]

## OpenCL总结 ##

![SouthEast 6][]

# 2. 《OpenCL Technical Overview》整理笔记 #

## **OpenCL设计需求** ##

**使用系统中的所有计算资源**

- GPU，CPU，和其它处理器作为对等体一起编程

- 同时支持数据/任务并行计算模型

**基于C的高效并行编程模型**

- 抽象底层硬件细节

**抽象是低层次、高性能的，但是设备可移植的**

- 容易上手，但主要正对开发专家

- 生态系统基础，没有中间件或“便携”函数

**在嵌入式，桌面和服务器系统范围内都可实现**

- HPC(高性能计算)，桌面，一个规范内的手持式型材(handheld profiles)

**　　推动未来的硬件需求**

- 浮点精度需求

- 同时适用消费和高性能计算应用

## OpenCL剖析 ##

**　　语言规范**

**　　　　-** 基于C的交叉平台编程接口

**　　　　-** ISO C99子集语言扩展，对开发者熟悉

**　　　　-** 良好定义的数值精度，IEEE 754规定最大错误的四舍五入行为

**　　　　-** 在线或离线编辑和构建的计算内核可执行文件

**　　　　-** 包含一个丰富的内建函数集合

**　　平台层API**

**　　　　-** 多种计算资源之上的硬件抽象层

**　　　　-** 查询、选择和初始化计算设备

**　　　　-** 创建计算环境和工作队列

**　　运行时API**

**　　　　-**执行计算内核

**　　　　-**管理调度、计算和内存资源

## 模型的层次结构 ##

**　** **　> 平台模型**

**　　> 内存模型**

**　　> 执行模型**

**　　> 编程模型**

## **OpenCL平台模型** ##

![SouthEast 7][]

**一个主机(Host) + 一个或更多的计算设备**  
　　　　- 每个**计算设备**由一个或多个**计算单元**组成  
　　　　　　- 每个**计算单元**进一步分成一个或多个**处理元件**

## OpenCL执行模型 ##

**OpenCL编程**

**- 内核**

**-** 执行代码的基本单元 --- 类似C函数

**-** 数据并行或任务并行

**- 主机编程**

**-** 计算内核和内部函数收集

**-** 类似于一个动态库

**内核执行**

**-** 主机程序通过一个叫做**NDRange**的索引空间调用内核

**-** NDRange =“N维范围”

**-** NDRange, 可以是1, 2, 或3维空间

**-** 一个内核实例在索引空间中的一个点称为一个**工作项**

**-** 工作项有来自索引空间的唯一的全局ID

**-** 工作项进一步组织成**工作组**

**-** 工作组具有一个唯一的工作组ID

**-** 工作项有在一个工作组中唯一的本地ID

## 内核执行 ##

**工作项的总数**：Gx X Gy

**每个工作组的大小**：Sx X Sy

**全局ID**可以从**工作组ID**和**本地ID**计算得到

![SouthEast 8][]

## 上下文和队列 ##

**> 上下文(Contexts)**用于控制和管理“世界”的状态

**> 内核(Kernels)**在由主机定义和操作的上下文中执行

**- 设备**

**- 内核**，OpenCL函数

**- 编程对象**，内核源代码及其可执行

**- 内存对象**

**> 命令队列**，协调内核的执行

**-** 内核执行命令

**-** 内存命令，传输或映射内存对象数据

**-** 同步命令，控制命令的顺序

**> 应用队列**计算内核执行实例

**-** 顺序排队

**-** 顺序或乱序执行

**-** 事件被用于实现执行实例的适当同步

## OpenCL内存模型 ##

**> 共享内存模型**

**-** 宽松的一致性

**> 多个不同的地址空间**

**-** 地址空间可以根据设备的内存子系统进行折叠

**> 地址空间**

**- 私有**，专用于一个工作项

**- 本地**，用于局部工作组

**- 全局**，所有工作组中的所有工作项都可访问

**- 常量**，只读全局空间

**> 实现这种映射层次**

**-** 对于可用物理内存

![SouthEast 9][]

## 内存一致性 ##

**·**“**OpenCL**使用**宽松的一致性内存模型**；例如，对一个工作项可见的内存状态，对于整个工作项的集合并不一直保证一致性。”

**· 工作项内部**，内存具有**加载/存储一致性**

**·同一屏障内的一个工作组**，**局部内存**对**所有工作项**具有**一致性**

**·全局内存**在同一屏障内的一个工作组内具有一致性，但对于**不同的工作组**并**不能保证一致性**

**　　·命令之间共享的内存一致性**，通过**同步**被强制执行

## 数据并行编程模型 ##

**　　· 定义N维计算域**

**　　　　-** 在一个N维域中的每个**独立执行单元**，被称为**工作项**

**　　　　-** N维域定义的并行执行的工作项总数 = 全局工作大小

**　　· 工作组可以被组织在一起 --- 工作组**

**　　　　-** 组中的工作项可以相互通信

**　　　　-** 可以在组内工作项之间同步执行，以协调内存访问

**　　· 并行执行多个工作组**

**　　　　-** 映射全局工作大小(global work size)到工作组，可以是显式的或隐式的

## 任务并行编程模型 ##

**　　· 数据并行执行模型**必须被**所有OpenCL计算设备实现**

**　　·** 一些计算设备，例如CPU也可以**执行任务并行计算内核**

**　　　　-** 像单个工作项那样执行

**　　　　-** 在OpenCL中编写的计算内核

**　　　　-** 本地C/C++函数

## 基本OpenCL编程结构 ##

**　　· 主机编程**

**　　　　平台层**

**　　　　-** 查询计算设备

**　　　　-** 创建上下文(Contexts)

**运行时**

**　　　　-** 创建上下文相关的内存对象

**　　　　-** 编译和创建内核编程对象

**　　　　-** 发出命令到命令队列

**　　　　-** 命令同步

**　　　　-** 清除OpenCL资源

**　　· 内核**

**　　　　语言**

**　　　　-** 带一些限制和扩展的C代码

## 例子：向量加法 ##

**　　· 计算 c = a + b**

**　　　　-** a，b和c是长度为N的向量

**　　· 基本的OpenCL概念**

**　　　　-** 简单内核代码

**　　　　-** 基本上下文管理

**　　　　-** 内存分配

**　　　　-** 内核调用

## **平台层** ##

**　　·** 平台层允许应用程序**查询平台的特定功能**

**　　· 查询平台信息**(例如，OpenCL配置文件)

**　　· 查询设备**

**　　　　****-****clGetDeviceIDs()**

**　　　　　　\-** 查找系统中存在哪些计算设备

**　　　　　　-** 设备类型包括CPU，GPU或加速器

**　　　　- clGetDeviceInfo()**

**　　　　　　-** 查询发现计算设备的能力，例如：

**　　　　　　　　\-** 计算核心数量

**　　　　　　　　-** NDRange限制

**　　　　　　　　-** 最大工作组大小

**　　　　　　　　-** 不同内存空间的大小(常量、局部变量和全局变量)

**　　　　　　　　-** 最大的内存对象大小

**　　· 创建上下文**

**　　　　-** 上下文被OpenCL运行时用于在一个或多个设备中管理对象和执行内核

**　　　　-** 上下文同一个或多个设备相关联

**　　　　　　\-** 多个上下文可以被关联到相同的设备

**　　　　-** clCreateContext()和clCreateContextFromType()返回一个处理句柄来创建上下文

## **命令队列** ##

**　　·** 命令队列存储一组操作集来执行

**　　·** 命令队列同上下文相关联

**　　·** 多命令队列可以被创建，用于处理不需要同步的独立命令

**　　·** 命令队列的执行在同步点被保证完成

## **向量加法：上下文、设备和队列** ##

![SouthEast 10][]

## **内存对象** ##

**　　· 缓冲区对象**

**　　　　-** 对象的一维集合(像C的数组)

**　　　　-** 有效的元素，包括**标量**和**矢量**类型以及**用户定义的结构**

**　　　　-** 缓冲对象可以**通过内核中的指针**被访问到

**　　· 图像对象**

**　　　　-** 二维或三维纹理，帧缓冲器或图像

**　　　　-** 必须通过内建函数编址

**　　· 采样对象**

**　　　　-** 描述如何在内核中采样一张图像

**　　　　　　\-** 寻址模式

**　　　　　　-** 过滤模式

## **创建内存对象** ##

**　　·** clCreateBuffer(), clCreateImage2D(), 和clCreateImage3D()

**　　· 内存对象**在一个**相关的上下文中创建**

**　　·** 内存可以被创建为只读，只写，或读写

**　　·** 对象在平台存储空间中的**创建位置可以被控制**

**　　　　-** 设备内存

**　　　　-** 存有从主机指针处拷贝的数据的设备内存

**　　　　-** 主机内存

**　　　　-** 与指针相关的主机内存

**　　　　　　-** 在那个点的内存在同步点保证是有效的

**　　· 图像对象**也用**通道格式创建**

**　　　　-** 通道顺序(例如，RGB，RGBA，等)

**　　　　-** 通道类型(例如，UNORM INT8，FLOAT，等)

## **操作对象数据** ##

**　　·** 对象数据可以被拷贝到内存，从主机内存，或到其它对象

**　　·** 内存命令在命令缓冲区中排队，当命令被执行时处理

**　　　　-** clEnqueueReadBuffer(), clEnqueueReadImage()

**　　　　-** clEnqueueWriteBuffer(), clEnqueueWriteImage()

**　　　　-** clEnqueueCopyBuffer(), clEnqueueCopyImage()

**　　·** 数据可以在图像和缓冲区对象之间拷贝

**　　　　-** clEnqueueCopyImageToBuffer()

**　　　　-** clEnqueueCopyBufferToImage()

**　　·** 对象数据的区域可以被访问，通过映射进主机地址空间

**　　　　-** clEnqueueMapBuffer(), clEnqueueMapImage()

**　　　　-** clEnqueueUnmapMemObject()

## **向量加法：创建内存对象** ##

![SouthEast 11][]

## **编程对象** ##

**　　· 程序对象封装**

**　　　　-** 一个相关的上下文

**　　　　-** 程序源代码或二进制文件

**　　　　-** 最新成功的程序构建，目标设备列表，构建选项

**　　　　-** 附加的内核对象的数量

**　　· 构建过程**

**　　　　1. 创建程序对象**

**　　　　　　\-** clCreateProgramWithSource()

**　　　　　　-** clCreateProgramWithBinary()

**　　　　2. 构建程序可执行**

**　　　　　　\-** 从源代码或二进制文件编译和连接，对于所有设备或相关上下文中的特定设备

**　　　　　　-** clBuildProgram()

**　　　　　　-** 构建选项

**　　　　　　　　\-** 预处理器

**　　　　　　　　-** 数学内部函数(浮点性能)

**　　　　　　　　-** 优化

## 内核对象 ##

**　　· 内核对象封装**

**　　　　-** 程序中声明的特定内核函数

**　　　　-** 用于内核执行的参数值

**　　· 创建内核对象**

**　　　　-** clCreateKernel()，在程序中为单个函数创建内核对象

**　　　　-** clCreateKernelsInProgram()，在程序中为所有内核创建一个对象

**　　· 设置参数**

**　　　　-** clSetKernelArg(<kernel>, <argument index>)

**　　　　-** 每个参数的数据必须在内核函数中进行设置

**　　　　-** 参数值被拷贝和存储在内核对象中

**　　· 内核与程序对象比较**

**　　　　-** 内核与程序执行相关

**　　　　-** 程序与程序源码相关

## **向量加法：程序与内核** ##

![SouthEast 12][]

## **向量加法：设置内核参数** ##

![SouthEast 13][]

## 内核执行 ##

**　　·** 一个执行内核的命令，必须排队到命令队列

**　　·** clEnqueueNDRangeKernel()

**　　　　-** 数据并行执行模型

**　　　　-** 描述内核执行的索引空间

**　　　　-** 需要NDRandge()维度和工作组大小的信息

**　　·** clEnqueueTask()

**　　　　-** 任务并行执行模型(多队列任务)

**　　　　-** 内核在单工作项上执行

**　　·** clEnqueueNativeKernel()

**　　　　-** 任务并行执行模型

**　　　　-** 执行一个未编译的本地C/C++函数，使用OpenCL编译器

**　　　　-** 此模式不使用内核对象，因此参数必须被传递

## 命令队列和同步 ##

**　　· 命令队列执行**

**　　　　-** 执行模型信号时，命令已完成或数据已准备就绪

**　　　　-** 命令队列可以明确地刷新到设备

**　　　　-** 命令队列顺序或乱序执行

**　　　　　　\-** 顺序，命令在顺序队列中完成，并且正确的内存是一致的

**　　　　　　-** 乱序，不保证当命令执行时或内存是一致性的，没有同步的话

**　　· 同步**

**　　　　-** 当对主机的命令完成时的信号，或其它在队列中的命令

**　　　　-** 阻塞调用

**　　　　　　\-** 命令未返回直到完成

**　　　　　　-** clEnqueueReadBuffer()可以被调用作为阻塞，将会阻塞直到完成

**　　　　-** 事件对象

**　　　　　　-** 跟踪一个命令的执行状态

**　　　　　　-** 一些命令可以阻塞指导事件对象发送一个先前命令完成的信号

**　　　　　　　　\-** clEnqueueNDRangeKernel()将一个事件对象作为参数，并等待直到一个先前命令(例如，clEnqueueWriteBuffer)完成

**　　　　　　　　-** 配置文件

**　　　　-** 队列屏障，队列命令可以阻塞命令的执行

## **向量加法：调用内核，读取输出** ##

![SouthEast 14][]

## **计算内核的OpenCL C**  ##

**　　· 由ISO C99派生**

**　　　　-** 一些限制：递归，函数指针，C99标准头文件中的函数

**　　　　-** 由C99定义的预处理指令被支持

**　　· 内建数据类型**

**　　　　-** 标量和常量数据类型，指针

**　　　　-** 数据类型转换函数：convert\_type<\_sat><\_roundingmode>

**　　　　-** 图像类型：image2d\_t, image3d\_t和sampler\_t

**　　· 内建函数 --- 必须的**

**　　　　-** 工作项函数，math.h，读写图像

**　　　　-** 关联式，几何函数，同步函数

**　　· 内建函数 --- 可选的**

**　　　　-** 双精度，原子到全局和局部内存

**　　　　-** 舍入模式的选择，写到image3d\_t上

## **OpenCL C语言函数要点** ##

**　　· 函数限定符**

**　　　　-** “\_\_kernel”限定符声明一个函数作为内核

**　　　　-** 内核可以调用其它内核函数

**　　· 地址空间限定符**

**　　　　-** \_\_global, \_\_local, \_\_constant, \_\_private

**　　　　-** 指针内核参数必须用地址空间限定符修饰来声明

**　　· 工作项函数**

**　　　　-** 查询工作项的标识符

**　　　　-** get\_work\_dim()

**　　　　-** get\_global\_id(), get\_local\_id(), get\_group\_id()

**　　· 图像函数**

**　　　　-** 图像必须通过内建函数来访问

**　　　　-** 读写执行通过从主机采样对象或在源代码中定义

**　　· 同步函数**

**　　　　-** 屏蔽，组内的所有工作项必须在任何工作项可继续之前执行屏蔽函数

**　　　　-** 内存栅栏，提供对存储器操作的顺序

## **向量加法内核** ##

![SouthEast 15][]

## **OpenCL C语言限制** ##

**　　·** 不支持 函数指针

**　　·** 在同一个内核中，允许使用指向指针的指针，但不能作为一个参数

**　　·** 不支持 位域

**　　·** 不支持 可变长数组和结构体

**　　·** 不支持 递归

**　　·** 不支持 写入的类型小于32位的指针

**　　·** 不支持 double类型，但保留

**　　·** 不支持 3D影像写操作

**　　·** 有些限制是通过扩展解决

## 可选的扩展 ##

**　　·** 扩展是可选功能，通过OpenCL接触

**　　·** OpenCL的工作组已经批准了由OpenCL的规格支持**多种扩展功能**：

**　　　　·** 双精度浮点类型

**　　　　·** 内建函数来支持doubles

**　　　　·** 原子函数

**　　　　·** 三维图像写入

**　　　　·** 字节寻址存储（写指针类型<32位）

**　　　　·** 内建函数支持半类型

## OpenGL互操作性 ##

**　　· 一个IP架构下两个标准**

**　　　　-** 使得非常密切的协作设计

**　　· 高效的跨API通信**

**　　-** 同时仍然允许这两个API来处理其设计工作负载的类型

**　　· OpenCL可高效地同OpenGL共享资源**

**　　-** 纹理，缓冲区对象，渲染缓冲

**　　-** 数据是共享的，但不可复制

**　　-** OpenCL的对象是从OpenGL的对象创建

**　　　　\-** clCreateFromGLBuffer(), clCreateFromGLTexture2D(), clCreateFromGLRenderbuffer()

**　　· 应用程序可选择计算设备来运行OpenCL和OpenGL**

**　　-** 高效地排队OpenCL和OpenGL命令到硬件

**　　-** 灵活的调度和同步

**　　· 举例**

**　　-** 顶点(Vertex)和图像数据由OpenCL的生成，然后用OpenGL渲染

**　　-** 图像由OpenGL渲染后，由OpenCL内核后处理

## OpenCL总结 ##

**　　· 跨供应商标准的便携式异构编程**

**　　-** 开放、免费的版权标准，来自主要供应商的临界质量支持

**　　· 创造显著的商业机会**

**　　-** 移除作为市场障碍的碎片化，以增加并行计算的增长

**　　· 在Khronos的API生态系统中的核心角色**

**　　-** 在Khronos，多个相关的API正被在一个IP框架下合作开发

**　　· 快速部署**

**　　-** 公开规则在六个月内产生，在2009年实现

**　　-**  将运行于目前最新一代的GPU硬件

**　　· 更多说明书和幻灯片在**[www.khronos.org/opencl/][www.khronos.org_opencl]

**　　-** 如果这里讲的与你公司相关，请加入Khronos并且动手！

[SouthEast]: /images/20220801/75965923df9b4592999c36b19b137624.png
[SouthEast 1]: /images/20220801/087ec2f9c658464ba875ad8f99f26dac.png
[SouthEast 2]: /images/20220801/3d83ad109ea245448805b231b068f1ea.png
[SouthEast 3]: /images/20220801/f7127c072b78435fa84510c4921fa0f0.png
[SouthEast 4]: /images/20220801/a61e0f059f8e4e4f990115607b3cf0cc.png
[SouthEast 5]: /images/20220801/dc18a95b8e794a90823b18135e2cfe55.png
[SouthEast 6]: /images/20220801/bfada05511124f579213cf0b9f6c6a7e.png
[SouthEast 7]: /images/20220801/623ef7e10e7446868db0f8bf42435cab.png
[SouthEast 8]: /images/20220801/94fe5309f07d4ec486895dc2eb401f8a.png
[SouthEast 9]: /images/20220801/55bf76bacfca4df6adbb5cf60eff4b2e.png
[SouthEast 10]: /images/20220801/348b08dca1594a3dbb39ace21b1917f2.png
[SouthEast 11]: /images/20220801/9e673e0fb83943c6ba1dfa758ad08459.png
[SouthEast 12]: /images/20220801/aed8956d6761405d998e7a9351d71633.png
[SouthEast 13]: /images/20220801/cb5a1b70dc334afba1f1ed29458bf378.png
[SouthEast 14]: /images/20220801/26dd357a44104b229ec9917c6370815a.png
[SouthEast 15]: /images/20220801/4f004ba4e47944cdac9d6bba63d02178.png
[www.khronos.org_opencl]: http://blog.csdn.net/cloud_desktop/article/details/www.khronos.org/opencl/