OpenCL专题01：开发配置_opencl配置

前言：为什么要用GPU编程

我们平时跑的大部分程序都是在CPU上的，CPU的每个核很强大，有很大的缓存，分支处理也比较快，功能全面，各种程序都可以运行，但带来的缺点就是一个CPU里面放不了太多核（最多几十个），并行计算能力不是特别高。

GPU和CPU不同。GPU的核很多，能做到几千个，所以并发计算能力很强。但是相比来说每个核的主频低，缓存少，单核比CPU的一个核差很多。而且GPU的架构是SIMT（Single instruction multiple thread），所以并不是每个核完全独立运行的。由于GPU的这些特点，它很适合做那种并行程度高的运算，比如矩阵运算。而机器学习领域，特别如人工神经网络这类模型，矩阵运算是最主要的运算，自然就会想到使用GPU来提高性能。

1. OpenCL简介

2008年6月，苹果首先提出了OpenCL（Open Computing Language）规范，旨在提供一个通用的开放API，在此基础上开发GPU通用计算软件。随后，Khronos Group宣布成立GPU通用计算开放行业标准工作组，以苹果的提案为基础创立OpenCL行业规范。业界最主要的图形或者计算相关的厂商都是Khronos的成员。

OpenCL是一个为异构平台编写并行计算程序的框架，此异构平台可由CPU，GPU或其他类型的处理器组成。OpenCL提供了基于任务分割和数据分割的并行计算机制，开发人员可以利用GPU和CPU的计算能力，把GPU和CPU异构的系统运用在很多并行计算的领域。

OpenCL程序分成两部分：一部分是在GPU上执行，另一部分是在CPU上运行。在GPU上执行的程序就是实现“异构”和“并行计算”的部分。为了能在GPU上执行代码，程序员需要写一个特殊的kernel函数，这个函数需要使用OpenCL语言编写。OpenCL语言是在C++语言的基础上增加了一些约束、关键字和数据类型。在CPU上运行的程序通过OpenCL的API管理GPU上运行的程序。CPU上的程序可以用C++、Java、Python等高级语言编写。

OpenCL开发面向异构平台的应用需要完成以下步骤：

• 发现构成异构系统的组件。
• 探查这些组件的特征， 使软件能够适应不同硬件单元的特定特性。
• 建立将在平台上运行的指令块（内核）。
• 建立并管理计算中涉及的内存对象。
• 在系统中正确的组件上按正确的顺序执行内核。
• 收集最终结果。

2. OpenCL的一些基本概念

2.1 Platform（平台）

CPU加上OpenCL框架管理下的若干GPU构成了这个平台，通过这个平台，应用程序可以与GPU共享资源，并在GPU上执行kernel。实际使用中，一般是一个厂商对应一个Platform，比如Intel，NVIDIA，AMD。

2.2 Device（设备）、计算单元、处理单元

计算单元的集合，比如GPU就是典型的Device。Intel和AMD的CPU也提供OpenCL接口，也可以作为Device。
Device进一步划分为计算单元，计算单元还可以进一步划分为一个或多个处理单元。Device上的计算都在处理单元中完成。

2.3 Contex（上下文）

虽然OpenCL应用的计算工作是在GPU上进行的，但是CPU定义了内核，而且为内核的建立创建了上下文，定义了NDRange和队列（队列存储了内核如何执行以及何时执行的细节），所有这些重要函数都包含在OpenCL定义的API中。
CPU的第一个任务是为 OpenCL应用定义上下文。顾名思义， 上下文定义了一个环境， 内核就在这个环境中定义和执行。更准确地说， 由以下资源定义上下文：

• 设备（device）：CPU使用的 GPU集合。
• 内核（kernel）：在GPU上运行的OpenCL函数。
• 程序对象（program object）：实现内核的程序源代码和可执行文件。
• 内存对象（memory object）：内存中对GPU可见的一组对象， 包含可以由内核实例处理的值。

2.4 Program Object（程序对象）

OpenCL程序，由kernel函数、其它函数和声明等组成。

2.5 Kernel（核函数）

可以从CPU调用，运行在GPU上的函数。

2.6 Memory Object（内存对象）

在CPU和GPU之间传递数据的对象，一般映射到OpenCL程序中的global memory。有两种具体的类型：Buffer Object（缓存对象）和Image Object（图像对象）。
OpenCL 定义了两种类型的内存对象：

• Buffer Object：内核可用的一个连续的内存区。程序员可以将数据结构映射到这个缓冲区， 并通过指针访问缓冲区。
• Image Object:：仅限于存储图像。图像存储格式可以进行优化来满足一个特定GPU的需要。

2.7 Command Queue（指令队列）

在指定GPU上管理多个指令，队列里指令执行顺序可以顺序也可以乱序。一个GPU可以对应多个指令队列。
CPU与GPU之间的交互是通过命令完成的， 这些命令由CPU提交给Command Queue。这些命令会在Command Queue中等待，直到在GPU上执行。Command Queue由CPU创建， 并在定义上下文之后关联到一个GPU。CPU将命令放入命令队列， 然后调度这些命令在关联GPU上执行。OpenCL支持 3 种类型的命令：

• 内核执行命令（kernel execution command）：在 OpenCL 设备的处理单元上执行内核。
• 内存命令（ memory command ）：在宿主机和不同内存对象之间传递数据， 在内存对象之间移动数据， 或者将内存对象映射到宿主机地址空间， 或者从宿主机地址空间解映射
• 同步命令（synchronization command ）对命令执行的顺序施加约束。

一个典型的CPU程序，需要定义上下文、命令队列、内存、程序对象和CPU所需要的数据结构。内存对象从CPU转移到GPU上，内核参数关联到内存对象，然后提交到Command Queue执行。内核完成工作时，计算中生成的内存对象可能会再复制到CPU。

2.8 NDRange（N维索引空间）

CPU发出命令，提交内核在一个GPU上执行。OpenCL运行时会创建一个整数索引空间，索引坐标对应执行内核的一个实例，各个执行内核的实例称为一个工作项，其对应的坐标就是工作项的全局ID。提交内核执行的命令相应地会创建一个工作项集合， 其中各个工作项使用内核定义的同样的指令序列。尽管指令序列是相同的， 但是由于代码中的分支语句或者通过全局ID选择的数据可能不同， 因此各个工作项的行为可能不同。
工作项组织为工作组。工作组提供了对索引空间更粗粒度的分解，跨越整个全局索引空间。工作项可由自身的全局ID或者工作组ID+其局部ID唯一标识。给定工作组中的工作项会在一个计算单元的处理单元上并发执行。
索引空间是一个N维的值网格，也称NDRange，下图给出一个具体的索引空间的例子：

2.9 编程模型

OpenCL 定义了两种不同的编程模型：任务并行和数据并行。

• 数据并行编程模型
  适合采用数据并行编程模型的问题都与数据结构有关， 这些数据结构的元素可以并发更新。OpenCL提供了层次结构的数据并行性：工作组中工作项的数据并行再加上工作组层次的数据并行。OpenCL规范讨论了这种数据并行形式的两个变种。在显式模式（explicit model）中， 程序员负责显式地定义工作组的大小。 利用第二个模型， 即隐式模型（implicit model）， 程序员只需定义NDRange空间， 由系统选择工作组。总之， 数据并行很自然地切合了OpenCL执行模型。 这个模型是层次结构， 因为数据并行计算（工作项）可能包括矢量指令（SIMD），作为更大规模的块级数据并行（工作组）的一部分。所有这些工作结合起来为表述数据并行算法创建了一个很好的环境。
• 任务并行编程模型
  OpenCL执行模型被设计为以数据并行作为主要目标。不过这个模型还支持大量任务并行算法。OpenCL将任务定义为单个工作项执行的内核， 而不考虑OpenCL应用中其他内核使用的NDRange。 如果程序员所希望的并发位来自于任务， 就会使用这个模型。 例如， 并发性可能只是通过矢量类型上的矢量操作来表述。

3. 编译运行第一个OpenCL程序

3.1 在Ubuntu上安装OpenCL

以下命令将安装OpenCL的开发环境：

apt install ocl-icd-* opencl-headers
1

安装之后，我们可以通过clinfo命令来查看当前计算机上存在的平台、设备、计算单元、处理单元等信息。clinfo可以通过以下命令来安装：

apt install clinfo
1

若计算机上安装了独立显卡，clinfo则可查看到存在若干个platform，以及每个platform下有的device、work group、work item等信息。
若没有安装独立显卡，则Intel的CPU也支持核心显卡，可通过如下命令安装核心显卡驱动：

apt purge beignet
apt install beignet
1
2

安装成功后，clinfo可查看到一个名为Intel Gen OCL Driver的platform。

疑难杂症
有些计算机上安装了独立显卡，但Ubuntu操作系统只启用了核心显卡。为了使用独立显卡，需要给独立显卡安装驱动。
输入命令：ubuntu-drivers devices，会显示系统推荐的驱动，
输入命令：apt install nvidia-推荐的数字，则可安装对应驱动。
驱动安装完后，再通过clinfo查看，则可看到增加了一个platform。

3.2 使用VSCode编译第一个OpenCL程序

3.2.1 手工编译

OpenCL程序是C/C++代码，需要使用g++编译器进行编译、链接和生成可执行程序。假设OpenCL程序已经写好（第4节给出一个基础OpenCL程序，先直接拷贝建立main.cpp文件即可，下一讲再介绍OpenCL程序的原理）。
我们首先要给计算机安装g++编译器，在OpenCL代码中需要#include <CL/cl.h>头文件，该头文件在安装OpenCL驱动时已经安装到g++默认搜索头文件的目录下：/usr/include/，所以直接通过#include <CL/cl.h>即可引入该头文件。
该头文件中定义的函数在链接库中定义，编译时需要链接对应的库，路径为：/lib/x86_64-linux-gnu，库名为OpenCL。

g++的链接命令
库文件名不等于库名，库文件名的命名规则为lib+<库名>+.so，所以OpenCL库对应的库文件为libOpenCL.so。
在编译时，为了链接指定的库，我们需要用到-L和-l参数。-L后指定库所在的目录，-l后跟库名，因此编译时，需要在g++命令中添加-L /lib/x86_64-linux-gnu -l OpenCL进行链接。

假设OpenCL代码文件名为main.cpp，则我们通过如下命令进行编译：

g++ -fdiagnostics-color=always -g main.cpp -o main -L /lib/x86_64-linux-gnu -l OpenCL -w
1

-g指定可调试；-o手动指定输出文件的路径，即后面跟着的路径；-L和-l参数前面已经解释过，指定库文件的路径和库名；-w指定编译时不显示告警信息。

3.2.2 VSCode配置默认生成任务

VSCode是目前使用最广泛的跨平台跨语言通用IDE环境。通过简单配置，可实现自动编译、断点调试等功能。主要是通过两个配置文件实现的。
在VSCode中，打开需要编译的cpp文件，点击“终端”>“配置默认生成任务”，则会在.vscode目录下新建一个tasks.json文件，可将如下内容拷贝覆盖：

{
	"version": "2.0.0",
	"tasks": [
		{
			"type": "cppbuild",
			"label": "C/C++: cpp 生成活动文件",
			"command": "/usr/bin/g++",
			"args": [
				"-fdiagnostics-color=always",
				"-g",
				"${workspaceFolder}/source/*.cpp",
				"-o",
				"${workspaceFolder}/out/${workspaceRootFolderName}",
				"-L",
				"/lib/x86_64-linux-gnu",
				"-l",
				"OpenCL",
				"-w"//编译时不显示任何警告，-Wall是显示所有警告
			]
		}
	]
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22

command字段指定g++编译器的位置，args字段指定g++命令的参数。
我个人开发项目的习惯是建立三个文件夹，include文件夹下放各种.h文件，source文件夹下放.cpp文件，out文件夹用于项目编译输出，另外还有一个VSCode的配置文件夹.vscode。
在上面这段tasks.jsp配置中，我们指定每次自动生成任务时，对${workspaceFolder}/source/下的所有cpp文件进行编译，并输出到${workspaceFolder}/out/目录下，以${workspaceRootFolderName}命名输出的文件。${workspaceFolder}符号在配置文件中会被替换为VSCode项目所在的目录名。

配置完毕后，我们点击“终端”>“运行生成任务”则可自动编译链接生成。

3.2.3 VSCode配置调试

我们可以配置VSCode使用gdb调试程序。首先需要在系统中安装gdb，通过apt install gdb即可完成安装。
在VSCode中，点击“运行”>“添加配置”则可在.vscode目录下新建launch.json文件，该文件则用于配置调试信息，可将如下内容复制覆盖。

{
   "version": "0.2.0",
   "configurations": [
       {
           "name": "(gdb) 启动",
           "type": "cppdbg",
           "request": "launch",
           "program": "${workspaceFolder}/out/${workspaceRootFolderName}",
           "args": [],
           "stopAtEntry": false,
           "cwd": "${fileDirname}",
           "environment": [],
           "externalConsole": false,
           "MIMode": "gdb",
           "setupCommands": [
               {
                   "description": "为 gdb 启用整齐打印",
                   "text": "-enable-pretty-printing",
                   "ignoreFailures": true
               },
               {
                   "description":  "将反汇编风格设置为 Intel",
                   "text": "-gdb-set disassembly-flavor intel",
                   "ignoreFailures": true
               }
           ],
           "preLaunchTask": "C/C++: cpp 生成活动文件"
       }
   ]
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

需要注意的是，"preLaunchTask"字段配置的是预执行任务，在调试前需要启动生成，因此配置的内容与tasks.json中的"label"字段内容一致，VSCode通过label来判断任务。

至此，我们可以通过VSCode设置断点和启动调试。

4. OpenCL Hello World

#define CL_TARGET_OPENCL_VERSION 220

#include <stdio.h>
#include <string.h>
#include <stdlib.h>
#include <iostream>
#include <CL/cl.h>

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;
    }
    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;
    }
    /** 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 */
    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);
    if (NULL != command_queue)
        clReleaseCommandQueue(command_queue);
    if (NULL != context)
        clReleaseContext(context);
    if (NULL != host_buffer)
        free(host_buffer);
    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168