OpenCL（壹）：了解OpenCL模型到编写第一个CL内核程序

1.前言

2.简单了解OpenCL

3.为什么要使用OpenCL

4.OpenCL架构

5.OpenCL中的平台模型(Platform Model)

6.OpenCL中的内存模型(Execution Model)

7.OpenCL中的执行模型(Memory Model)

8.OpenCL中的编程模型(Programmin Model)

9.OpenCL中的同步机制

10.编写第一个OpenCL程序

前言

大多数由程序员编写的代码都是执行在CPU上的，而现代计算机不只提供了CPU还提供了GPU，但是大多数场景往往都不会涉及到使用GPU进行运算，有个别业务如ffmpeg编解码，OpenCV中UMat，VTK中的模型渲染等等则会使用到GPU，但是这些都是第三方编写好的函数进行调用。而本系列文章将会对如何把算法迁移到GPU上执行，迁移的主要理由则是作者不满足于压榨语言特性以及编写更高效的算法来使程序获得更高效的执行效率

简单了解OpenCL

OpenCL全称Open Computing Language即开放计算语言。OpenCL为异构平台提供了一个编写程序，尤其是并行程序的开放的框架标准。OpenCL所支持的异构平台可由多核CPU、GPU或其他类型的处理器组成。OpenCL由两部分组成，一是用于编写内核程序(在OpenCL设备上运行的代码) 的语言，二是定义并控制平台的API。OpenCL提供了基于任务和基于数据两种并行计算机制，它极大地扩展了GPU 的应用范围，使之不再局限于图形领域。

为什么要使用OpenCL

高性能CPU的由于很难克服提高时钟频率后的散热问题转而使用增加运算核心的方法加速。作为图形渲染专用的处理器， GPU具有高度的并行特性。以下几点是GPU的特点，也是使用GPU运算的理由：

1.GPU的运算核心数量要远远超过高端 CPU的核心数量。GPU的每个运算核心并没有 CPU的运算核心工作频率高，但是其总体性能-芯片面积比和性能-功耗比都很高，在处理并行计算的相关任务中有很大优势。

2.GPU是通过大量并行线程之间交织运行隐藏全局访问的延迟，同时GPU还拥有大量的寄存器、局部存储器及Cache等来提升外部存储的访问性能。

而且基于 GPU或者其他并行运算设备的算法与传统的基于CPU的串行算法有很大差别：

1.并行算法中要有大量的线程在运行，而一般的串行算法都只有一个线程在运行。

2.并行算法中的每个线程的行为需要尽量保持一致，如果分支很多，各线程又选择不同路径执行，会严重降低 GPU运算的效率。在CPU中，即使有两个线程的行为高度不一致，也不会非常影响性能。

3.在程序不加特殊约束的情况下，并行运算设备是不保证每个线程看到的全局内存是一致的。程序员有责任维护线程同步以及内存管理等任务。

4.在传统的串行运算设备中，例如 CPU，线程之间切换的开销是比较大的。所以一般来说，是不鼓励程序员为一个算法开启大量线程的。而在类似于 GPU的并行运算设备中，线程之间的切换是非常廉价的。这些设备也正是通过线程之间切换来隐藏一些内存访问延迟的。与CPU相反，并行设备一般是不鼓励设备运行具有很少线程的算法。

OpenCL架构

在OpenCL中划分为以下四个模型，这四个模型共同构成了OpenCL框架的核心，使得开发人员能够编写出高效、可移植的并行计算程序：

1.平台模型（Platform Model）

2.内存模型（Execution Model）

3.执行模型（Memory Model）

4.编程模型（Programmin Model）

OpenCL中的平台模型(Platform Model)

平台模型主要是由一个主机 (host)连接一个或多个 OpenCL设备构成。其中每个OpenCL设备又可以分割成一个或多个计算单元 (CU)，每个计算单元又可以进一步分割成一个或多个处理单元 (PE)，各种计算操作都是在处理单元中完成的。所有由 OpenCL编写的应用程序都是从 host启动并在 host上结束的，host端管理着整个平台上的所有计算资源。应用程序会从host端向OpenCL设备的处理单元发送计算命令。

       |-----OpenCL设备
host----     
       |-----OpenCL设备--------计算单元CU
                          |
                          |----计算单元CU-----------处理单元PE
                                              |
                                              |----处理单元PE

图1.OpenCL中的平台模型架构

OpenCL中的内存模型(Execution Model)

内存模型主要分为以下四种：

1.全局内存：对所有工作组可见，主要存储数据和回传至host的数据

2.局部内存：只对同一个工作组内的工作节点可见，映射到片上物理内存，具有较短的访问延迟和较高的传输带宽

3.私有内存：只能由工作节点自己访问，局部变量和非指针内核传输通常都在私有内存是开辟存储空间

4.常量内存：存储只读数据，所有工作节点都可以访问，一个OpenCL设备对应一个全局内存，而在全局内存中又划分了一部分地址作为常量内存

       |-----全局内存
host----     
       |-----全局内存--------工作组
                       |
                       |----工作组-----------工作节点
                                       |
                                       |----工作节点

		全局内存	常量内存	本地内存	私有内存
Host	分配	动态	动态	动态	不可分配
Host	访问	可读写	可读写	不可访问	不可访问
Kerenl	分配	不可访问	静态	静态	静态
Kerenl	访问	可读写	只读	可读写	可读写

表1.host和kernel是如何管理和使用各类内存

而由于OpenCL设备和host之间的内存是互相独立的，所有host要与OpenCL设备进行数据的交互通常通过以下两种方式（这两种方式都分为non-block模式和block模式）：

1.拷贝数据法：host通过相应的OpenCL的API接口将数据从host写入到OpenCL设备的内存中或者从OpenCL设备内存读出数据到host内存中

2.内存映射法：通过相应OpenCL的API接口将OpenCL的内存对象映射到 host端可见的内存地址空间中。映射之后用户就可以在host端的映射地址读写该内存了，在读写完成之后用户必须使用对应 API解除这种映射关系

图2.OpenCL中的内存模型架构

OpenCL中的执行模型(Memory Model)

OpenCL中的执行模型主要用于管理kernel在OpenCL设备上的运行，它分为以下两种：

1.host上执行的主程序

2.OpenCL设备上执行的内核程序Kernel

对于Kernel程序的工作空间又划分为工作组空间，在Host创建Kernel程序时，必须先为该Kernel程序创建工作空间，该空间可以是一维（线），也可以是二维（矩阵）和三维（立体）。所有工作节点的执行同一个Kernel程序，而每一个节点在相应维度上的索引被定义为该节点在该维度上的全局ID。工作组的空间必须与Kernel程序的工作空间维度相同，整个工作空间可以划分为多个工作组空间，每一个工作组都有对应的工作组空间（这里其实就相当于内存模型中的全局内存内含多个工作组（局部内存），局部内存中又有多个工作节点（私有内存））

对应的上段讲到的全局ID是工作节点在Kernel工作空间相应维度的索引（简单理解为相对于全局内存的索引），局部ID则是工作节点对于该工作组空间的索引（相对于局部内存的索引）

图3.OpenCL中的执行模型架构

OpenCL中的编程模型(Programmin Model)

OpenCL中的编程模型主要分为以下两种：

1.数据并行模型：运行相同的内核程序，根据全局ID或者局部ID映射的内存导致处理的数据不同，而数据并行模型又分为以下两种：

1.显式分级模型：指定参与并行计算的工作节点的数目和工作节点所属的工作组

2.隐式分级模型：工作节点和工作组由OpenCL分配管理

2.任务并行模型：每个工作节点在执行内核程序时是独立的，但是工作节点间的数据可以通过全局内存或局部内存进行数据交互

为了方便理解后续的内容，此次对Kernel程序和Native Kernel程序进行补充：

1.Kernel程序指的是内核程序，运行在OpenCL设备中的程序

2.Native Kernel程序指的是运行在Host上的程序

在OpenCL中实习任务并行模型的方法主要是通过使用OpenCL设备支持的向量类型数据结构，同时执行或者选择性执行多个Kernel程序和执行Kernel程序的同时交叉执行Native Kernel程序

OpenCL中的同步机制

OpenCL中的同步主要分为两个概念，一个是同步领域，一个是同步方法。

同步领域：

1.工作组内同步：实现同一个工作组中所有工作节点之间的同步

2.命令队列间同步：实现同一个上下文中不同的命令队列之间和同一个命令队列的不同命令之间的同步

同步方法：

1.工作组内通过阻断函数实现

2.命令队列内同步通过OpenCL提供的命令队列的阻断函数

3.不同命令队列间同步通过每一个命令关联的事件实现同步

编写第一个OpenCL程序

在编写第一个OpenCL程序前我们需要先了解OpenCL程序的基本流程：

图4.OpenCL程序基本流程

#include <CL/cl.h>
#include <iostream>

#define KERNEL(...)#__VA_ARGS__		//使用#__VA_ARGS_宏将传入KERNEL宏的实参转为字符串

//编写内核函数，该内核函数将转换为字符串并由指针kernelSourseCode指向
const char* kernelSourceCode = KERNEL(
	__kernel void wildpointer(__global uint * buffer) {
		size_t gidx = get_global_id(0);
		size_t gidy = get_global_id(1);
		size_t lidx = get_local_id(0);
		buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
	}
);

int main() {
	std::cout << "WildPointer：Hello OpenCL" << std::endl;
	//获取可用的设备平台ID
	cl_int status = 0;
	size_t deviceListSize;
	cl_uint numPlatforms;
	cl_platform_id platform = NULL;
	status = clGetPlatformIDs(0, NULL, &numPlatforms);
	if (status != CL_SUCCESS) {
		std::cout << "错误：获取设备失败" << std::endl;
		return EXIT_FAILURE;
	}
	if (numPlatforms > 0) {
		cl_platform_id* platforms = (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));
		status = clGetPlatformIDs(numPlatforms, platforms, NULL);
		if (status != CL_SUCCESS) {
			std::cout << "错误：获取平台ID失败" << std::endl;
			return -1;
		}
		for (unsigned int i = 0; i < numPlatforms; ++i) {
			char pbuff[100];
			status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(pbuff), pbuff, NULL);
			platform = platforms[i];
			if (!strcmp(pbuff, "GPU")) {
				break;
			}
		}
		delete platforms;
	}

	//找到可用的设备平台
	cl_context_properties cps[3] = {CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 };
	cl_context_properties* cprops = (NULL == platform) ? NULL : cps;

	// 使用正确的设备类型创建上下文
	cl_context context = clCreateContextFromType(cprops, CL_DEVICE_TYPE_GPU, NULL, NULL, &status);
	if (status != CL_SUCCESS) {
		std::cout << "错误：为GPU生成上下文失败" << std::endl;
		return EXIT_FAILURE;
	}

	//寻找OpenCL设备(GPU)
	//获取设备个数
	status = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceListSize);
	if (status != CL_SUCCESS) {
		std::cout << "错误：寻找可用的设备个数失败" << std::endl;
		return EXIT_FAILURE;
	}
	cl_device_id* devices = (cl_device_id*)malloc(deviceListSize);
	if (devices == 0) {
		std::cout << "错误：当前可用设备数为0" << std::endl;
		return EXIT_FAILURE;
	}

	//获取设备列表
	status = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceListSize, devices, NULL);
	if (status != CL_SUCCESS) {
		std::cout << "错误：获取设备信息失败" << std::endl;
		return EXIT_FAILURE;
	}

	//装载内核程序，编译CL程序，生成OpenCL内核实例
	size_t sourceSize[] = { strlen(kernelSourceCode) };
	cl_program program = clCreateProgramWithSource(context, 1, &kernelSourceCode, sourceSize, &status);
	if (status != CL_SUCCESS) {
		std::cout << "错误：将二进制文件装载到内核程序失败" << std::endl;
		return EXIT_FAILURE;
	}

	//为指定的OpenCL设备生成CL程序
	status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
	if (status != CL_SUCCESS) {
		std::cout << "错误：编译CL程序失败" << std::endl;
		return EXIT_FAILURE;
	}

	//获取内核实例的句柄
	cl_kernel kernel = clCreateKernel(program, "wildpointer", &status);
	if (status != CL_SUCCESS) {
		std::cout << "错误：在程序上初始化内核程序失败" << std::endl;
		return EXIT_FAILURE;
	}

	//创建OpenCL命令队列
	cl_command_queue commandQueue = clCreateCommandQueue(context, devices[0], 0, &status);
	if (status != CL_SUCCESS) {
		std::cout << "错误：初始化命令队列失败" << std::endl;
		return EXIT_FAILURE;
	}

	//创建OpenCL缓冲区
	unsigned int* outbuffer = new unsigned int[4 * 4];
	memset(outbuffer, 0, 4 * 4 * 4);
	cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR, 4 * 4 * 4, NULL, &status);
	if (status != CL_SUCCESS) {
		std::cout << "错误：初始化OpenCL缓冲区失败" << std::endl;
		return EXIT_FAILURE;
	}

	//将参数传入内核程序
	status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&outputBuffer);
	if (status != CL_SUCCESS) {
		std::cout << "错误：设置内核参数失败" << std::endl;
		return EXIT_FAILURE;
	}

	//将内核程序插入命令队列
	size_t globalThreads[] = { 4, 4 };
	size_t localThread[] = { 2, 2 };
	status = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, globalThreads, localThread, 0, NULL, NULL);
	if (status != CL_SUCCESS) {
		std::cout << "错误：内核程序插入命令队列失败" << std::endl;
		return EXIT_FAILURE;
	}

	status = clFinish(commandQueue);
	if (status != CL_SUCCESS) {
		std::cout << "错误：完全命令队列" << std::endl;
		return EXIT_FAILURE;
	}

	status = clEnqueueReadBuffer(commandQueue, outputBuffer, CL_TRUE, 0, 4 * 4 * 4, outbuffer, 0, NULL, NULL);
	if (status != CL_SUCCESS) {
		std::cout << "错误：读取缓冲区队列失败" << std::endl;
		return EXIT_FAILURE;
	}

	std::cout << "内核程序输出：" << std::endl;
	for (int i = 0; i < 16; ++i) {
		std::cout << outbuffer[i] << " ";
		if ((i + 1) % 4 == 0) {
			std::cout << std::endl;
		}
	}

	status = clReleaseKernel(kernel);
	status = clReleaseProgram(program);
	status = clReleaseMemObject(outputBuffer);
	status = clReleaseCommandQueue(commandQueue);
	status = clReleaseContext(context);
	free(devices);
	delete outbuffer;
	return 0;
}

PS：针对代码调用的函数会在下一章讲解，由于内容过长就不放到同一篇文章中了，具体如果了解了OpenCL中的四大模型，并且熟悉C/C++语言则对于阅读上述程序不会太吃力