OpenCL 编程入 门
日期 ： 2014/11/6
OpenCL 编程入门
1. 概述
1.1 引言
OpenCL 的文章网络上有很 多， 从入门 到各种精 通都 有。 本文的目的， 是用尽量 简洁的
叙述， 建立对 OpenCL 相对全面的概念。 本文适 合刚 开始接触异 构计算和 OpenCL 的编程人
员，或想要 了解 OpenCL 的人。
如果只是想 知道如何 写一个 Hello World ，可以直接跳到：2.3 节。
1.2 OpenCL 概念
OpenCL 是由 Khronos Group 针对异构 计算设备（heterogeneous device ）进行并行 运算
所设计的标 准 API 以及程式语言。
OpenCL 程序分 成成两部分 ：一部 分是在设 备上执行 的（ 例如 GPU ） ，另一部分是 在主
机上运行的 。 在设备 上执行的 程序 就是 实现 “ 异构” 和 “并行计 算” 的部 分。 为 了能在设 备
上执行代码 ，程序员 需要写一 个特殊的 函数（kernel 函数） 。这个 函数需要 使用 OpenCL 语
言编写。OpenCL 语言采用了 C 语言的一部分 加上一 些约束、 关 键字和数 据类型 。 在主机 上
运行的程序 用 OpenCL 的 API 管理设备上运行的程序 。主机程序 的 API 用 C 语言编写，也
有 C++ 、Java 、Python 等高级语言接口 。
2. OpenCL 使用和编程
2.1 OpenCL 的使用模 型
网上找到下 面这张图 ， 比较简 洁地说明 了 OpenCL 的 使用。
1 OpenCL 编程入 门
从图中可以 看出：
1. 异构计算设 备，可以 是 CPU 或 GPU 。现在也有支持 OpenCL 的 FPGA 设备 和至强
融核协处理 设备（MIC ） 。
2. OpenCL 的 API 通过 Context （环境上下文）联 系在一 起。
3. 运行设备端 的程序， 经过了编 译-> 设置参数-> 运 行等 步骤。
2.2 OpenCL 的一些基 本概念
要理解 OpenCL 编程， 首 先需要理 解一些基 本概念。 当然不喜欢 阅读概念 定义的同 学可
以先跳过这 一节，直 接看编程 的内容 （2.3 节） 。
Platform ( 平台) ： 主机加 上 OpenCL 框架管理下的若 干设备构成 了这个平 台， 通过这 个
平台， 应用程序 可以与设 备共享 资源并在 设备上执 行 kernel 。 实际使用中基本上一个厂 商对
应一个 Platform ，比如 Intel, AMD 都是这样。
Device （设备） ： 官方的解 释是计算 单元 （Compute Units ） 的集合。 举例来说，GPU 是
典型的 device 。Intel 和 AMD 的多核 CPU 也提供 OpenCL 接口，所以也可以 作为 Device 。
Context （上下文） ：OpenCL 的 Platform 上共享和使用资源的环境， 包 括 kernel 、device 、
memory objects 、command queue 等。使用中一般一 个 Platform 对应一个 Context 。
Program ：OpenCL 程序，由 kernel 函数、其他函数 和声明等组 成。
Kernel （核函数） ：可以 从主机端 调用，运 行在设备 端的函数。
2
OpenCL 编程入 门
Memory Object （内存对象） ： 在 主机和 设备之间 传递 数据的对象 ， 一 般映射到 OpenCL
程序 中的 global memory 。 有两种具体的类型 ：Buffer Object （缓存对象） 和 Image Object （图
像对象） 。
Command Queue （指令队列） ： 在指定 设备上管 理多 个指令 （Command ） 。 队列里指令
执行可以顺 序也可以 乱序。一 个设备可 以对应多 个指 令队列。
NDRange ：主 机端运行 设备端 kernel 函数的主要接口 。
注： Platform, Device 和 Context 的对应关系在实际使用 中可以有所 变化。 理 论上 Platform
与 Context 是 一 对 多 的 关 系 ， Device 与 Context 可 以 是 多 对 多 的 关 系 。 如果用
clCreateContextFromType 创建 Context ，可以忽略具 体的 device id ；如果用 clCreateContext
创建，可以 忽略具体 的 platform id 。实际使用中一般 只用一个 Device 对一个 Context 。多种
对应组合的 实验我没 有做过， 有经验的 同学可以 补充 。
2.3 OpenCL 的编程步 骤
OpenCL 主机端的编程比一 般的 Hello World 要复杂一些 ， 这里归纳 为 12 个 步骤。这个
过程不太需 要变化， 所以写过 一次以后 基本上只 要复 制就可以了 。
3 OpenCL 编程入 门
2.3.1 获取 Platform
调用两次 clGetPlatformIDs 函数， 第一次获取可用 的 平台数量， 第二 次获取一 个可用的
平台。
/ step 1: get platform /
cl_uint num_platforms;
ret = clGetPlatformIDs(0, NULL, &num_platforms); // get platform number
if ((CL_SUCCESS != ret) || (num_platforms < 1))
{
cout << "Error getting platform number: " << ret << endl;
return 0;
}
cl_platform_id platform_id = NULL;
ret = clGetPlatformIDs(1, &platform_id, NULL); // get first platform id
if (CL_SUCCESS != ret)
{
cout << "Error getting platform id: " << ret << endl;
4
OpenCL 编程入 门
return 0;
}
2.3.2 获取 Device
调用两次 clGetDeviceIDs 函数， 第一次获取 可用的设 备数量， 第二 次获取一 个可用的 设
备。
/ step 2: get device /
cl_uint num_devices;
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices);
if ((CL_SUCCESS != ret) || (num_devices < 1))
{
cout << "Error getting GPU device number: " << ret << endl;
return 0;
}
cl_device_id device_id = NULL;
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error getting GPU device id: " << ret << endl;
return 0;
}
2.3.3 创建 Context
这里以 clCreateContext 为例。也可以用 clCreateContextFromType 。
/ step 3: create context /
cl_context_properties props[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id,
0 };
cl_context context = NULL;
context = clCreateContext(props, 1, &device_id, NULL, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == context))
{
cout << "Error creating context: " << ret << endl;
return 0;
}
2.3.4 创建 Command Queue
一个 Device 有多个 Command Queue 。 把有数据关联关 系的 Command 放到同一个 Queue
里执行， 没有关联 关系 的 Command 放到不同的 Queue 。Command Queue 之间并行执行， 由
具体的 Device 来安排。简单的用法 一个 Device 只用一个 Command Queue 。
/ step 4: create command queue /
cl_command_queue command_queue = NULL;
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
if ((CL_SUCCESS != ret) || (NULL == command_queue))
5 OpenCL 编程入 门
{
cout << "Error creating command queue: " << ret << endl;
return 0;
}
2.3.5 创建 Memory Object
Memory Object 有两种类型： 缓存 （Buffer ） 和图像 （Image）。 Buffer 类似数组。 理论上
Buffer 可以代 替 Image 。OpenCL 对 Image 有一些格式识别和处 理，有助 于简化代 码。有些
类型的 Device 可能对 Image 这种 Z 型的访存模式有硬件性能优 化。
Buffer 由上下文 context 创建，这样上下文管 理的多个 设备就会共 享 Buffer 中的数据。
注意示例中 用了 CL_MEM_USE_HOST_PTR ，即内存分配在 Host 端。理论上 Device
端分配的内 存运行速 度更快， 但实际应 用中往往 数据 拷贝消耗性 能比较多 ，而用 Host 端内
存能做到类 似“零拷 贝”的效 果，所以 建议将内 存分 配在 Host 端。另 外，有 些 OpenCL 设
备的 Host 与 Device 之间 Cache 交换做得比较好， 所 以 内存对象在 DEVICE 端还是 HOST 端
在 运行速度 上的差异 可能很小 。
/ step 5: create memory object /
cl_mem mem_obj = NULL;
int host_buffer = NULL;
const int ARRAY_SIZE = 1000;
const int BUF_SIZE = ARRAY_SIZE sizeof(int);
// create and init host buffer
host_buffer = (int )malloc(BUF_SIZE);
init_buf(host_buffer, ARRAY_SIZE);
// create opencl memory object using host ptr
mem_obj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, BUF_SIZE,
host_buffer, &ret);
if ((CL_SUCCESS != ret) || (NULL == mem_obj))
{
cout << "Error creating command queue: " << ret << endl;
return 0;
}
2.3.6 创建编译 Program
程序对象可 以从源文 件或二进 制文件创 建。不管 哪种 方式都要经 过编译之 后才能使 用。
对二进制文 件来说， 这里的 clBuildProgram 更像是程序加载的过 程。
/ step 6: create program /
char kernelSource =
"__kernel void test(__global int pInOut)\n"
"{\n"
" int index = get_global_id(0);\n"
6
OpenCL 编程入 门
" pInOut[index] += pInOut[index];\n"
"}\n";
cl_program program = NULL;
// create program
program = clCreateProgramWithSource(context, 1, (const char)&kernelSource, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == program))
{
cout << "Error creating program: " << ret << endl;
return 0;
}
// build program
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error building program: " << ret << endl;
return 0;
}
2.3.7 创建 Kernel
从 Program 对象生成 kernel （核函数）对象，表示设 备程序的入 口。
/ step 7: create kernel /
cl_kernel kernel = NULL;
kernel = clCreateKernel(program, "test", &ret);
if ((CL_SUCCESS != ret) || (NULL == kernel))
{
cout << "Error creating kernel: " << ret << endl;
return 0;
}
2.3.8 设置 Kernel 参数
用 clSetKernelArg 设置 Kernel 参数。
/ step 8: set kernel arguments /
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void)&mem_obj);
if (CL_SUCCESS != ret)
{
cout << "Error setting kernel argument: " << ret << endl;
return 0;
}
2.3.9 设置 Group Size
一个 OpenCL 运行任务中并行计算 的单位是 work-item。而 work-item 的组织形式就由维
数（dim ）,各维度尺寸（ global_work_size ）和分组方式（ local_work_size ）等参数决定。
7 OpenCL 编程入 门
这些参数的 含义在后 续章节中 描述。 这里简单 注意几 点： 一是 维数 ， 一般的 OpenCL 设备最
大支持维数 为 3 ， 可以查 询 CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS 来获取 ； 二是
local_work_size ，合适的 size 与具体 OpenCL 设备的最大并发资 源有关， 可以不设 置。
/ step 9: set work group size /
cl_uint work_dim = 3; // in most opencl device, max dimition is 3
size_t global_work_size[] = { ARRAY_SIZE, 1, 1 };
size_t local_work_size = NULL; // let opencl device determine how to break work items into
work groups
2.3.10 Kernel 入 队执行
用 clEnqueueNDRangeKernel 将 kernel 对象，以及 work-item 参数放入命令队列中进行
执行 。
/ step 10: run kernel /
ret = clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL, global_work_size,
local_work_size, 0, NULL, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error enqueue NDRange: " << ret << endl;
return 0;
}
2.3.11 读取结果
推荐用 clEnqueueMapBuffer 来读取结果。这样做的好 处是，当内 存分配在 HOST 端，
map 操作只需将 Device Cache 中的内容同步到 host 内存，不需 要内存拷 贝。所以 消耗的时
间要比 clEnqueueReadBuffer 小。
/ step 11: get result /
int device_buffer = (int )clEnqueueMapBuffer(command_queue, mem_obj, CL_TRUE, CL_MAP_READ
| CL_MAP_WRITE, 0, BUF_SIZE, 0, NULL, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == device_buffer))
{
cout << "Error map buffer: " << ret << endl;
return 0;
}
2.3.12 释放资源
这没什么好 讲的，注 意哪些资 源需要释 放就可以 了。
/ step 12: release all resources /
if (NULL != kernel) clReleaseKernel(kernel);
if (NULL != program) clReleaseProgram(program);
if (NULL != mem_obj) clReleaseMemObject(mem_obj);
if (NULL != command_queue) clReleaseCommandQueue(command_queue);
if (NULL != context) clReleaseContext(context);
if (NULL != host_buffer) free(host_buffer);
8
OpenCL 编程入 门
2.4 完整的 Hello World 示 例代码
完整的 Hello World 示例代码。
#include
#include
#include
#include
#include
using namespace std;
void check_result(const int buf, const int len)
{
int i;
for (i = 0; i < len; i++)
{
if (buf[i] != (i + 1) 2)
{
cout << "Result error!" << endl;
break;
}
}
if (i == len) cout << "Result ok." << endl;
}
void init_buf(int buf, int len)
{
int i;
for (i = 0; i < len; i++)
{
buf[i] = i + 1;
}
}
int main(void)
{
cl_int ret;
/ step 1: get platform /
cl_uint num_platforms;
ret = clGetPlatformIDs(0, NULL, &num_platforms); // get platform number
if ((CL_SUCCESS != ret) || (num_platforms < 1))
{
cout << "Error getting platform number: " << ret << endl;
return 0;
9 OpenCL 编程入 门
}
cl_platform_id platform_id = NULL;
ret = clGetPlatformIDs(1, &platform_id, NULL); // get first platform id
if (CL_SUCCESS != ret)
{
cout << "Error getting platform id: " << ret << endl;
return 0;
}
/ step 2: get device /
cl_uint num_devices;
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices);
if ((CL_SUCCESS != ret) || (num_devices < 1))
{
cout << "Error getting GPU device number: " << ret << endl;
return 0;
}
cl_device_id device_id = NULL;
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error getting GPU device id: " << ret << endl;
return 0;
}
/ step 3: create context /
cl_context_properties props[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id,
0 };
cl_context context = NULL;
context = clCreateContext(props, 1, &device_id, NULL, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == context))
{
cout << "Error creating context: " << ret << endl;
return 0;
}
/ step 4: create command queue /
cl_command_queue command_queue = NULL;
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
if ((CL_SUCCESS != ret) || (NULL == command_queue))
{
cout << "Error creating command queue: " << ret << endl;
return 0;
}
10
OpenCL 编程入 门
/ step 5: create memory object /
cl_mem mem_obj = NULL;
int host_buffer = NULL;
const int ARRAY_SIZE = 1000;
const int BUF_SIZE = ARRAY_SIZE sizeof(int);
// create and init host buffer
host_buffer = (int )malloc(BUF_SIZE);
init_buf(host_buffer, ARRAY_SIZE);
// create opencl memory object using host ptr
mem_obj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, BUF_SIZE,
host_buffer, &ret);
if ((CL_SUCCESS != ret) || (NULL == mem_obj))
{
cout << "Error creating command queue: " << ret << endl;
return 0;
}
/ step 6: create program /
char kernelSource =
"__kernel void test(__global int pInOut)\n"
"{\n"
" int index = get_global_id(0);\n"
" pInOut[index] += pInOut[index];\n"
"}\n";
cl_program program = NULL;
// create program
program = clCreateProgramWithSource(context, 1, (const char)&kernelSource, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == program))
{
cout << "Error creating program: " << ret << endl;
return 0;
}
// build program
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error building program: " << ret << endl;
return 0;
}
/ step 7: create kernel /
11 OpenCL 编程入 门
cl_kernel kernel = NULL;
kernel = clCreateKernel(program, "test", &ret);
if ((CL_SUCCESS != ret) || (NULL == kernel))
{
cout << "Error creating kernel: " << ret << endl;
return 0;
}
/ step 8: set kernel arguments /
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void)&mem_obj);
if (CL_SUCCESS != ret)
{
cout << "Error setting kernel argument: " << ret << endl;
return 0;
}
/ step 9: set work group size /
cl_uint work_dim = 3; // in most opencl device, max dimition is 3
size_t global_work_size[] = { ARRAY_SIZE, 1, 1 };
size_t local_work_size = NULL; // let opencl device determine how to break work items into
work groups
/ step 10: run kernel /
ret = clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL, global_work_size,
local_work_size, 0, NULL, NULL);
if (CL_SUCCESS != ret)
{
cout << "Error enqueue NDRange: " << ret << endl;
return 0;
}
/ step 11: get result /
int device_buffer = (int )clEnqueueMapBuffer(command_queue, mem_obj, CL_TRUE, CL_MAP_READ
| CL_MAP_WRITE, 0, BUF_SIZE, 0, NULL, NULL, &ret);
if ((CL_SUCCESS != ret) || (NULL == device_buffer))
{
cout << "Error map buffer: " << ret << endl;
return 0;
}
// check result
check_result(device_buffer, ARRAY_SIZE);
/ step 12: release all resources /
if (NULL != kernel) clReleaseKernel(kernel);
if (NULL != program) clReleaseProgram(program);
if (NULL != mem_obj) clReleaseMemObject(mem_obj);
12
OpenCL 编程入 门
if (NULL != command_queue) clReleaseCommandQueue(command_queue);
if (NULL != context) clReleaseContext(context);
if (NULL != host_buffer) free(host_buffer);
return 0;
}
与这个例子 运算结果 等价的串 行程序为 ：
for (int i = 0; i < ARRAY_SIZE; i++)
{
a[i] += a[i];
}
3. GPU 并 行运 行 的原 理
3.1 OpenCL 执行模型
前面的例子 中已经知 道，HOST 通过 clEnqueueNDRange 调用， 将执行 核函数 （Kernel ）
的命令放到 command queue 中，使得 OpenCL 程序在 DEVICE 端获得执行 。 一次执 行需要
设置两部分 参数，对 应例子中 的 step 8 和 step 9 ：先设置核函数 参数，再 设置分组 大小。
我们先来看 核函数。
__kernel void test(__global int pInOut)
{
int index = get_global_id(0);
pInOut[index] += pInOut[index];
}
例子中的核 函数完成 了一个简 单的运算 ，即对每个 work-item ，将数组中对应成员的值
放大一倍。 核函数的参数 只有 一个， 就是用 clSetKernelArg 设置的数组指针。 分组信息 在核
函数中表现 为 work-item 的序号， 在例子中通过 get_global_id 获取。 可以看到一个运 算任务
中，每个 work-item 有同样的参数，有不 同的序号 。Work-item 之间并行执行。
Work-item 的数量以及分组方式， 由 clEnqueueNDRange 中的 work_dim, global_work_size
和 local_work_size 等参数 确定。
假设 work_dim = 3, global_work_size = { x, y, z } ，则 work-item 总的数量就等于 x y
z 。 例子 中 x = ARRAY_SIZE = 1000, y = 1, z = 1 ，所以 总的 work-item 数量是 1000。
那 local_work_size 有什么用处呢 ？这里要 引入 OpenCL 里一个重要的 “Work Group”概
念。 一个 OpenCL 设备可能由多个 计算单元 组成， 而 每个计算单 元由多个 处理单元组成； 一
个 “work item ” ，是运行在某个 处理单元上的 kernel 的实例。 那 么，同 一个计算 单元上的 多
个处理单元 上的 work-item 组成一个 work-group 是就 很自然的事 情了。 如 果用 SIMD(Single
Instruction Multiple Data) 来解释 ， 同一个 work-group 里的多个 work-item ，用不同的操作数
13 OpenCL 编程入 门
并行运行同 一组指令 。而 local_work_size 就是在 work_dim 定义的维度上，给出构成 work
group 的每个维度上的分组 方式。Local_id 就是 work item 在每个 work group 里的序号。
引用 OpenCL 1.2 标准文档 里的图例（ 二维） ：
所以标识一个 work-item 的 方 法 有 两 种 ： 一 种 是 通 过 get_global_id ； 另 一 种 是 通 过
get_group_id 加 get_local_id 。 两种之间的换算 关系是 ：global_id = group_id local_work_size
+ local_id 。
使用 work-group 的意义是什么？个人 认为主要 还是优 化内存访问 。 这点 在 “OpenCL 内
存模型”中 解释。
clEnqueueNDRange 中的另一个参数 global_work_offset ，允许各个维 度在计算 global id
时，可以指 定起始的 id 。这个参 数 可以将 运算任务 映 射到数据的 某个片段 上 。
3.2 OpenCL 内存模型
内存模型看 下面这个 图例：
14
OpenCL 编程入 门
用核函数中 的内存变 量来简单 地解释：用 clCreateBuffer 创建、用 clSetKernelArg 传递
的数据在 global memory 和 constant memory 中； 核函数中的寄存器 变量在 private memory 中；
核函数的内 部变量、 缓存等， 在 local memory 中。
图例中可以 看到 Device 并不直接访问 global memory ， 而是通过 Cache 来访问。 可以想
象当同时运 行的 work-item ，使用的内存都在同 一块 cache 中，则内存吞吐的 效率最高 。对
应到 work group 中， 就是在程序设计上 尽量使 同一个 work group 中的 work item 操作连续的
内存，以提 高访存效 率。
4. OpenCL 的 一 些 重 要资 源
http://www.khronos.org/opencl/ 组织的主页
https://developer.nvidia.com/opencl/ N 家的主页
http://developer.amd.com/tools-and-sdks/opencl-zone/ A 家的主页
https://software.intel.com/en-us/intel-opencl I 家的主页
https://www.khronos.org/registry/cl/specs/opencl-1.2.pdf OpenCL 1.2 标准规范
15 |
|
