OpenCL(Open Computing Language，开放计算语言）杂谈

Posted 2023-06-21 吴建明

OpenCL(Open Computing Language，开放计算语言）杂谈

参考文献链接

https://zhuanlan.zhihu.com/p/602844623

Linux内核库：一文细说OpenCL框架

OpenCL编程指南-1.2OpenCL基本概念_阳光开朗男孩的博客-CSDN博客

OpenCL概念基础

面向异构平台的应用都必须完成以下步骤：

1）发现构成异构系统的组件。

2）探查这些组件的特征，使软件能够适应不同硬件单元的特定特性。

3）创建将在平台上运行的指令块（内核)。

4）建立并管理计算中涉及的内存对象。

5）在系统中正确的组件上按正确的顺序执行内核。

6）收集最终结果。

这些步骤通过OpenCL中的一系列API再加上一个面向内核的编程环境来完成。我们将采用一种“分而治之”的策略解释以上步骤的所有工作。我们把问题分解为以下模型：

1）平台模型 (platform model)：异构系统的高层描述。

2）执行模型 (execution model)：指令流在异构平台上执行的抽象表示。

3）内存模型 (memory model)：OpenCL中的内存区域集合以及一个OpenCL计算期间这些内存区域如何交互。

4）编程模型( programming model)：程序员设计算法来实现一个应用时使用的高层抽象。

执行模型

OpenCL应用由两个不同部分组成：一个宿主机程序 (host program) 以及一个或多个内核 (kernel) 组成的集合。宿主机程序在宿主机上运行。OpenCL并没有定义宿主机程序如何工作的具体细节，只是定义了它与OpenCL中定义的对象如何交互。

内核在OpenCL设备上执行。它们完成OpenCL应用的具体工作。内核通常是一些简单的函数，将输入内存对象转换为输出内存对象。OpenCL定义了两类内核：

1）OpenCL内核：用OpenCL C编程语言编写并用OpenCL编译器编译的函数。所有OpenCL实现都必须支持OpenCL内核。

2）原生内核：OpenCL之外创建的函数，在OpenCL中可以通过一个函数指针来访问。例如，这些函数可以是宿主机源代码中定义的函数，或者是从一个专用库导出的函数。需要说明的是，执行原生内核的能力是OpenCL的一个可选功能，原生内核的语义依赖于具体实现。

OpenCL执行模型定义了内核如何执行。为了详细解释OpenCL执行模型，我们将分部分来讨论。首先我们会解释单个内核如何在OpenCL设备上运行。由于编写OpenCL应用的重点就是执行内核，所以这个概念是理解OpenCL的基础。接下来我们会介绍宿主机如何定义上下文来执行内核以及内核如何排队等待执行。

内核如何在OpenCL设备上执行

内核在宿主机上定义。宿主机程序发出一个命令，提交内核在一个OpenCL设备上执行。由宿主机发出这个命令时，OpenCL运行时系统会创建一个整数索引空间。对应这个索引空间中的各个点将分别执行内核的一个实例。我们将执行内核的各个实例称为一个工作项 (workitem)，工作项由它在索引空间中的坐标来标识。这些坐标就是工作项的全局ID。

提交内核执行的命令相应地会创建一个工作项集合，其中各个工作项使用内核定义的同样的指令序列。尽管指令序列是相同的，但是由于代码中的分支语句或者通过全局ID选择的数据可能不同，因此各个工作项的行为可能不同。

工作项组织为工作组 (work- group)。工作组提供了对索引空间更粗粒度的分解，跨越整个全局索引空间。换句话说，工作组在相应维度的大小相同，这个大小可以整除各维度中的全局大小。为工作组指定一个唯一的ID，这个ID与工作项使用的索引空间有相同的维度。另外为工作项指定一个局部ID，这个局部ID在工作组中是唯一的，这样就能由其全局ID或者由其局部ID和工作组ID唯一地标识一个工作项。

给定工作组中的工作项会在一个计算单元的处理单元上并发执行。这是理解OpenCL并发性的关键。具体实现可能串行化内核的执行，甚至可能在一个内核调用中串行化工作组的执行。OpenCL只能确保一个工作组中的工作项并发执行 (和共享设备上的处理器资源)。因此，不要认为工作组或内核调用会并发执行。尽管实际上它们通常确实会并发执行，但是算法设计人员不能依赖这一点。

索引空间是一个N维的值网格，因此也称为NDRange。目前，这个N维索引空间中的N可以是1、2或3。在一个OpenCL程序中，NDRange由一个长度为N的整数数组定义，N指定索引空间各维度的大小。各个工作项的全局和局部ID都是一个N维元组。在最简单的情况下,全局ID中各个值的取值范围从О开始，到该维度元素个数减1。

与为工作项指定ID类似，仍采用这种方法为工作组指定ID。有一个长度为N的数组定义各个维度中工作组的个数。工作项指定到一个工作组，并给定一个局部ID，这个局部ID中各个值的取值范围也是从0开始，到该维度中工作组个数减1。因此，通过结合工作组ID和工作组中的局部ID就可以唯一地定义一个工作项。

下面进一步分析根据这个模型建立的不同索引，并研究它们之间如何关联。可以考虑一个2维的NDRange。我们使用小写字母g表示给定下标x或y时各维度中一个工作项的全局ID。大写字母G指示索引空间各维度的大小。因此，各工作项在全局NDRange索引空间中有一个坐标 (gx，gy)，全局索引空间的大小为(Gx，Gy)，工作项坐标取值范围为[0… (Gx, -1), 0… (Gy, -1)]。

OpenCL要求各个维度中工作组的数目能够整除NDRange索引空间各个维度的大小。这样可以保证所有工作组都是满的，而且大小相同。各个方向 (在这个2维的例子中，就是x和y方向) 的工作组大小要用来为各个工作项定义一个局部索引空间。我们把一个工作组内的索引空间称为局部索引空间 (local index space)。按照前面使用大小写字母的约定，局部索引空间中各个维度 (x和y) 的大小用大写字母L表示，工作组中的局部ID使用小写字母l表示。

• OpenCL(Open Computing Language，开放计算语言）：从软件视角看，它是用于异构平台编程的框架；从规范视角看，它是异构并行计算的行业标准，由Khronos Group来维护；
• 异构平台包括了CPU、GPU、FPGA、DSP，以及最近几年流行的各类AI加速器等；
• OpenCL包含两部分：

    1）用于编写运行在OpenCL device上的kernels的语言（基于C99）；    2）OpenCL API，至于Runtime的实现交由各个厂家，比如Intel发布的opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz

 以人工智能场景为例来理解一下，假如在某个AI芯片上跑人脸识别应用，CPU擅长控制，AI processor擅长计算，软件的flow就可以进行拆分，用CPU来负责控制视频流输入输出前后处理，AI processor来完成深度学习模型运算完成识别，这就是一个典型的异构处理场景，如果该AI芯片的SDK支持OpenCL，那么上层的软件就可以基于OpenCL进行开发了。

话不多说，看看OpenCL的架构吧。

2. OpenCL架构

OpenCL架构，可以从平台模型、内存模型、执行模型、编程模型四个角度来展开。

2.1 Platform Model

平台模型：硬件拓扑关系的抽象描述

 平台模型由一个Host连接一个或多个OpenCL Devices组成；

• OpenCL Device，可以划分成一个或多个计算单元Compute Unit（CU）；
• CU可以进一步划分成一个或多个处理单元Processing Unit（PE），最终的计算由PE来完成；
• OpenCL应用程序分成两部分：host代码和device kernel代码，其中Host运行host代码，并将kernel代码以命令的方式提交到OpenCL devices，由OpenCL device来运行kernel代码；

2.2 Execution Model

执行模型：Host如何利用OpenCL Device的计算资源完成高效的计算处理过程

Context

OpenCL的Execution Model由两个不同的执行单元定义：1）运行在OpenCL设备上的kernel；2）运行在Host上的Host program；其中，OpenCL使用Context代表kernel的执行环境：

 Context包含以下资源：

• Devices：一个或多个OpenCL设备；
• Kernel Objects：OpenCL Device的执行函数及相关的参数值，通常定义在cl文件中；
• Program Objects：实现kernel的源代码和可执行程序，每个program可以包含多个kernel；
• Memory Objects：Host和OpenCL设备可见的变量，kernel执行时对其进行操作；

NDrange

 kernel是Execution Model的核心，放置在设备上执行，当kernel执行前，需要创建一个索引空间NDRange（一维/二维/三维）；

• 执行kernel实例的称为work-item，work-item组织成work-group，work-group组织成NDRange，最终将NDRange映射到OpenCL Device的计算单元上；

有两种方式来找到work-item：

1. 通过work-item的全局索引；
2. 先查找到所在work-group的索引号，再根据局部索引号确定；

以一维为例：

 上图中总共有四个work-group，每个work-group包含四个work-item，所以local_size的大小为4，而local_id都是从0开始重新计数；

• global_size代表总体的大小，也就是16个work-item，而global_id则是从0开始计数；

以二维为例：

 二维的计算方式与一维类似，也是结合global和local的size，可以得出global_id和local_id的大小，细节不表了；

三维的方式也类似，略去。

2.3 Memory Model

内存模型：Host和OpenCL Device怎么来看待数据

 OpenCL的内存模型中，包含以下几类类型的内存：

• Host memory：Host端的内存，只能由Host直接访问；
• Global Memory：设备内存，可以由Host和OpenCL Device访问，允许Host的读写操作，也允许OpenCL Device中PE读写，Host负责该内存中Buffer的分配和释放；
• Constant Global Memory：设备内存，允许Host进行读写操作，而设备只能进行读操作，用于传输常量数据；
• Local Memory：单个CU中的本地内存，Host看不到该区域并无法对其操作，该区域允许内部的PE进行读写操作，也可以用于PE之间的共享，需要注意同步和并发问题；
• Private Memory：PE的私有内存，Host与PE之间都无法看到该区域；

2.4 Programming Model

 在编程模型中，有两部分代码需要编写：一部分是Host端，一部分是OpenCL Device端；

• 编程过程中，核心是要维护一个Context，代表了整个Kernel执行的环境；
• 从cl源代码中创建Program对象并编译，在运行时创建Kernel对象以及内存对象，设置好相关的参数和输入之后，就可以将Kernel送入到队列中执行，也就是Launch kernel的流程；
• 最终等待运算结束，获取计算结果即可；

3. 编程流程

 上图为一个OpenCL应用开发涉及的基本过程；

下边来一个实际的代码测试跑跑，Talk is cheap, show me the code!

4. 示例代码

• 测试环境：Ubuntu16.04，安装Intel CPU OpenCL SDK（opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz）；
• 为了简化流程，示例代码都不做容错处理，仅保留关键的操作；
• 整个代码的功能是完成向量的加法操作；
• 在Host程序中，创建program对象时会去读取kernel的源代码，本示例源代码位于：vector_add.cl文件中

4.1 Host端程序

#include <iostream>

#include <fstream>

#include <sstream>



#include <CL/cl.h>



const int DATA_SIZE = 10;



int main(void)



    /* 1. get platform & device information */

    cl_uint num_platforms;

    cl_platform_id first_platform_id;

    clGetPlatformIDs(1, &first_platform_id, &num_platforms);





    /* 2. create context */

    cl_int err_num;

    cl_context context = nullptr;

    cl_context_properties context_prop[] = 

        CL_CONTEXT_PLATFORM,

        (cl_context_properties)first_platform_id,

        0

    ;

    context = clCreateContextFromType(context_prop, CL_DEVICE_TYPE_CPU, nullptr, nullptr, &err_num);





    /* 3. create command queue */

    cl_command_queue command_queue;

    cl_device_id *devices;

    size_t device_buffer_size = -1;



    clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, nullptr, &device_buffer_size);

    devices = new cl_device_id[device_buffer_size / sizeof(cl_device_id)];

    clGetContextInfo(context, CL_CONTEXT_DEVICES, device_buffer_size, devices, nullptr);

    command_queue = clCreateCommandQueueWithProperties(context, devices[0], nullptr, nullptr);

    delete [] devices;





    /* 4. create program */

    std::ifstream kernel_file("vector_add.cl", std::ios::in);

    std::ostringstream oss;



    oss << kernel_file.rdbuf();

    std::string srcStdStr = oss.str();

    const char *srcStr = srcStdStr.c_str();

    cl_program program;

    program = clCreateProgramWithSource(context, 1, (const char **)&srcStr, nullptr, nullptr);





    /* 5. build program */

    clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr);





    /* 6. create kernel */

    cl_kernel kernel;

    kernel = clCreateKernel(program, "vector_add", nullptr);





    /* 7. set input data && create memory object */

    float output[DATA_SIZE];

    float input_x[DATA_SIZE];

    float input_y[DATA_SIZE];

    for (int i = 0; i < DATA_SIZE; i++) 

        input_x[i] = (float)i;

        input_y[i] = (float)(2 * i);

    



    cl_mem mem_object_x;

    cl_mem mem_object_y;

    cl_mem mem_object_output;

    mem_object_x = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_x, nullptr);

    mem_object_y = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_y, nullptr);

    mem_object_output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * DATA_SIZE, nullptr, nullptr);





    /* 8. set kernel argument */

    clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem_object_x);

    clSetKernelArg(kernel, 1, sizeof(cl_mem), &mem_object_y);

    clSetKernelArg(kernel, 2, sizeof(cl_mem), &mem_object_output);





    /* 9. send kernel to execute */

    size_t globalWorkSize[1] = DATA_SIZE;

    size_t localWorkSize[1] = 1;

    clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);





    /* 10. read data from output */

    clEnqueueReadBuffer(command_queue, mem_object_output, CL_TRUE, 0, DATA_SIZE * sizeof(float), output, 0, nullptr, nullptr);

    for (int i = 0; i < DATA_SIZE; i++) 

        std::cout << output[i] << " ";

    

    std::cout << std::endl;





    /* 11. clean up */

    clRetainMemObject(mem_object_x);

    clRetainMemObject(mem_object_y);

    clRetainMemObject(mem_object_output);

    clReleaseCommandQueue(command_queue);

    clReleaseKernel(kernel);

    clReleaseProgram(program);

    clReleaseContext(context);



    return 0;

4.2 OpenCL Kernel函数

内容如下：

__kernel void vector_add(__global const float *input_x,

 __global const float *input_y,

 __global float *output)



 int gid = get_global_id(0);

 

 output[gid] = input_x[gid] + input_y[gid];

4.3 输出

 参考文献

The OpenCL Specification

https://zhuanlan.zhihu.com/p/602844623

Linux内核库：一文细说OpenCL框架

OpenCL编程指南-1.2OpenCL基本概念_阳光开朗男孩的博客-CSDN博客

人工智能芯片与自动驾驶