Rockchip RK3588 - OpenCL环境搭建

一、OpenCL环境搭建
二、OpenCL环境搭建
三、OpenCL测试
- 3.1 platform.cpp
- 3.2 编译
  - 3.2.1 直接编译
  - 3.2.2 cmake编译
四、OpenCV测试
五、FloatVideo-TouchScreen优化
六、代码下载

在上一节《Rockchip RK3588 - 基于Qt的视频监控和控制系统》，我们介绍了实时监控的实现，在实时监控中我们需要将分辨率为1920x1080的图像缩放为指定窗口大小的图像，当采样帧率比较高时，会占用大量的CPU资源；

root@NanoPC-T6:/opt/qt-project/FloatVideo-TouchScreen# export DISPLAY=:0.0;./FloatVideo-TouchScreen -size 0.8
# 打开新的终端
root@NanoPC-T6:~# top
任务: 278 total,   2 running, 276 sleeping,   0 stopped,   0 zombie
%Cpu(s): 36.0 us,  1.9 sy,  0.0 ni, 62.1 id,  0.0 wa,  0.0 hi,  0.0 si,  0.0 st
MiB Mem :  15953.1 total,  14749.5 free,    662.4 used,    541.2 buff/cache
MiB Swap:      0.0 total,      0.0 free,      0.0 used.  14995.4 avail Mem

 进程号 USER      PR  NI    VIRT    RES    SHR    %CPU  %MEM     TIME+ COMMAND
   1513 root      20   0 2876120 127804  72488 S 270.9   0.8   0:27.46 FloatVideo-Touc
    864 root      20   0 3345456 238548 186388 S  14.9   1.5   0:03.11 Xorg
   1251 pi        20   0 1861028  76424  57116 S   2.6   0.5   0:00.96 xfwm4
   ......

那么我们是不是可以通过GPU来实现图像的缩放呢，在RK3588上可以使用OpenCL接口进行GPU加速。

回到顶部

一、`OpenCL`环境搭建

OpenCL（Open Computing Language开放计算语言）是一种开放的、免版税的标准，用于超级计算机、云服务器、个人计算机、移动设备和嵌入式平台中各种加速器的跨平台并行编程。

OpenCL是由Khronos Group创建和管理的。OpenCL使应用程序能够使用系统或设备中的并行处理能力，从而使应用程序运行得更快、更流畅。

1.1 工作原理

OpenCL是一种编程框架和运行时，它使程序员能够创建称为内核程序（或内核）的小程序，这些程序可以在系统中的任何处理器上并行编译和执行。处理器可以是不同类型的任意组合，包括CPU、GPU、DSP、FPGA或张量处理器，这就是为什么OpenCL经常被称为异构并行编程的解决方案。

OpenCL框架包含两个API：

platform layer API：在主机CPU上运行，首先用于使程序能够发现系统中可用的并行处理器或计算设备。通过查询哪些计算设备可用，应用程序可以在不同的系统上便携地运行—适应加速器硬件的不同组合。一旦发现了计算设备，platform layer API就允许应用程序选择并初始化它想要使用的设备；
Runtime API：它使应用程序的内核程序能够为它们将要运行的计算设备编译，并行加载到这些处理器上并执行。一旦内核程序完成执行，将使用Runtime API收集结果；

为了更好适用于不同的处理器，OpenCL抽象出来了四大模型：

平台模型：描述了OpenCL如何理解拓扑连接系统中的计算资源，对不同硬件及软件实现抽象，方便应用于不同设备；
内存模型：对硬件的各种内存器进行了抽象；
执行模型：程序是如何在硬件上执行的；
编程模型：数据并行和任务并行；

1.2 平台模型

OpenCL中，需要一个主机处理器（Host），一般为CPU。而其它的硬件处理器（多核CPU/GPU/DSP等）被抽象成Compute Device；

每个Compute Device包含多个Compute Unit；
每个Compute Unit又包含多个Processing Elements（处理单元）。

举例说明：计算设备可以是GPU，计算单元对应于GPU内部的流多处理器（streaming multiprocessors(SMs)），处理单元对应于每个SM内部的单个流处理器。处理器通常通过共享指令调度和内存资源，以及增加本地处理器间通信，将处理单元分组为计算单元，以提高实现效率。

1.3 内存模型

OpenCL内存模型定义了如何访问和共享不同内核和处理单元之间的数据。

1.3.1 内存类型

OpenCL支持以下内存类型：

Global memory：全局内存对在上下文中执行的所有工作项可访问，主机可以使用__global关键字读取、写入和映射命令访问全局内存，在单个工作组中，全局内存是一致的；
Constant memory：常量内存是用于主机分配和初始化的对象的内存区域，所有工作项都可以以只读方式访问常量内存；
Local memory：本地内存是特定于工作组的，工作组中的工作项可以访问本地内存；使用__local关键字进行访问，对于工作组中的所有工作项来说，本地内存是一致的；
Private memory：私有内存是特定于工作项的，其他工作项无法访问私有内存；

1.3.2 内存模型

OpenCL内存模型如下：

1.4 执行模型

OpenCL执行模型包括主机应用程序、上下文（context）和OpenCL内核的操作。

主机应用程序使用OpenCL命令队列将kernel和数据传输函数发送到设备以执行。

通过将命令入队到命令队列（Command Queues）中，kernel和数据传输函数可以与应用程序主机代码并行异步执行。

1.4.1 主机应用程序

主机应用程序在应用处理器上运行。主机应用程序通过为以下命令设置命令队列来管理内核的执行：

内存命令；
内核执行命令；
同步操作；

1.4.2 上下文

主机应用程序为内核定义上下文。上下文包括：

计算设备（Compute devices）；
内核（Kernels）：OpenCL核心计算部分，类似C语言的代码。在需要设备执行计算任务时，数据会被推送到Compute Device，然后Compute Device的计算单元会并发执行内核程序；
程序对象（Programs）：Kernels的集合，OpenCL中可以使用cl_program表示；
内存对象（Memory Objects.）；

1.4.3 `OpenCL`内核的操作

Kernels在计算设备上运行。kernel是一段代码，在计算设备上与其它内核并行执行。内核的操作按以下顺序进行：

Kernels在主机应用程序中定义；
主机应用程序将kernel提交给计算设备执行。计算设备可以是应用处理器、GPU或其它类型的处理器；
当主机应用程序发出提交kernel的命令时，OpenCL创建工作项的NDRange；
对于NDRange中的每个元素，创建kernel的一个实例。这使得每个元素可以独立并行地进行处理。

1.5 `OpenCL`计算流程

对于OpenCl，利用显卡计算时，需要经历如下步骤：

主机应用程序进行设备初始化（获取平台和设备id，创建上下文和命令队列）；
编写并编译kernel（读取内核文件->创建program对象->编译程序->创建内核）；
主机应用程序准备数据并传入设备（准备主机端数据，创建设备端内存对象并拷贝主机端数据）；
主机应用程序将kernel提交给设备执行（传入kernel函数参数，启动kernel函数）；
将结果拷贝回主机应用程序；
后续处理；
释放资源。

回到顶部

二、`OpenCL`环境搭建

一个完整的OpenCL框架，从内核层到用户层，可分为四部分：

内核层GPU驱动；
用户层动态库；
头文件；
应用程序；

在《Rockchip RK3399 - Mali-T860 GPU驱动（mesa+Panfrost）》文章中我们提到GPU驱动一般分为两部分：

一小部分在linux内核中；
另外一大部分在Userspace，在Userspace的部分向下操作内核中的驱动，向上对应用层提供标准的 API 接口，例如：
- OpenGL ES 1.1、2.0、3.0、3.1、3.2；
- OpenCL 1.1、1.2、2.0；
- Vulkan 1.0；
- RenderScript（受支持的API列表因二进制和GPU类型而异）；

Mail GPU IP提供商ARM公司只开放了内核部分驱动，而且这部分驱动还没有按照linux kernel的规范以DRM的框架去实现，此外Userspace部分ARM没有开源，只是以库的形式提供给购买了Mali GPU授权的SoC厂商。

Mali GPU User-Space Binary Drivers；

2.1 内核层`GPU`驱动

以RK3588为例，搭载了Mail-G610。内核层GPU驱动这一部分，不需要自己移植，我们开发板所使用的的友善linux kernel 6.1已移植。

接下来的内容仅仅是作为扩展，不感兴趣忽略即可。ARM Mail G610采用的Bifrost 驱动微架构，如果需要自己移植内核层GPU驱动的话。可以从Open Source Bifrost Mali 3rd Gen GPU Architecture Kernel Drivers下载最新版本内核层GPU驱动。比如这里我下载了BX304L01B-SW-99002-r47p0-01eac0.tar；

关于如何安装该驱动可以参考：

2.1.1 驱动入口

解压后，可以看到驱动入口文件在drivers/gpu/arm/midgard/mali_kbase_core_linux.c ，在该文件我们可以看到支持的GPU型号；

static const struct of_device_id kbase_dt_ids[] = { { .compatible = "arm,malit6xx" },
                                                    { .compatible = "arm,mali-midgard" },
                                                    { .compatible = "arm,mali-bifrost" },
                                                    { .compatible = "arm,mali-valhall" },
                                                    { /* sentinel */ } };
MODULE_DEVICE_TABLE(of, kbase_dt_ids);

static struct platform_driver kbase_platform_driver = {
        .probe = kbase_platform_device_probe,
        .remove = kbase_platform_device_remove,
        .driver = {
                   .name = kbase_drv_name,
                   .pm = &kbase_pm_ops,
                   .of_match_table = of_match_ptr(kbase_dt_ids),
                   .probe_type = PROBE_PREFER_ASYNCHRONOUS,
        },
};

module_platform_driver(kbase_platform_driver);

2.1.2 `gpu`设备节点

在 NanoPC上安装运行debian操作系统后，使用以下命令检查Mali GPU的设备树节点：

root@NanoPC-T6:~# apt install device-tree-compiler -y
root@NanoPC-T6:~# dtc -I fs /proc/device-tree | grep mali

在 arch/arm64/boot/dts/rockchip/rk3588s.dtsi中我们可以定位到gpu设备节点；

gpu: gpu@fb000000 {
		compatible = "arm,mali-bifrost";
		reg = <0x0 0xfb000000 0x0 0x200000>;
		interrupts = <GIC_SPI 94 IRQ_TYPE_LEVEL_HIGH>,
					 <GIC_SPI 93 IRQ_TYPE_LEVEL_HIGH>,
					 <GIC_SPI 92 IRQ_TYPE_LEVEL_HIGH>;
		interrupt-names = "GPU", "MMU", "JOB";

		clocks = <&scmi_clk SCMI_CLK_GPU>, <&cru CLK_GPU_COREGROUP>,
				 <&cru CLK_GPU_STACKS>, <&cru CLK_GPU>;
		clock-names = "clk_mali", "clk_gpu_coregroup",
					  "clk_gpu_stacks", "clk_gpu";
		assigned-clocks = <&scmi_clk SCMI_CLK_GPU>;
		assigned-clock-rates = <200000000>;
		power-domains = <&power RK3588_PD_GPU>;
		operating-points-v2 = <&gpu_opp_table>;
		#cooling-cells = <2>;
		dynamic-power-coefficient = <2982>;

		upthreshold = <30>;
		downdifferential = <10>;

		status = "disabled";
};

在arch/arm64/boot/dts/rockchip/rk3588-nanopi6-common.dtsi中可以定位到：

&gpu {
        mali-supply = <&vdd_gpu_s0>;
        mem-supply = <&vdd_gpu_mem_s0>;
        upthreshold = <60>;
        downdifferential = <30>;
        status = "okay";
};

其中：

compatible：说明了设备兼容的驱动名称，即"arm,mali-bifrost"；可以看到arm,mali-bifrost是和panfrost驱动相匹配的，因此会执行驱动的.probe函数，这里就不深入研究了；
reg：指定了寄存器的基地址和大小，即基地址0xfb000000，大小为0x200000；
interrupts和interrupt-names：分别指定了该设备所使用的中断号和中断的名称；
clocks：指定了使用哪个时钟控制器（CRU）提供GPU时钟；
power-domains：用于指定设备所属的电源域，即RK3588_PD_GPU；
mali-supply：指定了GPU设备使用的电源管脚；
status：指定GPU设备的状态（"okay" 表示设备正常工作）；

2.2 用户层动态库

寻找官方（Mali ARM/Rockchip ）提供的用户层动态库libmali.so。

2.2.1 `Mali ARM`官方下载安装`libmali.so`

通过浏览器进入Mali ARM官网：https://developer.arm.com/downloads/-/mali-drivers/user-space。

寻找官方提供的用户层动态库libmali.so，libmali.so一般会有不同的版本（X11，fbdev、Wayland等），其提供了OpenGL ES，EGL，OpenCL接口。

不过不幸的是：Mail ARM官网并没有看到适用于RK3588的用户层动态库，但是RK3288的倒是有，这里我们就以RK3288为例：

下载后，解压缩可以看到：

注意：上图中libEGL.so、libOpenCL.so、libGLESv2.so等库大小均为0，不难猜测libmail.so应该提供了OpenGL ES，EGL，OpenCL接口。

将libmali.so存放在ARM的/usr/lib/，同时建立软链接libOpenCL.so指向libmali.so；

root@NanoPC-T6:~# ln -s /usr/lib/libmali.so /usr/lib/libOpenCL.so

2.2.2 `Rockchip`官方提供的`libmali.so`

我们使用的友善提供的debian文件系统已经安装了libmali.so，该用户层动态库是由Rockchip官方提供的。

如何来查看是否已经安装了OpenCL库和驱动，可以通过如下命令检查是否已经安装了libmali.so ；

root@NanoPC-T6:~# find /usr -name libmali.so
/usr/lib/aarch64-linux-gnu/libmali.so
root@NanoPC-T6:~# strings /usr/lib/aarch64-linux-gnu/libmali.so | grep Mali-G610
Mali-G610

root@NanoPC-T6:~# strings /usr/lib/aarch64-linux-gnu/libmali.so | grep cl
.....
clReleaseCommandBufferKHR
clReleaseCommandQueue
clReleaseContext
clReleaseDevice
clReleaseEvent
clReleaseKernel
clReleaseMemObject
.....

root@NanoPC-T6:~# ls -l /usr/lib/aarch64-linux-gnu/libmali.so
lrwxrwxrwx 1 root root 12  7月 29  2020 /usr/lib/aarch64-linux-gnu/libmali.so -> libmali.so.1

其中/usr/lib/aarch64-linux-gnu/libmali.so是libmali.so库的路径，Mali-G610是Mali GPU驱动的版本号。

如果命令输出为空，则说明该库不是Mali GPU驱动库。如果输出包含Mali-G610 字符串，则说明该库是Mali GPU驱动库，并且版本号为Mali-G610。

此外在/usr/lib/aarch64-linux-gnu目录下包含单独的OpenGL ES，EGL，OpenCL库；

root@NanoPC-T6:/opt# ls -l /usr/lib/aarch64-linux-gnu/libOpenCL*
lrwxrwxrwx 1 root root    18  1月 12  2021 /usr/lib/aarch64-linux-gnu/libOpenCL.so.1 -> libOpenCL.so.1.0.0
-rw-r--r-- 1 root root 60856  1月 12  2021 /usr/lib/aarch64-linux-gnu/libOpenCL.so.1.0.0

root@NanoPC-T6:/opt# strings /usr/lib/aarch64-linux-gnu/libOpenCL.so.1.0.0 | grep cl
fclose
closedir
dlclose
clGetExtensionFunctionAddress
clGetPlatformIDs
clCreateContext
clCreateContextFromType
clGetGLContextInfoKHR
......

root@NanoPC-T6:/opt# ls -l /usr/lib/aarch64-linux-gnu/libEGL*
lrwxrwxrwx 1 root root     20  3月 25  2021 /usr/lib/aarch64-linux-gnu/libEGL_mesa.so.0 -> libEGL_mesa.so.0.0.0
-rw-r--r-- 1 root root 259072  3月 25  2021 /usr/lib/aarch64-linux-gnu/libEGL_mesa.so.0.0.0
lrwxrwxrwx 1 root root     11  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so -> libEGL.so.1
lrwxrwxrwx 1 root root     15  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so.1 -> libEGL.so.1.1.0
-rw-r--r-- 1 root root  84416  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so.1.1.0
......

也可以通过如下clinfo命令查看是否已经安装OpenCL库，如果出现下图所示界面，则系统已经安装；

root@NanoPC-T6:~# aptitude install clinfo
root@NanoPC-T6:~# clinfo
arm_release_ver: g13p0-01eac0, rk_so_ver: 10
Number of platforms                               1
  Platform Name                                   ARM Platform
  Platform Vendor                                 ARM
  Platform Version                                OpenCL 3.0 v1.g13p0-01eac0.a8b6f0c7e1f83c654c60d1775112dbe4
  Platform Profile                                FULL_PROFILE
  Platform Extensions                             cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics           
......

Number of devices                                 1
  Device Name                                     Mali-G610 r0p0
  Device Vendor                                   ARM
  Device Vendor ID                                0xa8670000
  Device Version                                  OpenCL 3.0 v1.g13p0-01eac0.a8b6f0c7e1f83c654c60d1775112dbe4
  Device UUID                                     000067a8-0100-0000-0000-000000000000
  Driver UUID                                     13833dc2-ecef-4e5b-0159-38fdaf75bfde
  Valid Device LUID                               No
  Device LUID                                     0000-000000000000
  Device Node Mask                                0
  Device Numeric Version                          0xc00000 (3.0.0)
  Driver Version                                  3.0
  Device OpenCL C Version                         OpenCL C 3.0 v1.g13p0-01eac0.a8b6f0c7e1f83c654c60d1775112dbe4
  Device OpenCL C all versions                    OpenCL C                                                         0x400000 (1.0
......

NULL platform behavior
  clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)  ARM Platform
  clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)   Success [ARM]
  clCreateContext(NULL, ...) [default]            Success [ARM]
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT)  Success (1)
    Platform Name                                 ARM Platform
    Device Name                                   Mali-G610 r0p0            # GPU型号
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU)  Success (1)
    Platform Name                                 ARM Platform
    Device Name                                   Mali-G610 r0p0
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL)  Success (1)
    Platform Name                                 ARM Platform
    Device Name                                   Mali-G610 r0p0

ICD loader properties
  ICD loader Name                                 OpenCL ICD Loader
  ICD loader Vendor                               OCL Icd free software
  ICD loader Version                              2.2.14
  ICD loader Profile                              OpenCL 3.0

其中arm_release_ver: g13p0-01eac0, rk_so_ver: 10为驱动版本信息。

接着我们需要将建立软链接libOpenCL.so指向libmali.so；

root@NanoPC-T6:~# ln -s /usr/lib/aarch64-linux-gnu/libmali.so /usr/lib/aarch64-linux-gnu/libOpenCL.so
root@NanoPC-T6:~# ls -l /usr/lib/aarch64-linux-gnu/libOpenCL.so
lrwxrwxrwx 1 root root 37  1月 16 23:43 /usr/lib/aarch64-linux-gnu/libOpenCL.so -> /usr/lib/aarch64-linux-gnu/libmali.so

2.3 安装头文件

从官网下载头文件OpenCL-Headers：

root@NanoPC-T6:/opt# git clone https://github.com/extdomains/github.com/KhronosGroup/OpenCL-Headers.git

运行以下命令来配置构建过程，并指定安装路径为/usr：

root@NanoPC-T6:/opt/OpenCL-Headers# cmake -S . -B build -DCMAKE_INSTALL_PREFIX=/usr
-- The C compiler identification is GNU 10.2.1
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /usr/bin/cc - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- The CXX compiler identification is GNU 10.2.1
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /usr/bin/c++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Found Python3: /usr/bin/python3.9 (found version "3.9.2") found components: Interpreter
-- Configuring done
-- Generating done
-- Build files have been written to: /opt/OpenCL-Headers/build

其中：

-S .：指定源代码目录的路径；
-B build：指定构建目录的路径；
-DCMAKE_INSTALL_PREFIX=/usr：指定cmake执行install 目标时，安装的路径前缀；

如上命令会让cmake在.目录下查找CMakeLists.txt文件，并在./build目录下生成Makefile文件。

接着运行以下命令在./build目录下执行构建操作，只构建install目标，将生成的文件安装到指定的位置；

root@NanoPC-T6:/opt/OpenCL-Headers# cmake --build build --target install
Scanning dependencies of target headers_c_200
[  0%] Building C object tests/lang_c/CMakeFiles/headers_c_200.dir/__/test_headers.c.o
[  0%] Linking C executable headers_c_200
[  0%] Built target headers_c_200
Scanning dependencies of target headers_c_120
[  1%] Building C object tests/lang_c/CMakeFiles/headers_c_120.dir/__/test_headers.c.o
[  1%] Linking C executable headers_c_120
[  1%] Built target headers_c_120
Scanning dependencies of target cl_version_h_c_300
[  1%] Building C object tests/lang_c/CMakeFiles/cl_version_h_c_300.dir/__/test_cl_version.h.c.o
[  2%] Linking C executable cl_version_h_c_300
.......
[ 99%] Built target cl_egl_h_cpp_100
Scanning dependencies of target cl_gl_h_cpp_120
[100%] Building CXX object tests/lang_cpp/CMakeFiles/cl_gl_h_cpp_120.dir/test_cl_gl.h.cpp.o
[100%] Linking CXX executable cl_gl_h_cpp_120
[100%] Built target cl_gl_h_cpp_120
Install the project...
-- Install configuration: ""
-- Installing: /usr/include/CL
-- Installing: /usr/include/CL/opencl.h
-- Installing: /usr/include/CL/cl_egl.h
-- Installing: /usr/include/CL/cl_ext_intel.h
-- Installing: /usr/include/CL/cl_layer.h
-- Installing: /usr/include/CL/cl_platform.h
-- Installing: /usr/include/CL/cl_d3d10.h
-- Installing: /usr/include/CL/cl_va_api_media_sharing_intel.h
-- Installing: /usr/include/CL/cl_icd.h
-- Installing: /usr/include/CL/cl.h
-- Installing: /usr/include/CL/cl_function_types.h
-- Installing: /usr/include/CL/cl_dx9_media_sharing.h
-- Installing: /usr/include/CL/cl_dx9_media_sharing_intel.h
-- Installing: /usr/include/CL/cl_gl_ext.h
-- Installing: /usr/include/CL/cl_d3d11.h
-- Installing: /usr/include/CL/cl_version.h
-- Installing: /usr/include/CL/cl_half.h
-- Installing: /usr/include/CL/cl_ext.h
-- Installing: /usr/include/CL/cl_gl.h
-- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersTargets.cmake
-- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersConfig.cmake
-- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersConfigVersion.cmake
-- Installing: /usr/share/pkgconfig/OpenCL-Headers.pc

头文件已经安装到/usr/include/CL目录下：

root@NanoPC-T6:/opt/OpenCL-Headers# ls -l /usr/include/CL
总用量 392
-rw-r--r-- 1 root root   8057  1月 15 00:10 cl_d3d10.h
-rw-r--r-- 1 root root   8095  1月 15 00:10 cl_d3d11.h
-rw-r--r-- 1 root root  12246  1月 15 00:10 cl_dx9_media_sharing.h
-rw-r--r-- 1 root root    959  1月 15 00:10 cl_dx9_media_sharing_intel.h
-rw-r--r-- 1 root root   5672  1月 15 00:10 cl_egl.h
-rw-r--r-- 1 root root 127490  1月 15 00:10 cl_ext.h
-rw-r--r-- 1 root root    902  1月 15 00:10 cl_ext_intel.h
-rw-r--r-- 1 root root  33387  1月 15 00:10 cl_function_types.h
-rw-r--r-- 1 root root    905  1月 15 00:10 cl_gl_ext.h
-rw-r--r-- 1 root root  12040  1月 15 00:10 cl_gl.h
-rw-r--r-- 1 root root  81631  1月 15 00:10 cl.h
-rw-r--r-- 1 root root  10430  1月 15 00:10 cl_half.h
-rw-r--r-- 1 root root  11505  1月 15 00:10 cl_icd.h
-rw-r--r-- 1 root root   3544  1月 15 00:10 cl_layer.h
-rw-r--r-- 1 root root  43430  1月 15 00:10 cl_platform.h
-rw-r--r-- 1 root root   7090  1月 15 00:10 cl_va_api_media_sharing_intel.h
-rw-r--r-- 1 root root   3125  1月 15 00:10 cl_version.h
-rw-r--r-- 1 root root    970  1月 15 00:10 opencl.h

回到顶部

三、`OpenCL`测试

此时已经有动态库和头文件，可以进行测试了。在/opt/目录下创建opencl-project文件夹;

root@NanoPC-T6:/opt# mkdir opencl-project

接着创建platform文件夹；

root@NanoPC-T6:/opt# cd opencl-project/
root@NanoPC-T6:/opt/opencl-project# mkdir platform
root@NanoPC-T6:/opt/opencl-project# cd platform

3.1 `platform.cpp`

在/opt/opencl-project/platform目录下编写测试代码platform.cpp；

#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>

#define MAX_PLATFORMS 10
#define MAX_DEVICES 10

int main() {
    cl_platform_id platforms[MAX_PLATFORMS];
    cl_device_id devices[MAX_DEVICES];
    cl_uint num_platforms, num_devices;
    cl_context context;
    cl_command_queue command_queue;
    cl_program program;
    cl_kernel kernel;
    cl_int ret;

    // 获取平台数量
    ret = clGetPlatformIDs(MAX_PLATFORMS, platforms, &num_platforms);
    if (ret != CL_SUCCESS) {
        printf("Failed to get platform IDs\n");
        return -1;
    }

    printf("Number of platforms: %u\n", num_platforms);

    // 遍历打印平台信息
    for (cl_uint i = 0; i < num_platforms; i++) {
        char platform_name[128];
        char platform_vendor[128];

        ret = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(platform_name), platform_name, NULL);
        if (ret != CL_SUCCESS) {
            printf("Failed to get platform name for platform %u\n", i);
        }

        ret = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, NULL);
        if (ret != CL_SUCCESS) {
            printf("Failed to get platform vendor for platform %u\n", i);
        }

        printf("Platform %u:\n", i);
        printf("    Name: %s\n", platform_name);
        printf("    Vendor: %s\n", platform_vendor);
        printf("\n");
    }

    // 获取设备数量
    ret = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, MAX_DEVICES, devices, &num_devices);
    if (ret != CL_SUCCESS) {
        printf("Failed to get device IDs\n");
        return -1;
    }

    // 创建OpenCL上下文
    context = clCreateContext(NULL, num_devices, devices, NULL, NULL, &ret);
    if (ret != CL_SUCCESS) {
        printf("Failed to create context\n");
        return -1;
    }

    // 创建命令队列
    command_queue = clCreateCommandQueue(context, devices[0], 0, &ret);
    if (ret != CL_SUCCESS) {
        printf("Failed to create command queue\n");
        return -1;
    }

    // 定义和构建OpenCL内核
    const char *kernel_source = "__kernel void hello_world() {\n"
                                "    printf(\"Hello, World!\\n\");\n"
                                "}\n";
    program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &ret);
    if (ret != CL_SUCCESS) {
        printf("Failed to create program\n");
        return -1;
    }

    ret = clBuildProgram(program, num_devices, devices, NULL, NULL, NULL);
    if (ret != CL_SUCCESS) {
        printf("Failed to build program\n");
        return -1;
    }

    // 创建OpenCL内核对象
    kernel = clCreateKernel(program, "hello_world", &ret);
    if (ret != CL_SUCCESS) {
        printf("Failed to create kernel\n");
        return -1;
    }

    // 执行内核函数
    ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
    if (ret != CL_SUCCESS) {
        printf("Failed to enqueue task\n");
        return -1;
    }

    // 等待执行完成
    ret = clFinish(command_queue);
    if (ret != CL_SUCCESS) {
        printf("Failed to finish execution\n");
        return -1;
    }

    printf("Kernel executed successfully\n");

    // 清理资源
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);

    return 0;
}

3.2 编译

这里我们介绍两种源码编译的方式。

3.2.1 直接编译

我们可以直接执行如下编译命令:

root@NanoPC-T6:/opt/opencl-project/platform# gcc platform.cpp -o platform -lmali

-lmail用于链接libmali.so库文件，-l选项指定要链接的库文件名，并在文件名前加上lib和.so的前缀和后缀。所以-lmali告诉编译器要链接的库文件名为libmali.so。

那么编译器如何知道libmali.so在哪里的呢？

首先搜索预定义的默认路径，如/usr/lib和/usr/local/lib等;
如果共享库没有在这些路径中找到，则会搜索在/etc/ld.so.conf和/etc/ld.so.conf.d目录中指定的路径。这些路径可以包含自定义共享库路径，比如：

root@NanoPC-T6:/opt/opencl-project/platform# ls -l /etc/ld.so.conf.d/
总用量 12
-rw-r--r-- 1 root root  32  7月 29  2020 00-aarch64-mali.conf
-rw-r--r-- 1 root root 103  4月 20  2023 aarch64-linux-gnu.conf
-rw-r--r-- 1 root root  44  9月 23  2022 libc.conf
root@NanoPC-T6:/opt/opencl-project/platform#  cat /etc/ld.so.conf.d/aarch64-linux-gnu.conf
# Multiarch support
/usr/local/lib/aarch64-linux-gnu
/lib/aarch64-linux-gnu
/usr/lib/aarch64-linux-gnu     # 该路径下有libmali.so库文件

3.2.2 `cmake`编译

当然也可以使用cmake进行编译platform.cpp，接下来我们介绍cmake编译配置。

(1) 在/opt/opencl-project/platform目录下创建CMakeLists.txt：

cmake_minimum_required(VERSION 3.0)
cmake_policy(VERSION 3.0...3.18.4)
project(proj)
add_executable(platform platform.cpp)
#寻找OpenCL库  /usr/share/cmake-3.18/Modules/FindOpenCL.cmake
find_package(OpenCL REQUIRED)
#打印调试信息
MESSAGE(STATUS "Project: ${PROJECT_NAME}")
MESSAGE(STATUS "OpenCL library status:")
MESSAGE(STATUS "    version: ${OpenCL_VERSION_STRING}")
MESSAGE(STATUS "    libraries: ${OpenCL_LIBRARY}")
MESSAGE(STATUS "    include path: ${OpenCL_INCLUDE_DIR}")

target_link_libraries(platform PRIVATE OpenCL::OpenCL)

(2) 配置构建过程：

root@NanoPC-T6:/opt/opencl-project/platform#  cmake -S . -B build
-- The C compiler identification is GNU 10.2.1
-- The CXX compiler identification is GNU 10.2.1
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /usr/bin/cc - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /usr/bin/c++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Looking for CL_VERSION_2_2
-- Looking for CL_VERSION_2_2 - found
-- Found OpenCL: /usr/lib/aarch64-linux-gnu/libOpenCL.so (found version "2.2")
-- Project: proj
-- OpenCL library status:
--     version: 2.2 
--     libraries: /usr/lib/aarch64-linux-gnu/libOpenCL.so         # 库文件路径
--     include path: /usr/include  # 头文件路径
-- Configuring done
-- Generating done
-- Build files have been written to: /opt/opencl-project/platform/build

其中：

-S .：选项指定源代码目录的路径，CMake将在该路径下查找CMakeLists.txt文件；
-B build：选项指定构建目录的路径；

实际上我们使用的版本是OpenCL 3.0，这里判定为2.2版本是因为cmake version 3.18.4 FindOpenCL.cmake能够识别的最大版本为2.2，其通过在CL/cl.h文件查找CL_VERSION_${VERSION}宏来判定安装的版本的。

可以通过修改/usr/share/cmake-3.18/Modules/FindOpenCL.cmake解决这个问题：

foreach(VERSION "3_0" "2_2" "2_1" "2_0" "1_2" "1_1" "1_0")

(3) 执行构建操作，生成可执行程序platform；

root@NanoPC-T6:/opt/OpenCL-Headers/exmaples# cmake --build build
Scanning dependencies of target platform
[ 50%] Building CXX object CMakeFiles/platform.dir/platform.cpp.o
In file included from /usr/include/CL/cl.h:20,
                 from /usr/include/CL/opencl.h:24,
                 from /opt/OpenCL-Headers/exmaples/platform.cpp:1:
/usr/include/CL/cl_version.h:22:104: note: ‘#pragma message: cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 300 (OpenCL 3.0)’
   22 | #pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 300 (OpenCL 3.0)")
      |                                                                                                        ^
[100%] Linking CXX executable platform
[100%] Built target platform

执行程序:

root@NanoPC-T6:/opt/opencl-project/platform# ls -l build/
总用量 48
-rw-r--r-- 1 root root 14229  1月 16 23:45 CMakeCache.txt
drwxr-xr-x 5 root root  4096  1月 16 23:46 CMakeFiles
-rw-r--r-- 1 root root  1632  1月 16 23:45 cmake_install.cmake
-rw-r--r-- 1 root root  5253  1月 16 23:45 Makefile
-rwxr-xr-x 1 root root 14248  1月 16 23:46 platform

root@NanoPC-T6:/opt/opencl-project/platform# ./build/platform
arm_release_ver: g13p0-01eac0, rk_so_ver: 10
Number of platforms: 1
Platform 0:
    Name: ARM Platform
    Vendor: ARM

Kernel executed successfully

回到顶部

四、`OpenCV`测试

4.1 `OCL`介绍

OpenCV在2011年开始与AMD合作加入OpenCL加速。因此，OpenCV-2.4.3版本包含了新的ocl模块，其中包含了一些现有OpenCV算法的OpenCL实现。也就是说，当客户端机器上有OpenCL运行时和兼容设备时，用户可以调用cv::ocl::resize()来使用加速的代码，而不是使用cv::resize()。在接下来的三年中，越来越多的函数和类被添加到ocl模块中；但它始终是一个独立的API，与OpenCV-2.x中的主要面向CPU的API并存。

在OpenCV-3.x中，架构概念已经改变为所谓的Transparent API（T-API）。在新架构中，一个单独的OpenCL加速的cv::ocl::resize() 已经从外部API中移除，而成为常规cv::resize()中的一个分支。这个分支在性能角度上自动调用，并在可行且有意义时进行优化。T-API的实现得到了AMD和Intel公司的赞助。

4.1.1 代码示例

Regular CPU code：

// initialization
VideoCapture vcap(...);
CascadeClassifier fd("haar_ff.xml");
Mat frame, frameGray;
vector<rect> faces;
for(;;){
  // processing loop
  vcap >> frame;
  cvtColor(frame, frameGray, BGR2GRAY);
  equalizeHist(frameGray, frameGray);
  fd.detectMultiScale(frameGray, faces, ...);
  // draw rectangles …
  // show image …
}

OpenCL-aware code OpenCV-2.x：

// initialization
VideoCapture vcap(...);
ocl::OclCascadeClassifier fd("haar_ff.xml");
ocl::oclMat frame, frameGray;
Mat frameCpu;
vector<rect> faces;
for(;;){
  // processing loop
  vcap >> frameCpu;
  frame = frameCpu;
  ocl::cvtColor(frame, frameGray, BGR2GRAY);
  ocl::equalizeHist(frameGray, frameGray);
  fd.detectMultiScale(frameGray, faces, ...);
  // draw rectangles …
  // show image …
}

OpenCL-aware code OpenCV-3.x：

// initialization
VideoCapture vcap(...);
CascadeClassifier fd("haar_ff.xml");
UMat frame, frameGray;
vector<rect> faces;
for(;;){
  // processing loop
  vcap >> frame;
  cvtColor(frame, frameGray, BGR2GRAY);
  equalizeHist(frameGray, frameGray);
  fd.detectMultiScale(frameGray, faces, ...);
  // draw rectangles …
  // show image …
}

相比于OpenCV-2.x，OpenCV-3.x封装了一个新的数据类型cv::UMat，这个数据类型能够无缝对接OpenCV的普通接口，从而最少的改动代码而最大的完成OpenCL平台的加速功能。

4.1.2 原理

如上面代码所示，只需要将原来的Mat格式换为UMat格式就可以实现OpenCV函数在OpenCL设备上加速运行，而这其中具体实施的基本原理是什么呢？接下来看一下其底层实现的基本原理，具体参看OpenCV 3.1中OpenCL部分实现的源代码（注意：每个版本代码可能略有差异）：

上图中表明了，当你使用的数据类型是UMat {dst.isUmat()}，并且开启了OpenCL使能{useOpenCL()}，那么OpenCV的接口将会跳转到OpenCL支持的设备中进行加速运行，当然你需要注意的是，在第一次使用OpenCL加速程序时，OpenCL需要编译生成对应平台的Kernel代码，而编译是需要花费大量的时间的，因此初次运行需要比较长的时间。

4.2 官方`opencv-ocl`

在/opt/opencl-project目录下新建opencv-ocl项目，源码位于：https://521github.com/opencv/opencv/tree/3.4.0/samples/opencl。

4.2.1 `main.cpp`

亲爱的读者和支持者们，自动博客加入了打赏功能，陆陆续续收到了各位老铁的打赏。在此，我想由衷地感谢每一位对我们博客的支持和打赏。你们的慷慨与支持，是我们前行的动力与源泉。

日期	姓名	金额
2023-09-06	*源	19
2023-09-11	*朝科	88
2023-09-21	*号	5
2023-09-16	*真	60
2023-10-26	*通	9.9
2023-11-04	*慎	0.66
2023-11-24	*恩	0.01
2023-12-30	I*B	1
2024-01-28	*兴	20
2024-02-01	QYing	20
2024-02-11	*督	6
2024-02-18	一*x	1
2024-02-20	c*l	18.88
2024-01-01	*I	5
2024-04-08	*程	150
2024-04-18	*超	20
2024-04-26	.*V	30
2024-05-08	D*W	5
2024-05-29	*辉	20
2024-05-30	*雄	10
2024-06-08	*:	10
2024-06-23	小狮子	666
2024-06-28	*s	6.66
2024-06-29	*炼	1
2024-06-30	*!	1
2024-07-08	*方	20
2024-07-18	A*1	6.66
2024-07-31	*北	12
2024-08-13	*基	1
2024-08-23	n*s	2
2024-09-02	*源	50
2024-09-04	*J	2
2024-09-06	*强	8.8
2024-09-09	*波	1
2024-09-10	*口	1
2024-09-10	*波	1
2024-09-12	*波	10
2024-09-18	*明	1.68
2024-09-26	B*h	10
2024-09-30	岁	10
2024-10-02	M*i	1
2024-10-14	*朋	10
2024-10-22	*海	10
2024-10-23	*南	10
2024-10-26	*节	6.66
2024-10-27	*o	5
2024-10-28	W*F	6.66
2024-10-29	R*n	6.66
2024-11-02	*球	6
2024-11-021	*鑫	6.66
2024-11-25	*沙	5
2024-11-29	C*n	2.88

创作不易，喜欢的话，请考虑支持一下！关注公众号查看全文！（微信扫码）

posted @ 2024-01-15 00:25 大奥特曼打小怪兽阅读(6643) 评论(0) 编辑收藏举报

刷新页面返回顶部

登录后才能查看或发表评论，立即登录或者逛逛博客园首页

相关博文：

· Rockchip RK3588 - linux下Qt和opencv交叉编译环境搭建

· Rockchip RK3399 - Mali-T860 GPU驱动（mesa+Panfrost）

· OpenCV 图像处理学习手册：6~7

· RK3588编译Opencv最新版

· opencv的一个小例子

阅读排行：
· 全程不用写代码，我用AI程序员写了一个飞机大战
· MongoDB 8.0这个新功能碉堡了，比商业数据库还牛
· DeepSeek 开源周回顾「GitHub 热点速览」
· 记一次.NET内存居高不下排查解决与启示
· 白话解读 Dapr 1.15：你的「微服务管家」又秀新绝活了

大奥特曼打小怪兽

Rockchip RK3588 - OpenCL环境搭建

一、OpenCL环境搭建

1.1 工作原理

1.2 平台模型

1.3 内存模型

1.3.1 内存类型

1.3.2 内存模型

1.4 执行模型

1.4.1 主机应用程序

1.4.2 上下文

1.4.3 OpenCL内核的操作

1.5 OpenCL计算流程

二、OpenCL环境搭建

2.1 内核层GPU驱动

2.1.1 驱动入口

2.1.2 gpu设备节点

2.2 用户层动态库

2.2.1 Mali ARM官方下载安装libmali.so

2.2.2 Rockchip官方提供的libmali.so

2.3 安装头文件

三、OpenCL测试

3.1 platform.cpp

3.2 编译

3.2.1 直接编译

3.2.2 cmake编译

四、OpenCV测试

4.1 OCL介绍

4.1.1 代码示例

4.1.2 原理

4.2 官方opencv-ocl

4.2.1 main.cpp

公告

创作不易，喜欢的话，请考虑关注一下公众号

积分与排名

随笔分类 (538)

随笔档案 (376)

阅读排行榜

推荐排行榜

最新评论

欢迎打赏支持我 ^_^

一、`OpenCL`环境搭建

1.4.3 `OpenCL`内核的操作

1.5 `OpenCL`计算流程

二、`OpenCL`环境搭建

2.1 内核层`GPU`驱动

2.1.2 `gpu`设备节点

2.2.1 `Mali ARM`官方下载安装`libmali.so`

2.2.2 `Rockchip`官方提供的`libmali.so`

三、`OpenCL`测试

3.1 `platform.cpp`

3.2.2 `cmake`编译

四、`OpenCV`测试

4.1 `OCL`介绍

4.2 官方`opencv-ocl`

4.2.1 `main.cpp`