OpenCL 嵌入式端算法加速
1 嵌入式端算法加速架构
1.1 OpenMP
1.1.1 平台是否支持
简单理解，我们使用 OpenMP 就可以非常简单地实现多线程有多简单呢，查看是否支持 OpenMP在使用 OpenMP 完成多线程任务时，首先得查看当前编译器是否支持 OpenMP，我在 Linux 上的 GCC 编译器是默认支持 OpenMP 的，只需在生成可执行文件的命令中加入 -fopenmp 即可下面的代码也可以查看当前编译器是否支持 OpenMP。
check_openmp.c
#include
int main()
{
#if _OPENMP
printf(“support openmp\n”);
#else
printf(“not support openmp\n”);
#endifreturn 0;
}
1.1.2 设置虚拟内核
输出结果如下，因为没有指定线程数，所以默认使用 CPU 核心数量的线程，这里在虚拟机(cpu四个核心)
当然我们也可以指定核心数量，代码如下
#include
int main(void)
{
pragma omp parallel num_threads(6)
{
printf(“Hello, world. \n”);
}
return 0;
}
输出结果如下，你可能会想为啥可以指定大于核心的线程数，不是只有四个核心吗，其实不影响的，CPU也不看你的核心数，只是知道该取指令进行计算，一般使用大于等于核心
这里的 omp.h 主要包括一些 openmp 的库函数，比如 omp_get_thread_num() 获取当前线程 id。
1.1.3 并行计算示例
下面使用多线程完成加法
loop_add.c
#include
#include
#include
int main(void) {
int sum = 0;
pragma omp parallel for
for (int i=1; i<=100; i++) {
sum += i;
}
printf(“%d”, sum);
return 0;
}
执行多次，其结果如下我们可以发现结果每次都不一样原因是线程对同一个资源产生了竞争，sum += i; 这一行如果多个线程同时写，可能会发生写冲突因为 sum+=i 可以展开为 sum=sum+i所以，如果一个线程计算完了 sum+i 但还没赋值过去，这时就会产生问题
那么 OpenMP 为了解决这个冲突，可以使用 reduction
#include
#include
#include
int main(void) {
int sum = 0;
pragma omp parallel for reduction(+:sum)
for (int i=1; i<=100; i++) {
sum += i;
}
printf(“%d”, sum);
return 0;
}
在RK3588 计算1 000,000,000次，主频2.4GHz 8核
耗时
CPU
不使用OpenMP
2.2s
12.5%
使用OpenMP
1.1s
87.5%
1.2 Neno
1.2.1 neon简介
Neon是基于Arm平台对SIMD(Single Instruction Multiple Data)技术的一种实现。Neon 技术为指令集架构提供了专门的扩展，提供了可以对多个数据流并行执行数学运算的附加指令。
NEON技术与Cortex-A8和Cortex-A9处理器相结合，已经被许多领先企业广泛采用。越来越多的机构正在IP设计中采用NEON技术，或提供为NEON技术优化的软件，构成了NEON生态系统的一部分。这一生态系统由大量的硅片领导厂商构成，如博通公司、Freescale Semiconductor、Matsush_ita、NEC、NVIDIA、松下、PMC-Sierra、三星电子、ST、TI和东芝美国电子元器件公司，这些公司都获得了Cortex处理器授权，其中提供NEON技术作为选择之一。此外，在这些ARM处理器和多媒体编解码器厂商基础上，构建开发和评测电路板的硬件设计合作伙伴正不断为NEON技术优化编解码器。
1.2.2 Neon 编程
目前，我们可以通过一下几个路径来使用Neon技术:
1. 基于Neon开发的第三方开源库，比如Arm Compute Library，Ne10，libYuv
2. 编译器的自动矢量化功能可以自动帮助我们使用Neno进行优化
3. Neon intrinsics是一组编译器用来替代Neon 指令的的内建函数，该内建函数风格跟C/C++函数一致，方便普通开发者调用。但该接口的性能并不是最好的。
4. 对于有经验的开发者来说，手写Neon的汇编代码可以将获得最好的性能实现。
在64位Arm架构中，寄存器主要是64位的，但Neon实现单元可以使用单独的128位寄存器进行SIMD处理。Neon的寄存器可以理解位一组寄存器的集合，可以作为8位，16位，32位，64位，128位寄存器使用。
/–1、Vector add(正常指令): vadd -> ri = ai + bi; r, a, b have equal lane sizes–/
int8x8_t vadd_s8 (int8x8_t __a, int8x8_t __b);//_mm_add_epi8
int16x4_t vadd_s16 (int16x4_t __a, int16x4_t __b);//_mm_add_epi16
int32x2_t vadd_s32 (int32x2_t __a, int32x2_t __b);//_mm_add_epi32
int64x1_t vadd_s64 (int64x1_t __a, int64x1_t __b);//_mm_add_epi64
1.2.3 neon 编程实例
编写一个 数列 和 😭 n=4k)
void add_int_c(int* dst, int* src1, int* src2, int count)
{
int i;
for (i = 0; i < count; i++)
dst[i] = src1[i] + src2[i];
}
void add_int_neon(int* dst, int* src1, int* src2, int count)
{
int i;
for (i = 0; i < count; i += 4)
{ //1.define
int32x4_t in1, in2, out; /
2.load
in1 = vld1q_s32(src1);
in2 = vld1q_s32(src2);
//3.calculate
out = vaddq_s32(in1, in2);
//4.send
vst1q_s32(dst, out);
//5.address ++
src1 += 4;
src2 += 4;
dst += 4;
}
}
在RK3588 计算1 000,000,000次，主频2.4GHz 8核
耗时
CPU
不使用neon
2.2s
12.5%
使用neon
1.1s
12.3%
1.3 OpenGL
1.3.1 OpenGL简介
OpenGL（全称：Open Graphics Library，译名：开放图形库或者“开放式图形库”）是用于渲染2D、3D矢量图形的跨语言、跨平台的一种应用程序编程接口（API）一种可以对图形硬件设备特性进行访问的软件库。也就是一种是个专业的图形程序接口，是一个功能强大，调用方便的底层图形库包含了超过500个不同函数调用，可以用于设置所需对象、图像、操作总的来说，OpenGL就是一个开放的图形应用编程接口，有一堆函数组成，以实现对于图形应用程序的开发
早先定义 OpenGL ES 是 OpenGL 的嵌入式设备版本，用于移动端平台（Android、iOS），但由于嵌入式设备要求的是高性能，所以一些其它纯追求高性能的设备也开始用这种 API 方式；OpenGL ES 是 OpenGL 的子集，区别在于 OpenGL ES 删减了 OpenGL 一切低效能的操作方式，有高性能的决不留低效能的，即只求效能不求兼容性（即：OpenGL ES 能实现的，OpenGL 也能实现；OpenGL 部分 API，OpenGL ES 不支持）
1.3.2 OpenGL 七大功能
作为独立于操作系统的开放的三维图形的软件开发包，在其基础上开发的应用程序能够简单方便的移植于各种平台。其具有七大功能：
1、建立3D模型：OpenGL除了能够处理一般的2D图形，即点、线、面的绘制外，主要任务是集合了3D立体的物体绘制函数。
2、图形变换：OpenGL利用基本变换以及投影变换处理图形。所谓的基本变换就是在处理2D平面图形时的平移、旋转、变比、镜像变换。投影变换就是在处理3D立体图形时的平行投影以及透视投影。通过变换方式，可以将2D的平面图形清晰明了的变换成3D的立体图形，从而在减少计算的时间的同时就能够提高了图形显示的速度。
3、颜色模式：OpenGL库中的颜色模型：使用较为广泛的RGBA模式以及颜色索引模式（color index）。
4、光照、材质的设置：OpenGL库中包含了多种光照的类型。材质是用光反射率来表示的。其原理是基于人眼的原理，场景中的物体是由光的红绿蓝的分量以及材质的红绿蓝的反射率的乘积后所形成的颜色值。
5、纹理映射：纹理指的是物体表面的花纹。OpenGL库中集合了对于物体纹理的映射处理方式，能够十分完整的复现物体表面的真实纹理。
6、图像增强功能和位图显示的扩展功能：OpenGL的功能包括像素的读写、复制外，以及一些特殊的图像处理功能：比如，融合、反走样、雾的等等特殊的处理方式。对于图像的重现和处理，可以使得效果更有真实感，逼真。
7、双缓存功能：OpenGL创新性的运用了双缓存形式。计算场景、生成画面图像、显示画面图像分别将其由前台缓存和后台缓存分开处理，大大提高了计算机的运算能力以及画面的显示速度。
总的来说，OpenGL总的主要功能就是绘图，实现对于各种图像的渲染。而这渲染图象细分就主要包括这几个子功能：绘制2D图像3D模型、上色（包括透明度）、光照阴影（设置材质）、映射纹理（绘制花纹）、图像增强处理（更具真实感）、前后双缓存（刷新显示更快）。
不适用原因
• 主要用于处理如3D渲染等比较高级的图像处理。
• 提供的函数接口较为高级
1.4 ISP
RK3588 的ISP ISP30 包含了一系列的图像处理算法模块，主要包括：暗电流矫正、坏点矫正、3A、HDR、镜头阴影矫正、镜头畸变矫正、3DLUT、去噪（包括RAW域去噪，多帧降噪，颜色去噪等）、锐化等。
ISP30包括硬件算法实现及软件逻辑控制部分，RkAiq即为软件逻辑控制部分的实现。
RkAiq软件模块主要实现的功能为：从ISP驱动获取图像统计，结合IQ Tuning参数，使用一系列算法计算出新的ISP、Sensor等硬件参数，不断迭代该过程，最终达到最优的图像效果。
不可用原因
• 库函数接口/服务 需要收取额外费用
• 仅支持部分相机
1.5 RGA
RGA (Raster Graphic Acceleration Unit)是RK3588一个独立的2D硬件加速器，可用于加速点/线绘制，执行图像缩放、旋转、格式转换等常见的2D图形操作。
1.5.1 Scale
缩放，对图像进行缩小放大。组水平、垂直最大支持32倍放大、缩小；通道水平、垂直最大支持32倍放大、缩小。
1.5.2 固定角度旋转
支持 0 度、90 度、180 度以及 270 度固定角度的旋转功能。
1.5.3 Mirror/Flip
Mirror 即水平镜像，Flip 即垂直翻转。可使用 Mirror+Flip 实现 180°旋转。
1.5.4 Overlay
视频叠加区域，调用 VGS 对 VPSS 通道的输出图像叠加位图。
不适合原因
• 支持功能单一，不能定制化开发(仅提供一些简单操作的函数接口)。
1.6 OpenCL
1.6.1 OpenCL简介
OpenCL（Open Computing Language）是一个开放的跨平台并行编程框架，可以用于CPU、GPU、FPGA等多种计算设备。它提供了一种标准化的方式，使得开发人员能够利用计算设备的并行性能，以加速计算密集型应用程序的执行。OpenCL定义了一个基于C语言的编程模型和API，开发人员可以在其中编写称为“内核”的程序，这些内核可以在计算设备上并行执行。在科学计算、图形处理、计算机视觉、机器学习等领域均有应用。
1.6.2 OpenCL具有以下特点
◦ 跨平台支持：OpenCL可以在Windows、Linux、macOS等多种操作系统上运行，而且可以使用多种计算设备。
◦ 灵活性：开发人员可以使用OpenCL编写自己的内核程序，以适应不同的硬件和应用需求。
◦ 性能：通过使用OpenCL，可以充分利用多核CPU和GPU等计算设备的并行性能，从而提高应用程序的执行效率。
◦ 标准化：OpenCL定义了一种通用的编程模型和API，使得开发人员可以更容易地编写并行程序。
2 OPENCL 版本和生态
2.1 发展历史
OpenCL（Open Computing Language）是一种开放的并行计算框架，用于在不同的计算设备上进行高性能计算，包括CPU、GPU、FPGA等。下面是OpenCL的发展历史：
-
2008年6月：OpenCL 1.0规范首次发布，由Khronos Group领导的OpenCL工作组开发。OpenCL 1.0支持在多个硬件平台上进行并行计算。
-
2009年12月：OpenCL 1.0规范正式发布，成为一个开放标准。
-
2011年11月：OpenCL 1.2规范发布，引入了许多新功能和扩展，包括支持动态并行性、共享虚拟内存和用户事件回调等。
-
2013年11月：OpenCL 2.0规范发布，引入了许多重要的新功能，如共享虚拟内存、通用地址空间和内存层次结构等。
-
2015年5月：OpenCL 2.1规范发布，增加了对C++11标准的支持，提供了更多的编程灵活性和功能。
-
2019年11月：OpenCL 2.2规范发布，引入了一些新的扩展和功能，如支持共享虚拟内存的子缓冲区和用于机器学习的新指令集。
2.2 生态
3 适合OpenCL 编程的场景
• 数据并行
• 计算密集
RK3588 GPU ARM Mali-G610，470GFlop
• 计算独立性能强，不需要数据同步
• 占用带宽大
• 小缓存
4．OpenCL 编程示例
4.1 串行代码转并行代码
4.2 Hello World 4.1.1 功能
编写一个 =
示意：
4.1.2 核函数 4.1.3 运行部件之间的关系 4.1.4 工作流 4.1.5 工作流对应函数
创建程序
内存拷贝
内存重映射
在RK3588 计算1 000,000,000次加法，主频2.4GHz 8核
耗时
CPU
不使用OpenMP
2.2s
12.5%
使用OpenMP
1.1s
87.5%
使用neon
1.1s
12.3%
OpenCL+数据拷贝
1.8s
6.25%
在RK3588 计算1 000,000,000次乘法，主频2.4GHz 8核
耗时
CPU
不使用OpenMP
4.9s
12.5%
OpenCL+数据拷贝
2.0s
6.25%
4.1.6 OpenCL 内存 4.3 矩阵相乘 4.3.1 C 语言 4.3.2 OpenCL矩阵相乘初阶 4.3.3 OpenCL矩阵相乘中阶 4 AVM 360拼接融合算法