OpenCL 认识一下

2024-04-02  本文已影响0人  buding_

先介绍一些基本概念

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

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

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

OpenCL是干啥的？

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

CUDA: 异构计算先行者~

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

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

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

OpenCL的基本认识

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在选择设备对象后才能确定运行环境;
运行之前需要对程序对象进行创建

创建程序对象.png 构建程序对象.png
6 内核对象
创建内核对象.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
__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的矢量数据

为什么有矢量数据类型？

矢量数据的使用与访问

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的使用限制

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

上一篇下一篇

猜你喜欢

热点阅读