Mali™-T600 系列 GPU OpenCL
开发者指南
版权所有 © 2012-2013 ARM.保留所有权利。
DUI0538E (ID021813)
Mali-T600 系列 GPU OpenCL
版本信息
本指南已进行了下列更改。
更改记录
日期
发布
保密性
更改
2012 年 7 月 12 日
A
机密
r1p1 第一版
2012 年 11 月 7 日
D
r1p2 第一版
2013 年 2 月 18 日
E
Mali T600 系列 OpenCL SDK 第一版
所有权通告
除非此所有权通告的下文中另有说明，否则标有 ™ 或 ® 的词语和徽标是 ARM® 在欧盟和其他国家/地区的注册商标或商标。此处提及的其他品牌和名称可能是其相关所有者的商标。
除非事先得到版权所有人的书面许可，否则不得以任何形式改编或复制本文档包含或所描述产品的全部或部分信息。
本文档中描述的产品将进行持续开发和改进。ARM 如实提供本文档包含的所有产品特性及其使用方法。但是，所有暗含或明示的保证（包括但不限于暗含的有关适销或适用于某种特定用途的保证）均不包括在内。
本文档的目的仅在于帮助读者使用产品。对因使用本文档中的任何信息、这些信息中的任何错误或遗漏或者任何不正确的使用产品而导致的任何损失或损害，ARM 概不负责。
术语 ARM 在其使用之处表示“ARM 或任何适当的子公司”。
保密状态
本文档是非机密文档。根据 ARM 与本文档的接收方之间协议的条款，使用、复制及披露本文档的权利受到许可限制。
产品状态
本文档中的信息是最终版，即用于开发完成的产品。
网址
http://www.arm.com
非机密
目录
Mali-T600 系列 GPU OpenCL 开发者指南
前言
关于本指南 ........................................................................................................... vi 反馈 .................................................................................................................. viii
第 1 章 简介
第 2 章 并行处理概念
第 3 章 OpenCL 概念
第 4 章 OpenCL 程序中的阶段
第 5 章 转换现有代码到 OpenCL
第 6 章 为 Mali GPU 重调现有 OpenCL 代码
第 7 章 为 Mali GPU 优化 OpenCL
第 8 章 Mali OpenCL SDK
附录 A OpenCL 数据类型
附录 B OpenCL 内置函数
附录 C OpenCL 扩展
本前言介绍《Mali-T600 系列 GPU OpenCL 开发者指南》。其中包含下列小节：
• 关于本指南（第 vi 页）。
• 反馈（第 viii 页）。
关于本指南
本指南适用于 Mali-T600 Linux OpenCL 软件开发包 (SDK)。
产品修订状态
rnpn 标识符表示本指南中所述产品的修订状态，其中：
rn 表示产品的大幅修订。
pn 表示产品的细微修订或修改状态。
目标读者
本指南专门为具有 C 或 C 类语言的经验、希望为 Mali-T600 系列 GPU 开发 OpenCL 应用程序的软件开发人员而撰写。
使用本指南
本指南分为下列章节：
本章提供 OpenCL 和 Mali-T600 系列 OpenCL SDK 的简介。
本章介绍并行处理的概念以及在 OpenCL 中的运作原理。
本章介绍 OpenCL 概念。
本章介绍 OpenCL 程序中的各个阶段。
本章介绍如何将现有代码转换到 OpenCL。
本章介绍如何为 Mali-T600 系列 GPU 重新调整现有 OpenCL 代码。
本章介绍如何为 Mali-T600 系列 GPU 优化 OpenCL。
本章介绍 Mali OpenCL SDK。
本章介绍 OpenCL 内存模型，以及可用的数据类型。
本章提供 Mali-T600 系列 Linux OpenCL 驱动程序中实施的 OpenCL 内置函数的列表。
本章提供 Mali-T600 系列 Linux OpenCL 驱动程序支持的扩展的列表。
术语表
ARM 术语表列举了 ARM 文档中使用的术语，还提供这些术语的定义。ARM 术语表不包含行业标准术语，除非 ARM 的含义与公认的含义不同。
请参见 ARM 术语表，网址为 http://infocenter.arm.com/help/topic/com.arm.doc.aeg0014-/index.html。
排版约定
本指南使用下列排版约定：
斜体
介绍特殊术语，以及表示交叉引用和引文。
粗体
突出显示界面元素，如菜单名称等。表示信号名称。适当时也用于描述性列表中的术语。
等宽字体
表示您可以在键盘上输入的文本，如命令、文件和程序名称，以及源代码。
表示命令或选项允许使用的缩写。您可以输入标有下划线的文本，而不输入完整的命令或选项名称。
等宽斜体
表示等宽字体文本中的参数，该参数可以替换为具体的值。
等宽粗体
表示语言关键字，在示例代码之外使用。
< 和 >
出现在代码或代码片段中时，用于括起汇编程序语法可替换的术语。例如：
MRC p15, 0 <Rd>, <CRn>, <CRm>, <Opcode_2>
其他读物
本节提供 ARM 和第三方的出版物列表。
若要访问 ARM 文档，请参见信息中心网站，网址为 http://infocenter.arm.com。
其他出版物
本节提供第三方出版的相关文档列表：
反馈
ARM 欢迎您就本产品和其文档提供反馈。
有关本产品的反馈
如果您对本产品有任何意见或建议，请联系您的供应商，并提供：
有关内容的反馈
如果您有内容方面的意见，请发送电子邮件到 errata@arm.com。请提供：
ARM 也欢迎您就增补和改进提供一般建议。
备注
ARM 仅在 Adobe Acrobat 和 Acrobat Reader 测试了本 PDF 文件，无法担保使用其他 PDF 阅读器时的文档显示质量。
第 1 章
简介
本章介绍 GPU 计算、OpenCL、Mali-T600 系列 Linux OpenCL 驱动程序，以及 Mali OpenCL SDK。其中包含下列小节：
GPU 计算，或者图形处理单元上通用计算 (GPGPU) 是使用 GPU 并行计算能力执行 3D 图形渲染以外任务的做法。
应用处理器设计为尽可能快速执行单个线程。此类处理通常包含标量运算和控制代码。
GPU 设计为能够同时执行许多个线程。它们能够并行计算包含相对较少控制代码的密集数据处理任务。GPU 通常包含数量远超应用处理器的处理元素，因而能够以比应用处理器快很多的速度进行计算。
OpenCL 是第一种开放标准语言，能够让开发者在 GPU、应用处理器和其他类型的处理器上运行通用计算任务。
开放计算语言 (OpenCL) 是一种开放标准，用于编写在异构多处理器系统上运行的程序。OpenCL 为能够在不同处理器上运行的程序提供单一编程环境。
OpenCL 包含独立于平台的 C99 型语言，可用于编写在 OpenCL 设备上执行的函数（称为内核），以及定义和控制平台的 API。
OpenCL 可以让您通过将密集数据处理例程移到 GPU 而非应用处理器，更加快速地执行某些程序。
OpenCL 是 Khronos 集团 (http://www.khronos.org) 制定的开放标准。
Mali-T600 系列 Linux OpenCL 驱动程序针对的是 Mali-T600 系列 GPU 的 OpenCL 实施。在本文中，它称为 Mali OpenCL 驱动程序。
Mali OpenCL 驱动程序：
Mali OpenCL 驱动程序适用于 Mali-T600 系列 GPU，不支持 Mali-300、Mali-400 或 Mali-450 GPU。
Mali-T600 系列 OpenCL SDK 包含示代码示例和教程，帮助您了解 OpenCL 开发。
请参见第 8 章 Mali OpenCL SDK。
第 2 章
并行处理概念
并行处理是同时在多个处理器上处理计算。OpenCL 可以使应用程序利用 GPU 等硬件资源通过并行处理加快计算速度。
本章介绍并行处理的主要概念。其中包含下列小节：
有下列几种并行类型：
数据并行
在数据并行应用程序中，数据分割为能够并行处理的数据元素。同时由不同的处理器读取和处理多个数据元素。
数据必须处于能够并行读取和写入的数据结构中。
渲染三维图形就是数据并行应用程序的一个示例。生成的像素相互独立，因此可以并行执行生成这些像素所需的计算。这种类型的并行粒度很细，同时存在数百或数千个活动线程。
OpenCL 主要用于数据并行处理。
任务并行
任务并行是指应用程序分割为多个任务，这些任务再并行执行。任务并行也称为功能并行。
应用程序可以使用任务并行的一个示例是播放在线视频。要显示网页，您的设备必须执行下列几项任务：
图 2-1 显示播放在线视频时同时运作的应用程序和操作系统的各个部分。
绘制视频帧
播放声音
解码声音
解码视频
解析数据
从外部服务器请求数据
从外部服务器读取数据
操作系统
网络堆栈
图 2-1 任务并行处理
流水线
流水线在一系列阶段中处理数据。在流水线中，各个阶段可以同时运作，但它们不处理相同的数据。流水线中的阶段数量通常相对较小。
作为流水线的一个示例，视频录制应用程序必须执行下列阶段：
这些阶段必须按秩序执行，但它们可以同时针对不同视频帧的数据运行。
图 2-2 显示流水线播放互联网视频时同时运作的应用程序的各个部分。
从图像传感器采集数据
针对镜头特效修正图像
修改：对比度、色彩平衡、曝光度
压缩图像
将数据添加到视频文件
将视频文件写入存储
图 2-2 流水线处理
并发应用程序同时进行多个运算。它们可以并行运算，或者通过时间共享系统串行运算。
在并发应用程序中，通常会有多个任务尝试共享同一数据。访问此类数据必须谨慎管理，否则会导致复杂的问题，例如：
竞争条件
当两个或更多个线程同时尝试修改一个变量的值时，就会出现竞争条件。变量的值应当始终相同，但出现竞争条件时，变量可以根据写入操作的顺序而获得不同的值。
死锁 当两个线程互相阻止、无法继续各自的运算时，会出现死锁。线程各自获得对方所需的锁定时就会出现此情况。
活锁 活锁与死锁相似，但线程保持运行。然而，线程由于被锁定而无法完成其任务。
并发数据结构是能够被多个任务访问而不造成并发问题的数据结构。
数据并行应用程序使用并发数据结构。它们是通常在 OpenCL 中使用的数据结构。
在开发并行应用程序时，您必须考虑并行处理的一些限制。
例如，如果应用程序完美并行，则在 10 个处理器上执行该应用程序，速度就能快 10 倍。
应用程序很少能完美并行，因为其存在串行的部分。此串行部分限制了应用程序可以利用的并行数量。
阿姆达尔定律描述了您可以从并行处理获得的加速。图 2-3 显示阿姆达尔定律的公式，其等式中的条件为：
S 应用程序中串行的部分。
P 应用程序中并行的部分。
N 处理器数量。
加速 =
图 2-3 阿姆达尔定律公式
图 2-4 显示不同数量的处理器为具有不同串行部分的应用程序提供的加速。
加速
完美缩放
5% 串行
10% 串行
20% 串行
处理器
图 2-4 具有不同串行部分的应用程序的加速
处理器数量相对较少时，加速最大。但是，随着处理器数量增加，增益会减少。
您无法在应用程序中避免阿姆达尔定律，但可以减轻其影响。请参见减小串行计算的影响（第 7-8 页）。
如果要在有大量处理器时获得较高的性能，应用程序中的串行部分必须非常小。此类应用程序被称为高度并行。请参见高度并行的应用程序（第 2-6 页）。
如果应用程序可以轻松在大量处理器之间并行，其被称为高度并行。
渲染三维图形就是高度并行应用程序的一个示例。像素是完全独立的，因此可以并行地计算和绘制它们。
OpenCL 完美适合开发和执行高度并行的应用程序。
图 2-5 高度并行的应用程序
图 2-5 显示一个图像被分割为多个小部分。这些部分全部可以同时处理。
您可以在应用程序中混合使用这些并行类型。例如，一个音频合成器可能使用所有这三种并行：
第 3 章
OpenCL 概念
本章介绍 OpenCL 概念。其中包含下列小节：
OpenCL 是一种开放标准，让您能够利用多种处理器的并行处理能力，如应用处理器、GPU 和其他计算设备。
OpenCL 为并行编程指定了一个 API，旨在获得移植能力：
在 OpenCL 设备上执行的函数称为内核。它们使用称为 OpenCL C 的语言编写，该语言基于 C99。
Mali-T600 系列 GPU 支持 OpenCL 1.1 Full Profile。
OpenCL 应用程序包含两个部分：
应用程序（或主机）端代码
OpenCL 内核
您必须正确编写这两个部分，才能获得最佳的性能。
OpenCL 执行模型包含：
主机程序
主机程序通过为下列命令设置命令队列，管理内核的执行：
上下文
主机程序定义内核的上下文。上下文包含：
OpenCL 内核的运算
内核是一个代码块，它在计算设备上与其他内核并行执行。内核按照下列顺序运算：
本节介绍 OpenCL 数据处理。其中包含下列小节：
由 OpenCL 处理的数据位于工作项的索引空间中。工作项排列在 N 维范围 (NDRange) 中，其中：
针对索引空间中的各个工作项执行一个内核实例。
图 3-1 显示具有一维、二维和三维的 NDRange。
一维 NDRange
工作项
二维 NDRange
三维 NDRange
图 3-1 NDRange 和工作项
您可以将工作项分到工作组中，以进行处理。图 3-2（第 3-6 页）显示了一个三维 NDRange，其拆分为 16 个工作组，各自包含 16 个工作项。
工作组
图 3-2 工作项和工作组
工作组具有多个属性和限制：
工作组属性
工作组限制
工作组通常不直接共享数据。它们使用全局内存共享数据。
不同的工作组之间不支持如下所列：
可以使用全局原子，但它们的速度比本地原子慢。
工作组中的工作项
一个工作组中的工作项可以执行如下所列：
• 执行屏障运算，以同步执行点。例如：
barrier(CLK_LOCAL_MEM_FENCE); // Wait for all kernels in // this work-group to catch up
OpenCL 中有多个标识符：
全局 ID 每个工作项具有一个唯一全球 ID，标识其在索引空间中的身份。
本地 ID 在各个工作组中，每个工作项具有一个唯一本地 ID，标识其在所属工作组中的身份。
工作组 ID
每个工作组具有一个唯一工作组 ID。
OpenCL 内存模型包含多个组件。图 3-3 显示 OpenCL 内存模型。
专用内存
本地内存
常量内存
全局内存
图 3-3 OpenCL 内存模型
Mali GPU 具有与桌面工作站不同的内存模型：
桌面 传统桌面工作站在物理上通常分隔全局、本地和
专用内存。
显卡通常拥有自己的本地内存。
数据必须复制到本地内存，并重新拷回。
Mali GPU Mali GPU 具有统一内存系统。
本地和专用内存在物理上是全局内存。
将数据从全局内存移到本地或专用内存通常不会提高性能。
不需要传统的数据拷贝操作。
每一计算设备（即着色器核心）拥有自己的数据缓存。
OpenCL 主要使用数据并行处理。OpenCL 使用下列术语：
第 4 章
OpenCL 程序中的阶段
本章介绍 OpenCL 程序中的阶段。其中包含下列小节：
要为 Mali GPU 开发 OpenCL 程序，您需要：
您可以使用搭载 OpenCL 实施的其他硬件平台进行开发，但不能用它们来估计 Mali-T600 系列 GPU 上的性能。
目前有针对下列操作系统的 OpenCL 实施：
下列为开发 OpenCL 应用程序的各个阶段：
在决定是否要使用 OpenCL 时，第一步是要查看您的程序的功能，并辨别程序的哪些部分可以并行运行。这常常是开发 OpenCL 程序时最艰难的部分。请参见为并行分析代码（第 5-3 页）。
仅在有可能带来益处的地方，将程序的相关部分转换到 OpenCL。对应用程序进行性能分析，以找出最为活跃的部分，再考虑对这些部分进行转换。
OpenCL 程序包含一组内核函数。您必须编写执行相关计算的内核。
利用 OpenCL 平台层函数进行查询，确定系统上有哪些可用的 OpenCL 设备。请参见查找可用计算设备（第 4-5 页）。
创建和设置 OpenCL 上下文和一个或多个命令队列，以调度内核的执行。请参见初始化和创建 OpenCL 上下文（第 4-6 页）。
编写程序中的代码，其应包含用于编译和生成源代码的命令，并可从编译的代码中提取内核对象。您必须遵循的命令序列为：
使用 OpenCL API 分配内存缓冲区。您可以使用 map() 和 unmap() 运算，使应用处理器和 Mali GPU 能够访问相关数据。
将控制内核执行顺序和同步的命令、读取和写入数据的命令，以及操纵内存对象的命令排入命令队列中。
要执行内核函数，您必须进行如下操作：
要设置 OpenCL，您必须选择计算设备。调用 clGetDeviceIDs() 查询 OpenCL 驱动程序，以获取机器上支持 OpenCL 的设备列表。您可以将搜索范围限定为某一设备类型，或者任何设备类型的组合。您也必须指定想要返回的设备 ID 数上限。
了解机器上可用的 OpenCL 设备，并且至少拥有一个有效设备 ID 后，您就可以创建 OpenCL 上下文。上下文将设备分组到一起，使内存对象可在不同计算设备之间共享。
将设备信息传递到 clCreateContext() 函数。例如：
// Create an OpenCL context
context = clCreateContext( NULL, 1, &device_id, notify_function, NULL, &err );
if (err != CL_SUCCESS)
{
Cleanup();
return 1;
}
您可以选择在创建 OpenCL 上下文时指定错误通知回调函数。将此参数留为 NULL 值会导致不注册任何错误通知。
如果想要针对特定的 OpenCL 上下文接收运行时错误，提供回调函数很有用处。例如：
// Optionally user_data can contain contextual information
// Implementation specific data of size cb, can be returned in private_info
void context_notify( const char *notify_message, const void *private_info, size_t cb, void *user_data )
printf("Notification:\n\t%s\n", notify_message);
在创建 OpenCL 上下文后，可使用 clCreateCommandQueue() 创建命令队列。例如：
// Create a command-queue on the first device available
// on the created context
commandQueue = clCreateCommandQueue( context, &device);
if (commandQueue == NULL) {
如果有多个 OpenCL 设备，如应用处理器和 GPU，您必须：
加载 OpenCL C 内核源代码，从中创建程序对象。程序对象与内核源代码一起加载，而后代码被编译，以便在与上下文关联的设备上执行。所有内核函数必须在程序源代码中使用 __kernel 限定符进行标识。OpenCL 程序也可以包含您可从内核函数中调用的函数。
程序对象包含：
要创建程序对象，请使用 clCreateProgramWithSource() 函数。例如：
// Create OpenCL program
program = clCreateProgramWithSource( context, device, “<kernel source>”);
if (program == NULL)
有不同的选项可用于生成 OpenCL 程序：
使用预生成程序对象的应用程序无法移植。
从二进制文件创建程序对象的过程与从源代码创建程序对象相似，除了您必须为希望在其上运行内核的各个设备提供二进制文件之外。使用 clCreateProgramWithBinary() 函数可执行此操作。
使用 clGetProgramInfo() 函数，以便在生成之后获取二进制文件。
在创建了程序对象后，您必须从程序对象的内容中生成程序可执行文件。使用 clBuildProgram() 函数来生成可执行文件。
编译程序对象中的所有内核。在下列代码中，只有一个内核：
err = clBuildProgram( program, 1, &device_id, "", NULL, NULL ); if (err == NULL)
本节介绍创建内核和内存对象。其中包含下列小节：
调用 clCreateKernel() 函数创建单个内核对象，或者调用 clCreateKernelsInProgram() 函数为 OpenCL 程序中的所有内核创建内核对象。例如：
// Create OpenCL kernel
kernel = clCreateKernel(program, “<kernel_name>", NULL);
if (kernel == NULL)
在创建并注册了内核后，可将程序数据发送到内核：
内存对象有两种：
缓冲对象
简单的内存块。
图像对象
不透明结构，尤其用于呈现 2D 或 3D 图像。
要创建缓冲对象，请使用 clCreateBuffer() 函数。要创建图像对象，请使用 clCreateImage2D() 或 clCreateImage3D() 函数。
本节介绍执行内核中的阶段。其中包含下列小节：
如果数据为 x 像素宽乘以 y 像素高的图像，则其为二维数据集。如果您处理的空间数据涉及节点的 x、y 和 z 位置，则其为三维数据集。
在 OpenCL 中，原始数据集中的维度数量不必相同。例如，您可以将三维数据集作为一维数据集在 OpenCL 中处理。
全局工作大小是所有维度组合起来所需的工作项总数。
您可以在单个工作项里处理多个数据项，以此更改全局工作大小。新的全局工作大小就是原始全局工作大小除以每个工作项中处理的数据项数目。
使用表 4-1 中的公式可以计算出不同维度的数据的全局工作大小，其中 n 为单个工作项中处理的数据元素数目。
表 4-1
维度
公式
一
x / n
二
( x * y ) / n
三
(x * y * z ) / n
如果您希望确保高性能，则最佳全局工作大小必须很大。其数目通常为几千，但理想数值取决于设备中着色器核心的数量。
要计算最佳全局工作大小，请使用下列等式：
全局工作大小 = <最大工作组大小> * <着色器核心数> * <常数>
其中，Mali-T604 GPU 的常数通常为 4 或 8。
您可以指定将设备上执行的核心排入队列时 OpenCL 所使用的工作组大小。要进行此操作，您必须了解工作项在其上执行的 OpenCL 设备所允许的最大工作组大小。要找出最大工作组大小，请使用 clGetKernelWorkGroupInfo() 函数，并请求 CL_KERNEL_WORK_GROUP_SIZE 属性。
如果应用程序不需要在工作项之间共享数据，请在将内核排入队列时将 local_work_size 参数设置为 NULL。这可让 OpenCL 驱动程序为您的核心确定最高效的工作组大小。
要获得每个维度的最大工作组大小，请调用 clGetDeviceInfo() 并使用 CL_DEVICE_MAX_WORK_ITEM_SIZES。要获得工作组大小总计，请调用 clGetKernelWorkGroupInfo() 并使用 CL_KERNEL_WORK_GROUP_SIZE。如果某一内核的最大工作组大小比 128 小，请尝试简化该内核。
每个维度的工作组大小必须均匀分入该维度的总数据大小。也就是说，工作组的 x 大小必须均匀分入总数据的 x 大小。如果此要求意味着需要用额外工作项来充填工作组，请确保这些额外工作项可以立即返回并且不执行任何工作。
确定了显示数据所需的维度、各个维度所需的工作项，以及相应的工作组大小后，请使用 clEnqueueNDRangeKernel() 将内核执行排入队列。例如：
size_t globalWorkSize[1] = { ARRAY_SIZE };
size_t localWorkSize[1] = { 4 } ;
// Queue the kernel up for execution across the array
errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorkSize,
localWorkSize, 0, NULL, NULL);
if (errNum != CL_SUCCESS)
printf( "Error queuing kernel for execution.\n" ); Cleanup(); return 1;
将内核执行排入队列并不意味着它会立即执行。内核执行被放入命令队列，由设备稍后处理。在调用 clEnqueueNDRangeKernel() 后，内核可能还未在设备上执行。
可以让内核等待执行，直到先前的事件结束为止。您可以指定特定的内核等待其他特定内核完成后再执行。
在内核完成执行后，您必须使结果能被主机访问。
要从内核访问结果，请使用 clEnqueueMapBuffer() 将缓冲区映射到主机内存。例如：
local_buffer = clEnqueueMapBuffer( queue, buffer, CL_FALSE, CL_MAP_READ, 0, ( sizeof(unsigned char) * buffer_size), num_deps, deps[1], NULL, &err );
ASSERT( CL_SUCCESS == err );
这一调用并不保证使缓冲区可用，直到您调用 clFinish() 为止。如果将第三个参数 CL_FALSE 更改为 CL_TRUE，则该调用变为阻止性调用，会立即予以执行。
当程序不再需要与 OpenCL 运行时和上下文相关的各种对象时，您必须释放这些资源。使用下列函数来释放您的 OpenCL 对象。这些函数会减小相关对象的引用计数：
在程序不再需要 OpenCL 对象时，请确保它们的引用计数都达到零。您可以通过查询对象来获得引用计数。例如，通过调用 clGetMemObjectInfo()。
第 5 章
转换现有代码到 OpenCL
本节介绍将现有代码转换到 OpenCL。其中包含下列小节：
对应用程序进行性能分析，找出计算最密集的部分。这些部分或许值得移植到 OpenCL。
应用程序中需要较高性能的部分通常是相对较小的代码部分。这一部分代码或许能够充分利用 OpenCL。将应用程序的任何更多部分移植到 OpenCL 可能无法带来益处。
您可以使用 DS-5™ 等性能分析程序对应用程序执行性能分析。DS-5 可以从 Mali 开发者网站下载，网址为 http://www.malideveloper.arm.com。
本节介绍如何为并行分析计算密集型代码。其中包含下列小节：
在确定了应用程序中计算最为密集的部分后，请分析其代码是否能够并行运行。
并行代码可以有以下情况：
轻松 并行该代码需要少许或不需要修改。
直观
并行该代码需要细微修改。请参见使用全局 ID 而非循环计数器（第 5-5 页）。
困难 并行该代码需要复杂的修改。请参见使用公式计算循环中的值，不要使用计数器（第 5-6 页）。
困难并且包含依赖项
并行该代码需要复杂的修改，而且要使用特殊的技巧来避免依赖项。请参见下列小节：
几乎不可能
如果并行该代码似乎不太可能，请研究该代码使用的算法和数据结构的并行替代方法。它们或许可使并行成为可能。请参见将并行处理用于不可并行的代码（第 5-10 页）。
不可能
这仅仅表示该实施无法并行。不要认为该代码是问题的绝对解决方案。代码仅仅是一个解决方案的一种可能实施。可能会存在多种不同的解决方案，其中一些或许可以并行。请参见将并行处理用于不可并行的代码（第 5-10 页）。
查找执行大量运算并满足以下条件的任务：
这些类型的运算属于数据并行型，因而完美适合 OpenCL。
如果任务基本没有依赖项，或许能够并行运行它们。
任务之间的依赖关系妨碍并行，因为它会强制任务按顺序执行。如果代码具有依赖项，请考虑：
循环是并行的良好目标，因为它们重复计算许多次，常常各自独立。
处理少量元素的循环
如果循环仅处理相对少量的元素，可能不适合数据并行处理。或许最好利用任务并行在一个或多个应用处理器上并行处理此类循环。
完美循环
寻找以下循环：
这些类型的循环属于数据并行型，因而完美适合 OpenCL。
简单循环并行
如果循环中包含根据上一迭代的值而递增的变量，这表示迭代之间存在妨碍并行的依赖关系。
查看您是否能够找出一个公式，可以根据主循环计数器计算该变量的值。
在 OpenCL 中，内核并行处理工作项。不存在可以引用的循环计数器，因为工作项不是在循环中处理的。
每个工作项都有一个唯一全局 ID 标识其身份。您可以使用此值，替代循环计数器。请参见使用全局 ID 而非循环计数器（第 5-5 页）。
循环需要来自上一迭代的数据
如果您的循环涉及基于上一迭代处理的数据的依赖项，则问题更为复杂。
是否能够重新构建循环以去除依赖关系？如果不行，则可能无法并行该循环。
有多种技巧可以分离依赖项。看看您是否可以使用这些技巧来并行循环。如需其中部分技巧的说明，请参见并行处理技巧（第 5-5 页）。
不可并行的循环
如果循环包含您无法去除的依赖项，请研究执行该计算的替代方法。它们或许可以并行。请参见将并行处理用于不可并行的代码（第 5-10 页）。
本节介绍您可以在 OpenCL 中使用的并行处理技巧。其中包含下列小节：
在 OpenCL 中，您使用内核来执行与循环迭代相当的任务。这意味着没有循环计数器可在计算中使用。
工作项的全局 ID 提供了与循环计数器相当的功能。使用全局 ID 来执行任何基于循环计数器的计算。
您可以在 OpenCL 内核中包含循环，但它们只能迭代该工作项的数据，而不是整个 NDRange。
下例演示了 C 语言中的一个简单循环，它将循环计数器的值分配给各个阵列元素。
C 语言循环示例：
下列循环使用数字填充阵列。
void SetElements(void) { int loop_count; int my_array[4096];
for (loop_count = 0; loop_count < 4096; loop_count++) { my_array[loop_count] = loop_count; }
printf("Total %d\n", loop_count); }
此循环可以并行，因为其循环元素都是独立的。OpenCL 内核中没有主循环计数器 loop_count，因此它替换为全局 ID。
OpenCL 内核中对应的代码：
__kernel void example(__global int * restrict my_array)
{ int id; id = get_global_id(0); my_array[id] = id;
OpenCL 内核执行与循环迭代相当的任务，所以无法根据循环迭代递增或递减变量。
为了替代递增和递减的变量，可设计一个公式，根据全局 ID 的值来计算变量的值。工作项的全局 ID 提供了与循环计数器相当的功能。
如果应用程序需要持续更新数据元素，并且它们之间存在依赖项，请尝试将计算分割为离散的单元，按照显示的每一图像帧执行一次迭代。
例如，图 5-1 中显示的图像是一个持续进行旗帜的物理学模拟的应用程序。
图 5-1 旗帜模拟
旗帜由一个节点网格组成，各个节点与相邻节点连接。它们显示在图 5-2（第 5-7 页）中。
图 5-2 旗帜模拟网格
该模拟作为一系列迭代运行。每次迭代中，所有节点都被更新，图像也重新绘制。
每次迭代中执行下列运算：
在这一情形中，将计算拆分为迭代也会拆分依赖项。一帧需要的数据在前一帧中计算。
一些类型的模拟需要许多迭代实现相对较小的运动。如果是这样，请尝试在绘制帧之前计算多个迭代。
如果应用程序需要持续更新数据元素，并且它们之间存在依赖项，请尝试将计算分割为离散的单元，并在多个阶段中执行这些计算。
此技巧通过进一步分割计算，扩展了按帧计算值（第 5-6 页）中所述的技巧。
将数据元素分割为奇数和偶数字段。这可以分离依赖项，让整个计算能够分阶段执行。处理过程在计算奇数和偶数字段之间轮换。
例如，这一技巧可以在神经网络模拟中使用。
各个神经元排列在一个三维网格中。计算一个神经元的状态涉及读取来自周围神经元的输入。这意味着每个神经元具有对周围神经元状态的依赖关系。
要执行该模拟，整个三维网格分割为图层，并以下列方式执行：
如果计算中有串行部分，请查看其是否可以去除并单独执行。
例如，调频 (FM) 这种音频合成技术通过读取称为载波的音频波形来运作。该波形的读取速率取决于称为调制器的另一波形。
载波值由指针读取，生成输出波形。通过获取上一个值，并根据调制器波形的值所确定的量移动它，以此计算出指针的位置。
指针的位置依赖于上一个值，那个值又依赖于它前面的值。这一依赖序列让该算法很难或无法并行。
另一种方法可以这样思考，指针在载波波形间以固定的速度移动，同时调制器增加或减去一个偏移量。这可以并行计算，但偏移量不正确，因为它们没有将对上一偏移的依赖关系考虑在内。
计算正确的偏移量是个串行过程。如果您预先计算这些值，那么剩余的计算就可并行。并行部分从生成的偏移量表读取数据，并用它来从载波波形中读取正确的值。
这一示例有一个潜在问题。调制波形每一次更改时，必须重新计算偏移量表。这就是阿姆达尔定律的一个示例。可能并行计算的量受到串行计算速度的限制。
软件流水线是并行处理技巧，它们可以将计算分割为一系列顺序阶段，实现同步处理多个数据元素。
流水线在硬件和软件中都很常见。例如，应用处理器和 GPU 使用硬件流水线。图形标准 OpenGL ES 则基于虚拟流水线。
在流水线中，一个完整的流程被分割为一系列阶段。数据元素在一个阶段中处理，然后其结果传递到下一阶段。
由于流水线的顺序性质，一个阶段在一个时间上仅被一个特定数据元素使用。这意味着其他阶段可以处理其他数据元素。
您可以在应用程序中使用软件流水线来处理不同的数据元素。
例如，游戏需要同时进行许多不同的运算。游戏可能采用类似如下的流水线：
任务或功能并行涉及将应用程序根据功能划分为不同的任务。
例如，在线游戏可以利用任务并行。要运行在线游戏，您的设备会执行下列几项功能：
这些任务需要同步，但大体上是互相独立的运算。这意味着您可以在不同的处理器上并行执行这些任务。
另一个任务并行示例是数字电视 (DTV)。在任何时间，电视机可能会执行下列几个操作：
如果无法并行您的代码，您依然可以使用并行处理。
大多数代码编写为在按顺序运行的应用处理器上运行。该代码使用串行算法和非并发数据结构。并行此类代码或许难度很高或无法实现。
代码无法并行这一情况仅仅意味着无法并行这一特定实施。并不表示该问题无法以并行方式解决。
调查下列方法：
使用数据结构和算法的并行版本
许多常见数据结构和使用它们的算法是非并发的。这妨碍您并行其代码。
许多常见数据结构和算法都有并行的版本。您或许能够使用它们替代原始版本来并行代码。
请参见使用并发数据结构（第 5-11 页）。
以不同方式解决问题
退后一步，思考代码解决的问题是什么。
审视这一问题，研究解决问题的替代方式。或许存在替代的方案，能使用可以并行的算法和数据结构。
为此，请从代码和数据结构的用途方面进行思考。
通常，代码的用途是处理或转换数据。它获取特定的输入，产生特定的输出。
如果以上的答案是肯定的，那么或许能够使用 OpenCL 解决您的问题。
本节介绍分割数据以便使用 OpenCL 处理。其中包含下列小节：
分割数据，这样它就能利用 OpenCL 并行计算。数据分割到下列级别的层次结构：
请参见第 3 章 OpenCL 概念。
OpenCL 执行数百或数千个内核实例，因此处理和数据结构必须能够并行到该程度。
这意味着，您必须使用允许同时独立读取和写入多个数据元素的数据结构。它们称为并发数据结构。
许多常见数据结构是非并发结构。这会妨碍并行其代码。例如，下列数据结构是非并发的：
许多常用数据结构都有并行的版本。
下列示例中是不同维度的数据，您可以使用 OpenCL 进行处理：
这些示例将问题映射到维度数量相同的 NDRange。OpenCL 并不要求您这么做。您可以将一维问题映射到二维或三维 NDRange。
一维数据
一维数据的一个示例是音频。音频通过一系列样本呈现。更改音频音量是并行任务，因为其运算是按照样本独立执行的。
在这一情形中，NDRange 是音频中的样本总数。每个工作项可以是一个样本，工作组则是样本的一个集合。
音频也可通过矢量进行处理。如果音频样本是 16 位的，您可以使一个工作项呈现 8 个样本，然后通过矢量指令一次处理 8个。
二维数据
图像自然适合 OpenCL，因为它是像素的二维阵列。您可以通过将 1600 x 1200 像素的图像映射到 1600 x 1200 的 NDRange 进行处理。
工作项的总数是图像中的像素总量，即 1920000。
NDRange 分割为工作组，其中每个工作组也是一个二维阵列。工作组的数量必须正好分入 NDRange 中。
如果每个工作项处理一个像素，则 8 x 16 大小的工作组的大小为 128。这一工作组大小在 x 和 y 轴上正好融入到 NDRange 中。要处理该图像，您需要 15000 个工作组，每个包含 128 个工作项。
您可以矢量化这一示例，在单个矢量中处理所有颜色通道。如果通道是 8 位值，您可以在单个矢量中处理多个像素。如果每个矢量处理 4 个像素，这表示每个工作项处理 4 个像素，处理整个图像所需的工作项数量就小 4 倍。这意味着，您的 NDRange 可以减小到 400 x 1200，仅需要 3750 个工作组就能处理该图像。
三维数据
您可以时候用三维数据来模拟真实世界中材料的行为。例如，您可以通过在三维数据集中模拟应力，以此模拟建筑用混凝土的行为。您可以使用生成的数据来确定要承受特定负载所需的混凝土的大小和设计。
您可以在游戏中使用这一技巧来模拟对象的物理学。当对象破裂时，物理学模拟将使得破裂过程更加真实。
第 6 章
为 Mali GPU 重调现有 OpenCL 代码
本章介绍如何为 Mali GPU 优化现有的 OpenCL 代码。其中包含下列小节：
OpenCL 是一种可移植语言，但并不始终在性能上可移植。这意味着 OpenCL 可以在不同类型的计算设备上运作，但性能无法维持。
现有的 OpenCL 通常针对桌面 GPU 等特定架构进行了调节。要在 Mali GPU 上获得更佳的性能，您必须针对 Mali GPU 重调代码。
转换 OpenCL 代码以便在 Mali GPU 上优化运行的过程如下：
本节介绍为 Mali GPU 优化现有 OpenCL 代码的步骤。其中包含下列小节：
如果不是自己编写的代码，您必须分析代码，以便了解它的确切行为。
尝试了解如下所列：
这些问题的答案可以充当向导，帮助您删除设备相关的优化。
这些问题可能难以回答，因为高度优化的代码可能会非常复杂。
存在针对其他计算设备的优化，它们对 Mali GPU 没有作用，或者会减低性能。
要针对 Mali GPU 优化代码，您必须首先删除下列所有类型的优化，从而创建不限设备的参考实施；
对本地或专用内存的使用
Mali GPU 使用缓存而非本地或专用内存，因此您不必手动将数据移入或移出其中。删除来往于本地或专用内存的所有分配和复制。
屏障 来往于本地内存的数据传输通常与屏障同步。删除所有此用途的屏障。
缓存大小优化
某些代码优化读取和写入操作，以便数据能装入缓存行。这是有用的优化，但缓存行大小可能与 Mali GPU 的不同。删除这些类别的优化。
对标量的使用
许多 GPU 使用标量，而 Mali GPU 使用的是矢量。如果有基于标量的优化，请将其删除。
针对库冲突的修改
某些代码包含可避免库冲突的优化。请删除这些优化。
扭曲或波前
Mali GPU 不使用扭曲或波前。请删除任何相关的优化。
针对分叉线程的优化
Mali GPU 上的线程是独立的，无法分叉。如果代码包含针对分叉线程的优化或变通，请将其删除。
要针对 Mali GPU 优化代码，请参见第 7 章 为 Mali GPU 优化 OpenCL。
第 7 章
为 Mali GPU 优化 OpenCL
本章介绍一系列可在为 Mali-T600 系列 GPU 编写 OpenCL 时使用的优化。其中包含下列小节：
ARM 有如下建议：
Mali-T600 系列 GPU 中的着色器核心包含 128 位宽的矢量寄存器。矢量化内核中的算法，以便充分利用 Mali GPU 硬件。
OpenCL 通常运行在应用处理器和 GPU 具有单独内存的系统中。要在这些系统中使用 OpenCL，您必须分配缓冲区，以便从/向单独的内存复制数据。
配有 Mali GPU 的系统通常具有共享内存，因此您不必复制数据。不过，OpenCL 会假设内存是单独的，缓冲区分配也会涉及内存拷贝。这可造成浪费，因为拷贝耗时又耗电。
为了避免拷贝，可使用 OpenCL API 分配内存缓冲区，并使用 map() 和 unmap() 运算。这些运算可以让应用处理器和 Mali GPU 不经拷贝就能访问数据。
Mali GPU 无法访问由 malloc() 创建的内存缓冲区，因为它们不会映射到 Mali GPU 的内存空间中。其显示在图 7-1 中。
应用处理器
Application processor
malloc() 创建的缓冲区
Buffer created by malloc()
Global memory
Mali GPU
Mali GPU 无法访问内存缓冲区
Mali GPU cannot access memory buffer
图 7-1 malloc() 创建的内存缓冲区
Mali GPU 可以访问 clCreateBuffer(CL_MEM_USE_HOST_PTR) 创建的内存缓冲区，但这种方式创建的缓冲区必须由应用处理器将数据拷贝到其中。这些拷贝运算属于计算密集型，因此要尽可能避免这种缓冲区分配方式。图 7-2（第 7-5 页）中显示了这种缓冲区分配方式。
内存缓冲区要求内存拷贝
Memory buffer Requires memory copy
拷贝
Copy
clCreateBuffer() 创建的缓冲区
Buffer created by clCreateBuffer()
CL_MEM_USE_HOST_PTR
图 7-2 clCreateBuffer(CL_MEM_USE_HOST_PTR) 创建的内存缓冲区
Mali GPU 可以访问由 clCreateBuffer(CL_MEM_ALLOC_HOST_PTR) 创建的内存缓冲区。这是分配缓冲区的首选方式，因为不需要拷贝数据。图 7-3 中显示了这种缓冲区分配方式。
Mali GPU 和应用处理器都可访问内存缓冲区
Mali GPU and Application processor can both access memory buffer
CL_MEM_ALLOC_HOST_PTR
图 7-3 clCreateBuffer(CL_MEM_ALLOC_HOST_PTR) 创建的内存缓冲区
如果不这么做，回调可能会发生在更大一批工作结束时，晚于根据实际完成的工作而应发生的时间。
您可以通过减少和优化相关计算，来减小应用程序中串行部分的影响：
第 8 章
Mali OpenCL SDK
Mali OpenCL SDK 包含下列教程，帮助您了解 OpenCL 开发：
Hello World 教程
此教程提供了 OpenCL 和矢量化的基本介绍。
模板教程
此教程提供了 OpenCL 模板，您可以用作开发 OpenCL 应用程序的起点。
内存优化
内存优化目录中包含数据共享教程，该教程演示如何在 Mali-T600 系列 GPU 和应用处理器之间高效共享内存。
索贝尔滤波器教程
此教程演示如何使用索贝尔图像过滤器。这是一种简单的卷积滤波器，主要用于边缘检测算法。
FIR 浮点滤波器教程
此教程演示如何使用浮点有限输入响应 (FIR) 图像滤波器。可将它用于像素化或降噪。
Mandelbrot 教程
此教程演示如何利用计算 Mandelbrot 集合来产生分形图案。
SGEMM 教程
此教程演示如何使用 OpenCL 中的单精度普通矩阵乘法 (SGEMM)。
该 OpenCL SDK 可从 Mali 开发者中心获得。
http://www.malideveloper.arm.com
附录 A
OpenCL 数据类型
此附录列出了 OpenCL 中可用的数据类型。表 A-1 显示内置的标量数据类型。
表 A-1 内置标量数据类型
OpenCL 类型
API 类型
说明
bool
-
真 (1) 或假 (0)
char
cl_char
8 位有符号
unsigned char, uchar
cl_uchar
8 位无符号
short
cl_short
16 位有符号
unsigned short, ushort
cl_ushort
16 位无符号
int
cl_int
32 位有符号
unsigned int, uint
cl_uint
32 位无符号
long
cl_long
64 位有符号
unsigned long, ulong
cl_ulong
64 位无符号
float
cl_float
32 位浮点
half
cl_half
16 位浮点，仅用于存储
size_t
32 位或 64 位无符号整数
ptrdiff_t
表 A-1 内置标量数据类型（续）
intptr_t
有符号整数
uintptr_t
无符号整数
void
作废
表 A-2 显示内置的矢量数据类型。
表 A-2 内置矢量数据类型
charna
cl_charn
ucharn
cl_ucharn
shortn
cl_shortn
ushortn
cl_ushortn
intn
cl_intn
uintn
cl_uintn
longn
cl_longn
ulongn
cl_ulongn
floatn
cl_floatn
表 A-3 显示其他内置数据类型。
表 A-3 其他内置数据类型
image2d_t
2D 图像句柄
image3d_t
3D 图像句柄
sampler_t
采样器句柄
event_t
事件句柄
表 A-4 显示保留数据类型。
表 A-4 保留数据类型
booln
布尔矢量
double, doublen
64 位浮点，矢量
halfn
16 位浮点，矢量
quad, quadn
128 位浮点，矢量
表 A-4 保留数据类型（续）
complex half, complex halfn, imaginary half, imaginary halfn
16 位复杂，矢量
complex float, complex floatn, imaginary float, imaginary floatn
32 位复杂，矢量
complex double, complex doublen, imaginary double, imaginary doublen
64 位复杂，矢量
complex quad, complex quadn, imaginary quad, imaginary quadn
128 位复杂，矢量
floatnxm
n*m 矩阵，32 位浮点
doublenxm
n*m 矩阵，64 位浮点
long double, long doublen
64 位 - 128 位浮点，矢量
long long, long longnb
128 位有符号
unsigned long long, ulong long, ulonglongn
128 位无符号
附录 B
OpenCL 内置函数
此附录列出了 OpenCL 内置的函数。其中包含下列小节：
所列的函数具有相对速度等级。等级范围为 A 到 C，其中 A 最快。为获得最高的性能，请务必使用矢量，并尝试使用具有 A 或 B 等级的函数。
内存访问与算术运算的等级分开。A 级内存运算可能相当于 C 级算术运算。
表 B-1 列出工作项函数。
表 B-1 工作项函数
函数
速度
get_work_dim()
get_global_size()
get_global_id()
get_local_size()
get_local_id()
get_num_groups()
get_group_id()
get_global_offset()
表 B-2 列出数学函数。
表 B-2 数学函数
fabs()
acos()
B
acosh()
C
ceil()
acospi()
asinh()
fdim()
asin()
atanh()
fmax()
asinpi()
copysign()
fmin()
atan()
erfc()
mad()
atan2()
erf()
maxmag()
atanpi()
fmod()
minmag()
atan2pi()
fract()
rint()
cbrt()
frexp()
round()
cos()
hypot()
trunc()
cosh()
ilogb()
cospi()
ldexp()
exp()
lgamma()
exp2()
lgamma_r()
exp10()
log()
expml()
log10()
floor()
log1p()
fma()
logb()
log2()
modf()
pow()
nan()
pown()
nextafter()
powr()
remainder()
rsqrt()
remquo()
sin()
rootn()
sincos()
sinh()
sinpi()
tan()
sqrt()
tanh()
tanpi()
tgamma()
通常，大多数架构中都存在准确性和速度之间的权衡。Mali-T600 系列 GPU 全速实施数学函数的完整精度变体，因此您不必作此权衡。
数学函数的 half_ 和 native_ 变体供移植使用。它们的速度并不比精确变体快。请参见数学函数（第 B-3 页）。
表 B-3 列出了 half_ 和 native_ 数学函数。
表 B-3 half_ 数学函数
half_ 函数
native_ 函数
half_cos()
native_cos()
half_divide()
native_divide()
half_exp()
native_exp()
half_exp2()
native_exp2()
half_exp10()
native_exp10()
half_log()
native_log()
half_log2()
native_log2()
half_log10()
native_log10()
half_powr()
native_powr()
half_recip()
native_recip()
half_rsqrt()
native_rsqrt()
half_sin()
native_sin()
half_sqrt()
native_sqrt()
half_tan()
native_tan()
表 B-4列出整数函数。
表 B-4 整数函数
abs()
abs_diff()
add_sat()
hadd()
rhadd()
clz()
max()
min()
sub_sat()
mad24()，与 32 位乘法累积相同
mul24()，与 32 位乘法相同
clamp()
mad_hi()
mul_hi()
mad_sat()
rotate()
upsample()
表 B-5 列出公共函数。
表 B-5 公共函数
step()
degrees()
mix()
radians()
smoothstep()
sign()
表 B-6 列出几何函数。
表 B-6 几何函数
dot()
normalize()
fast_distance()
fast_length()
fast_normalize()
cross()
distance()
length()
表 B-7 列出关系函数。
表 B-7 关系函数
any()
all()
bitselect()
select()
isequal()
isnotequal()
isgreater()
isgreaterequal()
isless()
islessequal()
islessgreater()
isfinite()
isinf()
isnan()
isnormal()
isordered()
isunordered()
signbit()
表 B-8 列出矢量数据加载和存储函数。它们都是 A 级速度。
表 B-8 矢量数据加载和存储函数
vload()
vstore()
vload_half()
vstore_half()
vloada_half()
vstorea_half()
表 B-9 列出同步函数。它们没有速度等级，因为同步函数需要等待多个线程完成。其花费的时间决定了这些函数在应用程序中的时间长度。
表 B-9 同步函数
barrier()
mem_fence()
read_mem_fence()
write_mem_fence()
表 B-10 列出异步复制函数。它们没有速度等级，因为复制速度取决于所复制的数据的大小。
表 B-10 异步复制函数
async_work_group_copy()
async_work_group_strided_copy()
wait_group_events()
prefetch()
表 B-11 列出原子函数。
表 B-11 原子函数
atomic_add()
atomic_sub()
atomic_xchg()
atomic_inc()
atomic_dec()
atomic_cmpxchg()
atomic_min()
atomic_max()
atomic_and()
atomic_or()
atomic_xor()
表 B-12 列出其他矢量函数。
表 B-12 其他矢量函数
vec_step()
shuffle()
shuffle2()
表 B-13 列出图像读取和写入函数。
表 B-13 图像读取和写入函数
read_imagef()
read_imagei()
read_imageui()
write_imagef()
write_imagei()
write_imageui()
get_image_width()
get_image_height()
get_image_depth()
get_image_channel_data_type()
get_image_channel_order()
get_image_dim()
附录 C
OpenCL 扩展
Mali OpenCL 驱动程序支持下列扩展：