OpenCL
1. OpenCL简介
OpenCL 是由 Khronos Group 针对异构计算设备（heterogeneous device）进行并行运算所设计的标准 API 以及程式语言。
OpenCL 程序分成两部分：一部分是在设备上执行的（例如 GPU），另一部分是在主机上（CPU）运行的。在设备上执行的程序就是实现“异构”和“并行计算”的部分。 为了能在设备上执行代码，程序员需要写一个特殊的程序（kernel 程序）。这个程序需要使用 OpenCL 语言编写。 OpenCL 语言采用了 C 语言的一部分加上一些约束、关键字和数据类型。在主机上运行的程序用 OpenCL 的 API 管理设备上运行的程序。主机程序的 API 用 C 语言编写，也有 C++、 Java、 Python 等高级语言接口。
设计OpenCL的目的主要有两个：
- 并行计算
- 同一份代码在不同的平台（异构）都能执行
OpenCL的根本目的就是针对并行计算提供一个统一的编程方式，从而方便对使用不同类型的并行计算硬件进行编程，这些硬件设备有多核CPU、GPU、FPGA、DSP、其他形式硬件等。如下所示：
2. Kernel程序
在OpenGL规范中，针对在GPU硬件上的编程提出了GLSL语言规范，同样的，在OpenCL规范中，针对在不同的并行计算设备上编程提出了“OpenCL C”语言规范。此语言设计参考了C语言，使用“OpenCL C”语言写的程序叫做kernel程序。通过一个简单的示例来直观的感受下如何在OpenCL程序中使用kernel程序。
void Init (void)
{
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);
// build program
clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
}
假设以上的代码截取自demo.c源文件，对这个示例代码的说明如下:
- kernelSource字符串表示的就是OpenCL kernel程序源码
- 通过OpenCL接口clCreateProgramWithSource来加载kernel源码
- kernel源码加载好之后需要进行编译，使用OpenCL接口clBuildProgram来实现这一目的
3. OpenCL驱动
OpenCL规范其实包含两个部分：
- OpenCL API规范，使用OpenCL API编写的程序叫做OpenCL程序，比如上面的示例demo.c，是在CPU上运行的
- OpenCL C规范，符合OpenCL C规范的程序叫做kernel程序，是在并行计算设备上运行的
OpenCL的驱动主要就有两大功能：
- 实现OpenCL API
- 实现OpenCL C编译器，用于将kernel程序编译成具体的并行计算设备硬件指令
如下图所示：
4. OpenCL编译器
这里主要讨论OpenCL 2.1版本之后的编译框架。OpenCL 2.1相比于之前的版本增加了两大功能：
- 制定了OpenCL C++语言规范，之前的Kernel程序都是用OpenCL C语言编写的，现在可以使用OpenCL C++语言编写
- OpenCL接口可以直接识别加载SPIR-V IR，然后编译成kernel程序在并行计算设备上执行
详细的框架如下所示：
从上图中可以看出以下几点信息：
- 用OpenCL C++语言编写的kernel程序必须编译成SPIR-V格式的IR给OpenCL接口使用
- 用OpenCL C语言编写的kernel程序可以直接传给OpenCL接口使用（主要是为了兼容考虑），也可以编译成SPIR-V格式的IR给OpenCL接口使用
- OpenCL 2.1版本的驱动能处理OpenCL C和SPIR-V格式的kernel程序
从OpenCL 2.1版本的编译框架改动来看，Khronos组织希望使用SPIR-V这一IR来统一Vulkan和OpenCL的中间表示语言。Khronos提供了一个名为SPIR-V generator/Clang的基于LLVM的编译器，用于将用OpenCL C和OpenCL C++语言编写的kernel源码编译成SPIR-V格式的二进制文件。
评论