OpenCL 认识一下

作者: buding_ | 来源:发表于2024-04-02 15:29 被阅读0次

先介绍一些基本概念

主处理器（CPU）：负责主要流程控制和常规计算
协处理器（GPU、FPGA、ASIC等芯片）：负责并行计算，速度快、低延时

FPGA（Field-Programmable Gate Array）是一种可重构的硬件平台，它具有高度的灵活性和可编程性，使其在多个领域中都有广泛的应用。以下是FPGA擅长的一些领域：

通信系统。FPGA用于无线和有线通信系统的信号处理、网络处理以及调制解调器，能够灵活地适应不断变化的通信协议需求。
数字信号处理。特别是在处理重复性任务如FIR滤波时，FPGA的并行处理机制可以高效地完成这些任务。
视频图像处理。FPGA能够处理高清晰度图像数据，实现视频压缩编码，同时处理多个视频流，保证图像输出的实时性和质量。
高速接口设计。由于其高速处理能力和多达成百上千的IO，FPGA在实现高速接口设计方面具有独特优势，减少了对多个接口芯片的需求。
人工智能。FPGA可以用于深度学习、大数据分析和数据库处理等人工智能应用中，尤其是在需要快速执行算法的场景中。
嵌入式系统。在航空航天、医疗器械等领域，FPGA能够处理高速度、高精度的信号数据，实现实时控制系统和硬件调试。

ASIC（Application-Specific Integrated Circuit，特定应用集成电路）是一种定制的、专门为特定应用设计的集成电路。与通用处理器（如CPU和GPU）不同，ASIC 被定制为执行特定任务或应用，通常以高度优化的方式执行特定的计算操作。ASIC 的设计旨在提供最佳性能和效率，因为它专门用于执行一种或几种紧密相关的任务。
在与CPU、GPU和NPU的比较中，ASIC 有一些独特的特点：

定制性： ASIC 的设计是为了满足特定的应用需求，因此它可以在硬件级别上进行高度的定制，以实现最佳性能。
性能和效率：由于 ASIC 专注于特定应用，它通常能够提供比通用处理器更高的性能和效率。这使得 ASIC 在特定领域的计算任务中表现出色。
电力效率： ASIC 的设计优化通常使其在特定任务上更为节能，因为它仅执行特定的计算操作，避免了通用处理器可能面临的冗余计算。
应用范围： ASIC 可以用于各种应用，包括密码学、信号处理、网络通信和各种专门的计算任务。

OpenCL是干啥的？

当你遇到针对并行度不高、带宽不够、延时较高等问题时，可以采取异构计算来进行优化；
常用异构计算组合: CPU+GPU、 CPU+ FPGA

CUDA: 异构计算先行者~

计算统一设备架构
以c/c++为基础
使用广泛的并行计算平台
只是针对Nvidia产品（平台绑定）
闭源，未被其他厂商接受；

OpenCL（Open Computing Language），也是一套并行计算框架！

覆盖CPU、GPU、FPGA以及其他多种处理器芯片；
支持不同平台，支持在不同平台移植，通用性强；

OpenCL的发展：
-> 2008年由由Apple设计、提出, 同年发布OpenCL 1.0
-> 2010年OpenCL 1.1技术规范发布
-> 2011年OpenCL 1.2技术规范发布（增加与OpenGL的交互性，同步事件，设备划分）
-> 2013年OpenCL 2.0技术规范发布（共享虚拟内存、内核嵌套并行、通用地址空间）

OpenCL的应用场景：
分子动力学模拟、计算流体力学、图形图像处理、视频音频处理、计算机视觉、人工智能、智能硬件、调度多核的ARM处理器、云计算

OpenCL的基本认识

选择OpenCL平台是OpenCL程序的第一步
OpenCL平台指的是OpenCL设备和OpenCL框架的组合
一套异构系统上可以同时存在多个OpenCL平台
需要显式的选择使用的平台, #include "CL/cl.h"

1. 获取系统中的OpenCL平台

获取OpenCL平台.png

获取OpenCL平台code.png

2 获取OpenCL设备

常见的OpenCL设备有 GPU、FPGA、CPU(作为特殊OpenCL设备处理)等；

选择OpenCL设备.png

1.获取设备数量
2.根据数量去申请一块内存
3.把设备的对象结构提存放到申请的内存中

[图片上传中...(OpenCL设备Type.png-8e7a3b-1712028461884-0)]

OpenCL设备Type.png

3 OpenCL上下文

关联设备
关联内存对象
关联命令队列
管理程序对象
管理内核对象
驱动应用程序与设备通信

上下文可以指定一个或多个设备作为当前的操作对象，上下文context用来管理command-queue，memory，program和kernel，以及指定kernel在上下文中的一个或多个设备上运行

上下文指定设备的创建，默认使用所有的设备.png

上下文根据类型选择设备的创建，选择第一个找到的设备.png

4 命令队列

沟通：1-命令是主机发给设备的消息 2-通知设备执行操作
内情：1-主机与设备的数据搬运 2-设备内/设备间的数据搬运和内核执行
流程：
1-命令提交到命令队列
2-命令队列把需要执行的命令发送给设备
3-设备接收到的命令
4-命令只能是从主机发送给设备
5-每个命令队列只能管理一个设备

OpenCL命令队列的创建.png

命令队列中的参数.png

5 程序对象

包含内核函数的集合
为关联设备编译内核
可以由OpenCL C源代码文本创建
可以使用程序二进制代码创建

内核对象就是在设备上执行的函数
程序对象就是内核对象的一个容器
内核对象由程序对象创建和管理
一个程序对象可以包含多个内核对象

OpenCL在选择设备对象后才能确定运行环境;
运行之前需要对程序对象进行创建

创建程序对象.png

构建程序对象.png

6 内核对象

本质是一段代码（函数）
需要编译为可执行文件（aocx）
在OpenCL设备上执行

创建内核对象.png

7 内存对象

内核参数如何设置？
主机端如何向内核传递数据？
主机端如何获取内核的执行结果？

而内存对象正是解决以上的3个问题的关键；

内存对象的创建.png

创建内存对象的参数.png

8 设置内核参数

设置内核参数.png

9 执行内核

执行内核.png

上面函数适用于无依赖的并行任务，而往往我们的场景都会存在数据依赖，那么可以采用下面的：

执行内核2.png

10 读取结果

读取结果.png

11 资源回收

1 内核对象依赖程序对象
2 程序对象依赖设备和上下文
3 命令队列依赖设备和上下文
4 内存对象依赖于上下文
5 上下文依赖于设备和平台

资源回收.png

ps：

错误码s.png

OpenCL C的语法认识

1 OpenCL C

专门用于编写OpenCL内核程序
类似于一个动态库，内核类似于动态库的一个导出函数
基于C99 C语言规范，在C99规范上进行扩充

__kernel void adder(__global float *a, __global float *b, __global float *result)
{
  int tid = get_global_id(0);
  result[tid] = a[tid] + b[tid];
}
/*
兼容C99的关键字
地址空间限定符： __global, global, __local, local, __constant, constant, __private, private
__constant init_value = 1;
函数限定符：__kernel, kernel
访问限定符：__read_only, read_only, __write_only, write_only, __read_write, read_write 只能修复函数的参数

*/

OpenCL C的数据类型.png

类型表示在OpenCL C语言中使用， API类型则是在主机端使用

2 OpenCL C的矢量数据

为什么有矢量数据类型？

使用矢量数据是提升性能的重要方式
现代处理器一条指令可以处理多条数据

矢量数据的使用与访问

矢量数据实际上是由多个标量数据组成的
可以采用<vector>.xyzw进行访问，分别访问0-3个数据，最多4个
同样可以采用索引方式, <vector>.s0, <vector>.s1 .... <vector>.sf, 最多到f（即16进制的15）
hi, 访问矢量数据的高半部分(N/2 , N/2+1 ... N)
lo, 访问矢量数据的低半部分(0, 1 ... N/2-1)
支持直接通过索引方式进行复制和修改
支持交换和复制
支持通用的四则运算、支持三目运算

float8 v8 = (float8)(1,2,3,4,5,6,7,8);
float a = v8.w; //4
float b = v8.s6 //7
float4 v4_high = v8.hi; // 1 2 3 4
float4 v4_low = v8.lo; // 5 6 7 8 

float4 v4_sum = v4_high + v4_low; // 1+5 2+6 3+7 +4+8
float4 v4_mul_2 = v4_high * 2;  // 1*2 2*2 3*2 4*2
float res = v4_high > v4_low ? v4_high : v4_low;
/*
res.x = v4_high.x > v4_low.x ? v4_high.x : v4_low.x ;
res.y = v4_high.y > v4_low.y ? v4_high.y : v4_low.y ;
res.z = v4_high.z > v4_low.z ? v4_high.z : v4_low.z ;
res.w = v4_high.w > v4_low.w ? v4_high.w : v4_low.w ;

*/

3 简短示例

__kernel void add(__global float * restrict a, __global float * restrict b, __global float *restrict c, __global int N) {
  int i = 0;
  for (i = 0; i < N; i++)  {
    c[i] = a[i] + b[i];
  }
}
/*
C语言中的restrict关键字是C99标准中引入的一个关键字，用于指明指针的独占访问权限。
它的作用在于告诉编译器，一个指针指向的内存区域没有别的指针可以访问。

在C语言中，restrict关键字可以应用于指针类型的参数、函数返回值和局部变量。
它的使用可以提供一些性能优化，因为编译器可以基于restrict关键字作出一些假设，从而进行更好的优化。
*/

OpenCL C的使用限制

内核函数必须以 __kernel或者kernel 作为前置修饰符
内核参数不能是指向指针的指针（即不能使用二维指针或多维指针作为参数，但是在内核函数的内部是可以使用，但不建议，因为会影响OpenCL设备的搜索效果）
所有内核函数必须返回void
结果通常以指针方式返回主机程序

OpenCL Host端开发流程

1、导入OpenCL库

2、【平台模型】操作

a)获取系统中的OpenCL平台，确定使用的平台

cl_int
clGetPlatformIDs(
         cl_uint num_entries,
         cl_platform_id *platforms,
         cl_uint *num_platforms)

b)获取OpenCL设备，确定使用的设备【平台模型】
➢常见的OpenCL设备有 GPU、FPGA、CPU(作为特殊OpenCL设备处理)等；

cl_int
clGetDeviceIDs(
             cl_platform_id platform,
  cl_device_type device_type,
  cl_uint num_entries,
  cl_device_id *devices,
  cl_uint *num_devices)

3.【执行模型】操作

a) 创建OpenCL上下文
➢上下文可以指定一个或多个设备作为当前的操作对象，上下文(context)用来持续跟踪管理命令队列(command-queue)，内存对象(memory)，程序对象(program)和内核对象(kernel)，以及指定kernel在上下文中的一个或多个设备上运行

cl_context
clCreateContext(
  const cl_context_properties *properties,
  cl_uint num_devices,
  const cl_device_id *devices,
  void (CL_CALL_BACK *pfn_notify)(
    const char *errinfo,
    const void *private_info,
    size_t cb,
    void *user_data),
    void *user_data,
    cl_int *errcode_ret)

b) 创建命令队列（command-queue)
➢作为一种通信机制，可以让host发请求到对应的device
➢主机与设备的数据搬运
➢命令队列需要在每个设备上都进行创建，并且命令队列要在上下文的基础上进行创建

cl_command_queue
clCreateCommandQueueWithProperties(
  cl_context context,
  cl_device_id device,
  cl_command_queue_properties peoperties,
  cl_int *errcode_ret)

➢任何以clEnqueue开头的OpenCL API都能向命令队列提交一个命令，并且这些API都需要一个命令队列对象作为输入参数。例如，clEnqueueReadBuffer()将device上的数据传递到host，clEnqueueNDRangeKernel()申请一个内核在对应device执行

c) 事件
➢OpenCL API中，用来指定命令之间依赖关系的对象称为事件(event)
➢通过事件对命令执行的状态随时进行查询 —— clGetEventInfo()

4.【编程模型】操作

a) 创建程序对象
➢包含内核函数的集合
➢为关联设备编译内核
➢可以由OpenCL C源代码文本创建
➢可以使用程序二进制代码创建
b) 创建内核对象（kernel）
➢本质是一段代码（函数）
➢需要编译为可执行文件（aocx）
➢在OpenCL设备上执行

cl_kerenl
clCreateKernel(
  cl_program program,
  const char *kernel_name,
  cl_int *errcode_ret)

c) 参数设置
d) 执行内核

5.【内存模型】操作

OpenCL定义了三种内存类型：数组、图像和管道
a) 数组缓存
➢Buffer类型类似于C语言中的数据(使用malloc函数开辟)，这种类型中数据在内存上是连续的

cl_mem
clCreateBuffer(
  cl_context context,
  cl_mem_flags flags,
  size_t size,
  void *host_ptr,
  cl_int *errcode_ret)

b) 图形对象

➢图像也是OpenCL内存对象，其抽象了物理数据的存储，以便“设备指定”的访存优化。与数组不同，图像数组数据不能直接访问。因为相邻的数据并不保证在内存上连续存储。使用图像的目的就是为了发挥硬件空间局部性的优势，并且可以利用设备硬件加速的能力

cl_mem
clCreateImage(
  cl_context context,
  cl_mem_flags flags,
  const cl_image_format *image_format,
  const cl_image_desc *image_desc,
  void *host_ptr,
  cl_int *errcode_ret)

c) 管道对象
➢管道内存对象就是一个数据元素(被称为packets)队列，其和其他队列一样，遵循FIFO(先进先出)的方式。一个管道对象具有一个写入末尾点，用于表示元素由这里插入；并且，有一个读取末尾点，用于表示元素由这里移除。

d) 数据转移命令
➢在内存对象被内核调用之前，为其写入相应的数据到设备端；以及，在内存对象最后一次使用之后，读取其数据到主机端
clEnqueueWriteBuffer()和clEnqueueReadBuffer()

e) 内存区域
➢全局内存对于执行内核中的每个工作项都是可见的(类似于CPU上的内存)。当数据从主机端传输到设备端，数据就存储在全局内存中。有数据需要从设备端传回到主机端，那么对应的数据需要存储在全局内存中。其关键字为global或__global
➢常量内存并非为只读数据设计，但其能让所有工作项同时对该数据进行访问。这里存储的值通常不会变化，使用关键字constant或__constant
➢局部内存中的数据，只有在同一工作组内的工作项可以共享。通常情况下，局部内存会映射到片上的物理内存，局部内存具有更短的访问延迟，以及更高的传输带宽。使用local或__local关键字。
➢私有内存只能由工作项自己进行访问。局部变量和非指针内核参数通常都在私有内存上开辟。

参考：
https://www.khronos.org/opencl/
https://www.bookstack.cn/read/Heterogeneous-Computing-with-OpenCL-2.0/content-chapter3-3.1-chinese.md