devbox
我家小花儿
这个屌丝很懒，什么也没留下！
关注作者
热门标签
热门文章
当前位置:   article > 正文

Rockchip RK3588 - OpenCL环境搭建

作者：我家小花儿 | 2024-05-04 17:19:41

赞

踩

firefly rk3588 opencl

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

  1. root@NanoPC-T6:/opt/qt-project/FloatVideo-TouchScreen# export DISPLAY=:0.0;./FloatVideo-TouchScreen -size 0.8
  2. root@NanoPC-T6:~# top
  3. 任务: 278 total,   2 running, 276 sleeping,   0 stopped,   0 zombie
  4. %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
  5. MiB Mem :  15953.1 total,  14749.5 free,    662.4 used,    541.2 buff/cache
  6. MiB Swap:      0.0 total,      0.0 free,      0.0 used.  14995.4 avail Mem
  7.  
  8.  进程号 USER      PR  NI    VIRT    RES    SHR    %CPU  %MEM     TIME+ COMMAND
  9.    1513 root      20   0 2876120 127804  72488 S 270.9   0.8   0:27.46 FloatVideo-Touc
  10.     864 root      20   0 3345456 238548 186388 S  14.9   1.5   0:03.11 Xorg
  11.    1251 pi        20   0 1861028  76424  57116 S   2.6   0.5   0:00.96 xfwm4
  12.    ......

那么我们是不是可以通过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驱动；
  • 用户层动态库；
  • 头文件；
  • 应用程序；
2.1 内核层GPU驱动

以RK3588为例，搭载了Mail-G610 GPU，Linux内核提供了针对Mali-T6xx / Mali-T7xx / Mali-T8xx GPU和GXX系列的Panfrost驱动，具体可以参考《Rockchip RK3399 - Mali-T860 GPU驱动》；

注意： 内核层GPU驱动这一部分，不需要自己移植，我们开发板所使用的的友善linux kernel 6.1已移植；

2.2 用户层动态库

用户层动态库有多种途径可以获得，比如以下两种：

  • 寻找官方(Mali ARM/Rockchip )提供的用户层动态库libmali.so；

  • 下载KhronosGroup OpenCL-SDK源码，并编译，可以得到libOpenCL.so；

下面我们分别介绍这两种方式， 对于这两种方式我们选择一种即可，对于我使用的NanoPC-T6采用第一种方式(默认已经支持)。

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等)，其提供了opengles，egl，opencl接口。

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

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

注意：上图中libEGL.so、libOpenCL.so、libGLESv2.so等库大小均为0，不难猜测libmail.so应该提供了opengles，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 ；

  1. root@NanoPC-T6:~# find /usr -name libmali.so
  2. /usr/lib/aarch64-linux-gnu/libmali.so
  3. root@NanoPC-T6:~# strings /usr/lib/aarch64-linux-gnu/libmali.so | grep Mali-G610
  4. Mali-G610
  5.  
  6. root@NanoPC-T6:~# strings /usr/lib/aarch64-linux-gnu/libmali.so | grep cl
  7. .....
  8. clReleaseCommandBufferKHR
  9. clReleaseCommandQueue
  10. clReleaseContext
  11. clReleaseDevice
  12. clReleaseEvent
  13. clReleaseKernel
  14. clReleaseMemObject
  15. .....
  16.  
  17. root@NanoPC-T6:~# ls -l /usr/lib/aarch64-linux-gnu/libmali.so
  18. 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目录下包含单独的opengles，egl，opencl库；

  1. root@NanoPC-T6:/opt# ls -l /usr/lib/aarch64-linux-gnu/libOpenCL*
  2. lrwxrwxrwx 1 root root    18  1月 12  2021 /usr/lib/aarch64-linux-gnu/libOpenCL.so.1 -> libOpenCL.so.1.0.0
  3. -rw-r--r-- 1 root root 60856  1月 12  2021 /usr/lib/aarch64-linux-gnu/libOpenCL.so.1.0.0
  4.  
  5. root@NanoPC-T6:/opt# strings /usr/lib/aarch64-linux-gnu/libOpenCL.so.1.0.0 | grep cl
  6. fclose
  7. closedir
  8. dlclose
  9. clGetExtensionFunctionAddress
  10. clGetPlatformIDs
  11. clCreateContext
  12. clCreateContextFromType
  13. clGetGLContextInfoKHR
  14. ......
  15.  
  16. root@NanoPC-T6:/opt# ls -l /usr/lib/aarch64-linux-gnu/libEGL*
  17. lrwxrwxrwx 1 root root     20  3月 25  2021 /usr/lib/aarch64-linux-gnu/libEGL_mesa.so.0 -> libEGL_mesa.so.0.0.0
  18. -rw-r--r-- 1 root root 259072  3月 25  2021 /usr/lib/aarch64-linux-gnu/libEGL_mesa.so.0.0.0
  19. lrwxrwxrwx 1 root root     11  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so -> libEGL.so.1
  20. lrwxrwxrwx 1 root root     15  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so.1 -> libEGL.so.1.1.0
  21. -rw-r--r-- 1 root root  84416  7月 29  2020 /usr/lib/aarch64-linux-gnu/libEGL.so.1.1.0
  22. ......

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

  1. root@NanoPC-T6:~# aptitude install clinfo
  2. root@NanoPC-T6:~# clinfo
  3. arm_release_ver: g13p0-01eac0, rk_so_ver: 10
  4. Number of platforms                               1
  5.   Platform Name                                   ARM Platform
  6.   Platform Vendor                                 ARM
  7.   Platform Version                                OpenCL 3.0 v1.g13p0-01eac0.a8b6f0c7e1f83c654c60d1775112dbe4
  8.   Platform Profile                                FULL_PROFILE
  9.   Platform Extensions                             cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics           ......
  10. NULL platform behavior
  11.   clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)  ARM Platform
  12.   clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)   Success [ARM]
  13.   clCreateContext(NULL, ...) [default]            Success [ARM]
  14.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT)  Success (1)
  15.     Platform Name                                 ARM Platform
  16.     Device Name                                   Mali-G610 r0p0            # GPU型号
  17.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU)  No devices found in platform
  18.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU)  Success (1)
  19.     Platform Name                                 ARM Platform
  20.     Device Name                                   Mali-G610 r0p0
  21.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR)  No devices found in platform
  22.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM)  No devices found in platform
  23.   clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL)  Success (1)
  24.     Platform Name                                 ARM Platform
  25.     Device Name                                   Mali-G610 r0p0
  26.  
  27. ICD loader properties
  28.   ICD loader Name                                 OpenCL ICD Loader
  29.   ICD loader Vendor                               OCL Icd free software
  30.   ICD loader Version                              2.2.14
  31.   ICD loader Profile                              OpenCL 3.0

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

  1. root@NanoPC-T6:~# ln -s /usr/lib/aarch64-linux-gnu/libmali.so /usr/lib/aarch64-linux-gnu/libOpenCL.so
  2. root@NanoPC-T6:~# ls -l /usr/lib/aarch64-linux-gnu/libOpenCL.so
  3. lrwxrwxrwx 1 root root 37  1月 16 23:43 /usr/lib/aarch64-linux-gnu/libOpenCL.so -> /usr/lib/aarch64-linux-gnu/libmali.so
2.2.3 OpenCL SDK编译安装

如果没有安装，请按照如下步骤安装：下载OpenCL SDK进行编译安装，具体可以参考《OpenCL安装过程记录》。

Khronos Group的OpenCL SDK是一个通用的官方开发工具包，适用于多个硬件平台，而AMD和Intel等硬件供应商提供的OpenCL SDK则更专注于其特定硬件平台的优化和支持。根据您的需求和使用的硬件平台，选择适合的OpenCL SDK可以帮助您获得最佳的性能和开发体验。

下载源码：

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

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

  1. root@NanoPC-T6:/opt/OpenCL-SDK# cmake -S . -B build -DCMAKE_INSTALL_PREFIX=/opt/OpenCL
  2. -- The C compiler identification is GNU 10.2.1
  3. -- The CXX compiler identification is GNU 10.2.1
  4. -- Detecting C compiler ABI info
  5. -- Detecting C compiler ABI info - done
  6. -- Check for working C compiler: /usr/bin/cc - skipped
  7. -- Detecting C compile features
  8. -- Detecting C compile features - done
  9. -- Detecting CXX compiler ABI info
  10. -- Detecting CXX compiler ABI info - done
  11. -- Check for working CXX compiler: /usr/bin/c++ - skipped
  12. -- Detecting CXX compile features
  13. -- Detecting CXX compile features - done
  14. -- No build type selected, default to Release
  15. -- Looking for pthread.h
  16. -- Looking for pthread.h - found
  17. -- Performing Test CMAKE_HAVE_LIBC_PTHREAD
  18. -- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
  19. -- Looking for pthread_create in pthreads
  20. -- Looking for pthread_create in pthreads - not found
  21. -- Looking for pthread_create in pthread
  22. -- Looking for pthread_create in pthread - found
  23. -- Found Threads: TRUE
  24. -- Looking for secure_getenv
  25. -- Looking for secure_getenv - found
  26. -- Looking for __secure_getenv
  27. -- Looking for __secure_getenv - not found
  28. -- Check if compiler accepts -pthread
  29. -- Check if compiler accepts -pthread - yes
  30. -- Could NOT find Doxygen (missing: DOXYGEN_EXECUTABLE)
  31. -- cargs (https://521github.com/likle/cargs) not found. To self-host, set cargs_INCLUDE_PATH and cargs_LIBRARY to point to the headers and library respectively adding '-D cargs_INCLUDE_PATH=/path/to/cargs/include/dir -D cargs_LIBRARY/path/to/cargs/libcargs' to the cmake command. (missing: cargs_INCLUDE_PATH cargs_LIBRARY)
  32. -- Fetching cargs.
  33. -- Adding cargs subproject: /opt/OpenCL-SDK/build/_deps/cargs-external-src
  34. -- TCLAP (http://tclap.sourceforge.net/) not found. To self-host, set TCLAP_INCLUDE_PATH to point to the headers adding '-DTCLAP_INCLUDE_PATH=/path/to/tclap' to the cmake command. (missing: TCLAP_INCLUDE_PATH)
  35. -- Fetching TCLAP.
  36. -- Found TCLAP: /opt/OpenCL-SDK/build/_deps/tclap-external-src/include
  37. -- Stb (https://521github.com/nothings/stb) not found. To self-host, set Stb_INCLUDE_PATH to point to the headers adding '-D Stb_INCLUDE_PATH=/path/to/stb' to the cmake command. (missing: Stb_INCLUDE_PATH)
  38. -- Fetching Stb.
  39. -- Found Stb: /opt/OpenCL-SDK/build/_deps/stb-external-src
  40. -- Found X11: /usr/include
  41. -- Looking for XOpenDisplay in /usr/lib/aarch64-linux-gnu/libX11.so;/usr/lib/aarch64-linux-gnu/libXext.so
  42. -- Looking for XOpenDisplay in /usr/lib/aarch64-linux-gnu/libX11.so;/usr/lib/aarch64-linux-gnu/libXext.so - found
  43. -- Looking for gethostbyname
  44. -- Looking for gethostbyname - found
  45. -- Looking for connect
  46. -- Looking for connect - found
  47. -- Looking for remove
  48. -- Looking for remove - found
  49. -- Looking for shmat
  50. -- Looking for shmat - found
  51. -- Looking for IceConnectionNumber in ICE
  52. -- Looking for IceConnectionNumber in ICE - found
  53. -- Could NOT find glm (missing: glm_DIR)
  54. -- Fetching glm.
  55. -- Adding glm subproject: /opt/OpenCL-SDK/build/_deps/glm-external-src
  56. CMake Warning (dev) at /usr/share/cmake-3.18/Modules/FindOpenGL.cmake:305 (message):
  57.   Policy CMP0072 is not set: FindOpenGL prefers GLVND by default when
  58.   available.  Run "cmake --help-policy CMP0072" for policy details.  Use the
  59.   cmake_policy command to set the policy and suppress this warning.
  60.  
  61.   FindOpenGL found both a legacy GL library:
  62.  
  63.     OPENGL_gl_LIBRARY: /usr/lib/aarch64-linux-gnu/libGL.so
  64.  
  65.   and GLVND libraries for OpenGL and GLX:
  66.  
  67.     OPENGL_opengl_LIBRARY: /usr/lib/aarch64-linux-gnu/libOpenGL.so
  68.     OPENGL_glx_LIBRARY: /usr/lib/aarch64-linux-gnu/libGLX.so
  69.  
  70.   OpenGL_GL_PREFERENCE has not been set to "GLVND" or "LEGACY", so for
  71.   compatibility with CMake 3.10 and below the legacy GL library will be used.
  72. Call Stack (most recent call first):
  73.   cmake/Dependencies/OpenGL/OpenGL.cmake:1 (find_package)
  74.   cmake/Dependencies.cmake:17 (include)
  75.   CMakeLists.txt:50 (include)
  76. This warning is for project developers.  Use -Wno-dev to suppress it.
  77.  
  78. -- Found OpenGL: /usr/lib/aarch64-linux-gnu/libOpenGL.so
  79. -- Could NOT find GLEW (missing: GLEW_INCLUDE_DIRS GLEW_LIBRARIES)
  80. -- Fetching GLEW.
  81. -- Adding GLEW subproject: /opt/OpenCL-SDK/build/_deps/glew-external-src
  82. CMake Warning (dev) at build/_deps/glew-external-src/CMakeLists.txt:2 (project):
  83.   Policy CMP0048 is not set: project() command manages VERSION variables.
  84.   Run "cmake --help-policy CMP0048" for policy details.  Use the cmake_policy
  85.   command to set the policy and suppress this warning.
  86.  
  87.   The following variable(s) would be set to empty:
  88.  
  89.     PROJECT_VERSION
  90.     PROJECT_VERSION_MAJOR
  91.     PROJECT_VERSION_MINOR
  92.     PROJECT_VERSION_PATCH
  93. This warning is for project developers.  Use -Wno-dev to suppress it.
  94.  
  95. -- Found Freetype: /usr/lib/aarch64-linux-gnu/libfreetype.so (found version "2.10.4")
  96. -- Fetching SFML.
  97. -- Adding SFML subproject: /opt/OpenCL-SDK/build/_deps/sfml-external-src
  98. -- libudev stable: 1
  99. -- Found UDev: /usr/lib/aarch64-linux-gnu/libudev.so
  100. --    include: /usr/include
  101. -- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY
  102. -- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY - Success
  103. -- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY
  104. -- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY - Success
  105. -- Performing Test COMPILER_HAS_DEPRECATED_ATTR
  106. -- Performing Test COMPILER_HAS_DEPRECATED_ATTR - Success
  107. -- Looking for sin in m
  108. -- Looking for sin in m - found
  109. -- Configuring done
  110. -- Generating done
  111. -- Build files have been written to: /opt/OpenCL-SDK/build
  112. root@NanoPC-T6:/opt/OpenCL-SDK#

其中：

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

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

root@NanoPC-T6:/opt/OpenCL-SDK# cmake --build build --target install

编译完成之后，我们查看安装目录：

  1. root@NanoPC-T6:/opt/OpenCL-SDK# ls /opt/OpenCL -l
  2. 总用量 16
  3. drwxr-xr-x 2 root root 4096  1月 16 20:44 bin  
  4. drwxr-xr-x 5 root root 4096  1月 16 20:44 include # 头文件
  5. drwxr-xr-x 4 root root 4096  1月 16 20:44 lib   # 库文件
  6. drwxr-xr-x 5 root root 4096  1月 16 20:44 share
  7.  
  8. root@NanoPC-T6:/opt/OpenCL-SDK# ls -l /opt/OpenCL/lib
  9. 总用量 4564
  10. drwxr-xr-x 4 root root    4096  1月 16 20:44 cmake
  11. -rw-r--r-- 1 root root    4842  1月 16 20:41 libcargs.a
  12. -rw-r--r-- 1 root root 1215506  1月 16 20:42 libglew.a
  13. lrwxrwxrwx 1 root root      23  1月 16 20:44 libglew-shared.so -> libglew-shared.so.2.2.0
  14. -rw-r--r-- 1 root root  961392  1月 16 20:42 libglew-shared.so.2.2.0
  15. -rw-r--r-- 1 root root   90550  1月 16 20:43 libOpenCLExt.a
  16. -rw-r--r-- 1 root root 1269528  1月 16 20:43 libOpenCLSDKCpp.so
  17. -rw-r--r-- 1 root root  205392  1月 16 20:43 libOpenCLSDK.so
  18. lrwxrwxrwx 1 root root      14  1月 16 20:44 libOpenCL.so -> libOpenCL.so.1
  19. lrwxrwxrwx 1 root root      16  1月 16 20:44 libOpenCL.so.1 -> libOpenCL.so.1.2
  20. -rw-r--r-- 1 root root   74744  1月 16 20:41 libOpenCL.so.1.2
  21. -rw-r--r-- 1 root root   61152  1月 16 20:42 libOpenCLUtilsCpp.so
  22. -rw-r--r-- 1 root root   27096  1月 16 20:42 libOpenCLUtils.so
  23. lrwxrwxrwx 1 root root      23  1月 16 20:44 libsfml-graphics.so -> libsfml-graphics.so.2.5
  24. lrwxrwxrwx 1 root root      25  1月 16 20:44 libsfml-graphics.so.2.5 -> libsfml-graphics.so.2.5.1
  25. -rw-r--r-- 1 root root  456128  1月 16 20:42 libsfml-graphics.so.2.5.1
  26. lrwxrwxrwx 1 root root      21  1月 16 20:44 libsfml-system.so -> libsfml-system.so.2.5
  27. lrwxrwxrwx 1 root root      23  1月 16 20:44 libsfml-system.so.2.5 -> libsfml-system.so.2.5.1
  28. -rw-r--r-- 1 root root   71592  1月 16 20:42 libsfml-system.so.2.5.1
  29. lrwxrwxrwx 1 root root      21  1月 16 20:44 libsfml-window.so -> libsfml-window.so.2.5
  30. lrwxrwxrwx 1 root root      23  1月 16 20:44 libsfml-window.so.2.5 -> libsfml-window.so.2.5.1
  31. -rw-r--r-- 1 root root  202536  1月 16 20:42 libsfml-window.so.2.5.1
  32. drwxr-xr-x 2 root root    4096  1月 16 20:44 pkgconfig
  33. root@NanoPC-T6:/opt/OpenCL-SDK# ls -l /opt/OpenCL/include/
  34. 总用量 20
  35. -rw-r--r-- 1 root root 4553  1月 16 20:40 cargs.h
  36. drwxr-xr-x 3 root root 4096  1月 16 20:44 CL
  37. drwxr-xr-x 2 root root 4096  1月 16 20:44 GL
  38. drwxr-xr-x 7 root root 4096  1月 16 20:44 SFML
  39. root@NanoPC-T6:/opt/OpenCL-SDK# ls -l /opt/OpenCL/include/CL/
  40. 总用量 788
  41. -rw-r--r-- 1 root root    786  1月 16 20:38 cl2.hpp
  42. -rw-r--r-- 1 root root   8057  1月 16 20:38 cl_d3d10.h
  43. -rw-r--r-- 1 root root   8095  1月 16 20:38 cl_d3d11.h
  44. -rw-r--r-- 1 root root  12246  1月 16 20:38 cl_dx9_media_sharing.h
  45. -rw-r--r-- 1 root root    959  1月 16 20:38 cl_dx9_media_sharing_intel.h
  46. -rw-r--r-- 1 root root   5672  1月 16 20:38 cl_egl.h
  47. -rw-r--r-- 1 root root 127490  1月 16 20:38 cl_ext.h
  48. -rw-r--r-- 1 root root    902  1月 16 20:38 cl_ext_intel.h
  49. -rw-r--r-- 1 root root  33387  1月 16 20:38 cl_function_types.h
  50. -rw-r--r-- 1 root root    905  1月 16 20:38 cl_gl_ext.h
  51. -rw-r--r-- 1 root root  12040  1月 16 20:38 cl_gl.h
  52. -rw-r--r-- 1 root root  81631  1月 16 20:38 cl.h
  53. -rw-r--r-- 1 root root  10430  1月 16 20:38 cl_half.h
  54. -rw-r--r-- 1 root root  11505  1月 16 20:38 cl_icd.h
  55. -rw-r--r-- 1 root root   3544  1月 16 20:38 cl_layer.h
  56. -rw-r--r-- 1 root root  43430  1月 16 20:38 cl_platform.h
  57. -rw-r--r-- 1 root root   7090  1月 16 20:38 cl_va_api_media_sharing_intel.h
  58. -rw-r--r-- 1 root root   3125  1月 16 20:38 cl_version.h
  59. -rw-r--r-- 1 root root    970  1月 16 20:38 opencl.h
  60. -rw-r--r-- 1 root root 396735  1月 16 20:38 opencl.hpp
  61. drwxr-xr-x 2 root root   4096  1月 16 20:44 Utils

接着我们将库文件和头文件放置到/usr路径下：

  1. sudo ln -s /opt/OpenCL/include/CL /usr/include
  2. sudo ln -s /opt/OpenCL/include/GL /usr/include
  3. sudo ln -s /opt/OpenCL/include/SFML /usr/include
  4. sudo ln -s /opt/OptnCL/lib/libOpenCL.so /usr/lib
2.3 安装头文件

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

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

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

  1. root@NanoPC-T6:/opt/OpenCL-Headers# cmake -S . -B build -DCMAKE_INSTALL_PREFIX=/usr
  2. -- The C compiler identification is GNU 10.2.1
  3. -- Detecting C compiler ABI info
  4. -- Detecting C compiler ABI info - done
  5. -- Check for working C compiler: /usr/bin/cc - skipped
  6. -- Detecting C compile features
  7. -- Detecting C compile features - done
  8. -- The CXX compiler identification is GNU 10.2.1
  9. -- Detecting CXX compiler ABI info
  10. -- Detecting CXX compiler ABI info - done
  11. -- Check for working CXX compiler: /usr/bin/c++ - skipped
  12. -- Detecting CXX compile features
  13. -- Detecting CXX compile features - done
  14. -- Found Python3: /usr/bin/python3.9 (found version "3.9.2") found components: Interpreter
  15. -- Configuring done
  16. -- Generating done
  17. -- 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目标，将生成的文件安装到指定的位置；

  1. root@NanoPC-T6:/opt/OpenCL-Headers# cmake --build build --target install
  2. Scanning dependencies of target headers_c_200
  3. [  0%] Building C object tests/lang_c/CMakeFiles/headers_c_200.dir/__/test_headers.c.o
  4. [  0%] Linking C executable headers_c_200
  5. [  0%] Built target headers_c_200
  6. Scanning dependencies of target headers_c_120
  7. [  1%] Building C object tests/lang_c/CMakeFiles/headers_c_120.dir/__/test_headers.c.o
  8. [  1%] Linking C executable headers_c_120
  9. [  1%] Built target headers_c_120
  10. Scanning dependencies of target cl_version_h_c_300
  11. [  1%] Building C object tests/lang_c/CMakeFiles/cl_version_h_c_300.dir/__/test_cl_version.h.c.o
  12. [  2%] Linking C executable cl_version_h_c_300
  13. .......
  14. [ 99%] Built target cl_egl_h_cpp_100
  15. Scanning dependencies of target cl_gl_h_cpp_120
  16. [100%] Building CXX object tests/lang_cpp/CMakeFiles/cl_gl_h_cpp_120.dir/test_cl_gl.h.cpp.o
  17. [100%] Linking CXX executable cl_gl_h_cpp_120
  18. [100%] Built target cl_gl_h_cpp_120
  19. Install the project...
  20. -- Install configuration: ""
  21. -- Installing: /usr/include/CL
  22. -- Installing: /usr/include/CL/opencl.h
  23. -- Installing: /usr/include/CL/cl_egl.h
  24. -- Installing: /usr/include/CL/cl_ext_intel.h
  25. -- Installing: /usr/include/CL/cl_layer.h
  26. -- Installing: /usr/include/CL/cl_platform.h
  27. -- Installing: /usr/include/CL/cl_d3d10.h
  28. -- Installing: /usr/include/CL/cl_va_api_media_sharing_intel.h
  29. -- Installing: /usr/include/CL/cl_icd.h
  30. -- Installing: /usr/include/CL/cl.h
  31. -- Installing: /usr/include/CL/cl_function_types.h
  32. -- Installing: /usr/include/CL/cl_dx9_media_sharing.h
  33. -- Installing: /usr/include/CL/cl_dx9_media_sharing_intel.h
  34. -- Installing: /usr/include/CL/cl_gl_ext.h
  35. -- Installing: /usr/include/CL/cl_d3d11.h
  36. -- Installing: /usr/include/CL/cl_version.h
  37. -- Installing: /usr/include/CL/cl_half.h
  38. -- Installing: /usr/include/CL/cl_ext.h
  39. -- Installing: /usr/include/CL/cl_gl.h
  40. -- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersTargets.cmake
  41. -- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersConfig.cmake
  42. -- Installing: /usr/share/cmake/OpenCLHeaders/OpenCLHeadersConfigVersion.cmake
  43. -- Installing: /usr/share/pkgconfig/OpenCL-Headers.pc

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

  1. root@NanoPC-T6:/opt/OpenCL-Headers# ls -l /usr/include/CL
  2. 总用量 392
  3. -rw-r--r-- 1 root root   8057  1月 15 00:10 cl_d3d10.h
  4. -rw-r--r-- 1 root root   8095  1月 15 00:10 cl_d3d11.h
  5. -rw-r--r-- 1 root root  12246  1月 15 00:10 cl_dx9_media_sharing.h
  6. -rw-r--r-- 1 root root    959  1月 15 00:10 cl_dx9_media_sharing_intel.h
  7. -rw-r--r-- 1 root root   5672  1月 15 00:10 cl_egl.h
  8. -rw-r--r-- 1 root root 127490  1月 15 00:10 cl_ext.h
  9. -rw-r--r-- 1 root root    902  1月 15 00:10 cl_ext_intel.h
  10. -rw-r--r-- 1 root root  33387  1月 15 00:10 cl_function_types.h
  11. -rw-r--r-- 1 root root    905  1月 15 00:10 cl_gl_ext.h
  12. -rw-r--r-- 1 root root  12040  1月 15 00:10 cl_gl.h
  13. -rw-r--r-- 1 root root  81631  1月 15 00:10 cl.h
  14. -rw-r--r-- 1 root root  10430  1月 15 00:10 cl_half.h
  15. -rw-r--r-- 1 root root  11505  1月 15 00:10 cl_icd.h
  16. -rw-r--r-- 1 root root   3544  1月 15 00:10 cl_layer.h
  17. -rw-r--r-- 1 root root  43430  1月 15 00:10 cl_platform.h
  18. -rw-r--r-- 1 root root   7090  1月 15 00:10 cl_va_api_media_sharing_intel.h
  19. -rw-r--r-- 1 root root   3125  1月 15 00:10 cl_version.h
  20. -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文件夹；

  1. root@NanoPC-T6:/opt# cd opencl-project/
  2. root@NanoPC-T6:/opt/opencl-project# mkdir platform
  3. root@NanoPC-T6:/opt/opencl-project# cd platform
3.1 platform.cpp

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

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <CL/cl.h>
  4.  
  5. #define MAX_PLATFORMS 10
  6. #define MAX_DEVICES 10
  7.  
  8. int main() {
  9.     cl_platform_id platforms[MAX_PLATFORMS];
  10.     cl_device_id devices[MAX_DEVICES];
  11.     cl_uint num_platforms, num_devices;
  12.     cl_context context;
  13.     cl_command_queue command_queue;
  14.     cl_program program;
  15.     cl_kernel kernel;
  16.     cl_int ret;
  17.  
  18.     // 获取平台数量
  19.     ret = clGetPlatformIDs(MAX_PLATFORMS, platforms, &num_platforms);
  20.     if (ret != CL_SUCCESS) {
  21.         printf("Failed to get platform IDs\n");
  22.         return -1;
  23.     }
  24.  
  25.     printf("Number of platforms: %u\n", num_platforms);
  26.  
  27.     // 遍历打印平台信息
  28.     for (cl_uint i = 0; i < num_platforms; i++) {
  29.         char platform_name[128];
  30.         char platform_vendor[128];
  31.  
  32.         ret = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(platform_name), platform_name, NULL);
  33.         if (ret != CL_SUCCESS) {
  34.             printf("Failed to get platform name for platform %u\n", i);
  35.         }
  36.  
  37.         ret = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, NULL);
  38.         if (ret != CL_SUCCESS) {
  39.             printf("Failed to get platform vendor for platform %u\n", i);
  40.         }
  41.  
  42.         printf("Platform %u:\n", i);
  43.         printf("    Name: %s\n", platform_name);
  44.         printf("    Vendor: %s\n", platform_vendor);
  45.         printf("\n");
  46.     }
  47.  
  48.     // 获取设备数量
  49.     ret = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, MAX_DEVICES, devices, &num_devices);
  50.     if (ret != CL_SUCCESS) {
  51.         printf("Failed to get device IDs\n");
  52.         return -1;
  53.     }
  54.  
  55.     // 创建OpenCL上下文
  56.     context = clCreateContext(NULL, num_devices, devices, NULL, NULL, &ret);
  57.     if (ret != CL_SUCCESS) {
  58.         printf("Failed to create context\n");
  59.         return -1;
  60.     }
  61.  
  62.     // 创建命令队列
  63.     command_queue = clCreateCommandQueue(context, devices[0], 0, &ret);
  64.     if (ret != CL_SUCCESS) {
  65.         printf("Failed to create command queue\n");
  66.         return -1;
  67.     }
  68.  
  69.     // 定义和构建OpenCL内核
  70.     const char *kernel_source = "__kernel void hello_world() {\n"
  71.                                 "    printf(\"Hello, World!\\n\");\n"
  72.                                 "}\n";
  73.     program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &ret);
  74.     if (ret != CL_SUCCESS) {
  75.         printf("Failed to create program\n");
  76.         return -1;
  77.     }
  78.  
  79.     ret = clBuildProgram(program, num_devices, devices, NULL, NULL, NULL);
  80.     if (ret != CL_SUCCESS) {
  81.         printf("Failed to build program\n");
  82.         return -1;
  83.     }
  84.  
  85.     // 创建OpenCL内核对象
  86.     kernel = clCreateKernel(program, "hello_world", &ret);
  87.     if (ret != CL_SUCCESS) {
  88.         printf("Failed to create kernel\n");
  89.         return -1;
  90.     }
  91.  
  92.     // 执行内核函数
  93.     ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
  94.     if (ret != CL_SUCCESS) {
  95.         printf("Failed to enqueue task\n");
  96.         return -1;
  97.     }
  98.  
  99.     // 等待执行完成
  100.     ret = clFinish(command_queue);
  101.     if (ret != CL_SUCCESS) {
  102.         printf("Failed to finish execution\n");
  103.         return -1;
  104.     }
  105.  
  106.     printf("Kernel executed successfully\n");
  107.  
  108.     // 清理资源
  109.     ret = clReleaseKernel(kernel);
  110.     ret = clReleaseProgram(program);
  111.     ret = clReleaseCommandQueue(command_queue);
  112.     ret = clReleaseContext(context);
  113.  
  114.     return 0;
  115. }
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目录中指定的路径。这些路径可以包含自定义共享库路径，比如：
  1. root@NanoPC-T6:/opt/opencl-project/platform# ls -l /etc/ld.so.conf.d/
  2. 总用量 12
  3. -rw-r--r-- 1 root root  32  7月 29  2020 00-aarch64-mali.conf
  4. -rw-r--r-- 1 root root 103  4月 20  2023 aarch64-linux-gnu.conf
  5. -rw-r--r-- 1 root root  44  9月 23  2022 libc.conf
  6. root@NanoPC-T6:/opt/opencl-project/platform#  cat /etc/ld.so.conf.d/aarch64-linux-gnu.conf
  7. # Multiarch support
  8. /usr/local/lib/aarch64-linux-gnu
  9. /lib/aarch64-linux-gnu
  10. /usr/lib/aarch64-linux-gnu     # 该路径下有libmali.so库文件
3.2.2 cmake编译

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

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

  1. cmake_minimum_required(VERSION 3.0)
  2. cmake_policy(VERSION 3.0...3.18.4)
  3. project(proj)
  4. add_executable(platform platform.cpp)
  5. #寻找OpenCL库  /usr/share/cmake-3.18/Modules/FindOpenCL.cmake
  6. find_package(OpenCL REQUIRED)
  7. #打印调试信息
  8. MESSAGE(STATUS "Project: ${PROJECT_NAME}")
  9. MESSAGE(STATUS "OpenCL library status:")
  10. MESSAGE(STATUS "    version: ${OpenCL_VERSION_STRING}")
  11. MESSAGE(STATUS "    libraries: ${OpenCL_LIBRARY}")
  12. MESSAGE(STATUS "    include path: ${OpenCL_INCLUDE_DIR}")
  13.  
  14. target_link_libraries(platform PRIVATE OpenCL::OpenCL)

(2) 配置构建过程：

  1. root@NanoPC-T6:/opt/opencl-project/platform#  cmake -S . -B build
  2. -- The C compiler identification is GNU 10.2.1
  3. -- The CXX compiler identification is GNU 10.2.1
  4. -- Detecting C compiler ABI info
  5. -- Detecting C compiler ABI info - done
  6. -- Check for working C compiler: /usr/bin/cc - skipped
  7. -- Detecting C compile features
  8. -- Detecting C compile features - done
  9. -- Detecting CXX compiler ABI info
  10. -- Detecting CXX compiler ABI info - done
  11. -- Check for working CXX compiler: /usr/bin/c++ - skipped
  12. -- Detecting CXX compile features
  13. -- Detecting CXX compile features - done
  14. -- Looking for CL_VERSION_2_2
  15. -- Looking for CL_VERSION_2_2 - found
  16. -- Found OpenCL: /usr/lib/aarch64-linux-gnu/libOpenCL.so (found version "2.2")
  17. -- Project: proj
  18. -- OpenCL library status:
  19. --     version: 2.2 
  20. --     libraries: /usr/lib/aarch64-linux-gnu/libOpenCL.so         # 库文件路径
  21. --     include path: /usr/include  # 头文件路径
  22. -- Configuring done
  23. -- Generating done
  24. -- 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；

  1. root@NanoPC-T6:/opt/OpenCL-Headers/exmaples# cmake --build build
  2. Scanning dependencies of target platform
  3. [ 50%] Building CXX object CMakeFiles/platform.dir/platform.cpp.o
  4. In file included from /usr/include/CL/cl.h:20,
  5.                  from /usr/include/CL/opencl.h:24,
  6.                  from /opt/OpenCL-Headers/exmaples/platform.cpp:1:
  7. /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)’
  8.    22 | #pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 300 (OpenCL 3.0)")
  9.       |                                                                                                        ^
  10. [100%] Linking CXX executable platform
  11. [100%] Built target platform

执行程序:

  1. root@NanoPC-T6:/opt/opencl-project/platform# ls -l build/
  2. 总用量 48
  3. -rw-r--r-- 1 root root 14229  1月 16 23:45 CMakeCache.txt
  4. drwxr-xr-x 5 root root  4096  1月 16 23:46 CMakeFiles
  5. -rw-r--r-- 1 root root  1632  1月 16 23:45 cmake_install.cmake
  6. -rw-r--r-- 1 root root  5253  1月 16 23:45 Makefile
  7. -rwxr-xr-x 1 root root 14248  1月 16 23:46 platform
  8.  
  9. root@NanoPC-T6:/opt/opencl-project/platform# ./build/platform
  10. arm_release_ver: g13p0-01eac0, rk_so_ver: 10
  11. Number of platforms: 1
  12. Platform 0:
  13.     Name: ARM Platform
  14.     Vendor: ARM
  15.  
  16. Kernel executed successfully

四、OpenCV测试用例

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

4.1 OCL介绍
4.2 项目源码
4.2.1 main.c
点击查看代码 
  1. /*
  2. // The example of interoperability between OpenCL and OpenCV.
  3. // This will loop through frames of video either from input media file
  4. // or camera device and do processing of these data in OpenCL and then
  5. // in OpenCV. In OpenCL it does inversion of pixels in left half of frame and
  6. // in OpenCV it does bluring in the right half of frame.
  7. */
  8. #include <cstdio>
  9. #include <cstdlib>
  10. #include <iostream>
  11. #include <fstream>
  12. #include <string>
  13. #include <sstream>
  14. #include <iomanip>
  15. #include <stdexcept>
  16.  
  17. #define CL_USE_DEPRECATED_OPENCL_2_0_APIS // eliminate build warning
  18.  
  19. #if __APPLE__
  20. #include <OpenCL/cl.h>
  21. #else
  22. #include <CL/cl.h>
  23. #endif
  24.  
  25. #include <opencv2/core/ocl.hpp>
  26. #include <opencv2/core/utility.hpp>
  27. #include <opencv2/video.hpp>
  28. #include <opencv2/highgui.hpp>
  29. #include <opencv2/imgproc.hpp>
  30.  
  31.  
  32. using namespace std;
  33. using namespace cv;
  34.  
  35. namespace opencl {
  36.  
  37. class PlatformInfo
  38. {
  39. public:
  40.     PlatformInfo()
  41.     {}
  42.  
  43.     ~PlatformInfo()
  44.     {}
  45.  
  46.     cl_int QueryInfo(cl_platform_id id)
  47.     {
  48.         query_param(id, CL_PLATFORM_PROFILE, m_profile);
  49.         query_param(id, CL_PLATFORM_VERSION, m_version);
  50.         query_param(id, CL_PLATFORM_NAME, m_name);
  51.         query_param(id, CL_PLATFORM_VENDOR, m_vendor);
  52.         query_param(id, CL_PLATFORM_EXTENSIONS, m_extensions);
  53.         return CL_SUCCESS;
  54.     }
  55.  
  56.     std::string Profile()    { return m_profile; }
  57.     std::string Version()    { return m_version; }
  58.     std::string Name()       { return m_name; }
  59.     std::string Vendor()     { return m_vendor; }
  60.     std::string Extensions() { return m_extensions; }
  61.  
  62. private:
  63.     cl_int query_param(cl_platform_id id, cl_platform_info param, std::string& paramStr)
  64.     {
  65.         cl_int res;
  66.  
  67.         size_t psize;
  68.         cv::AutoBuffer<char> buf;
  69.  
  70.         res = clGetPlatformInfo(id, param, 0, 0, &psize);
  71.         if (CL_SUCCESS != res)
  72.             throw std::runtime_error(std::string("clGetPlatformInfo failed"));
  73.  
  74.         buf.resize(psize);
  75.         res = clGetPlatformInfo(id, param, psize, buf, 0);
  76.         if (CL_SUCCESS != res)
  77.             throw std::runtime_error(std::string("clGetPlatformInfo failed"));
  78.  
  79.         // just in case, ensure trailing zero for ASCIIZ string
  80.         buf[psize] = 0;
  81.  
  82.         paramStr = buf;
  83.  
  84.         return CL_SUCCESS;
  85.     }
  86.  
  87. private:
  88.     std::string m_profile;
  89.     std::string m_version;
  90.     std::string m_name;
  91.     std::string m_vendor;
  92.     std::string m_extensions;
  93. };
  94.  
  95.  
  96. class DeviceInfo
  97. {
  98. public:
  99.     DeviceInfo()
  100.     {}
  101.  
  102.     ~DeviceInfo()
  103.     {}
  104.  
  105.     cl_int QueryInfo(cl_device_id id)
  106.     {
  107.         query_param(id, CL_DEVICE_TYPE, m_type);
  108.         query_param(id, CL_DEVICE_VENDOR_ID, m_vendor_id);
  109.         query_param(id, CL_DEVICE_MAX_COMPUTE_UNITS, m_max_compute_units);
  110.         query_param(id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, m_max_work_item_dimensions);
  111.         query_param(id, CL_DEVICE_MAX_WORK_ITEM_SIZES, m_max_work_item_sizes);
  112.         query_param(id, CL_DEVICE_MAX_WORK_GROUP_SIZE, m_max_work_group_size);
  113.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, m_preferred_vector_width_char);
  114.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, m_preferred_vector_width_short);
  115.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, m_preferred_vector_width_int);
  116.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, m_preferred_vector_width_long);
  117.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, m_preferred_vector_width_float);
  118.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, m_preferred_vector_width_double);
  119. #if defined(CL_VERSION_1_1)
  120.         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF, m_preferred_vector_width_half);
  121.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR, m_native_vector_width_char);
  122.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT, m_native_vector_width_short);
  123.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, m_native_vector_width_int);
  124.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, m_native_vector_width_long);
  125.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT, m_native_vector_width_float);
  126.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE, m_native_vector_width_double);
  127.         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF, m_native_vector_width_half);
  128. #endif
  129.         query_param(id, CL_DEVICE_MAX_CLOCK_FREQUENCY, m_max_clock_frequency);
  130.         query_param(id, CL_DEVICE_ADDRESS_BITS, m_address_bits);
  131.         query_param(id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, m_max_mem_alloc_size);
  132.         query_param(id, CL_DEVICE_IMAGE_SUPPORT, m_image_support);
  133.         query_param(id, CL_DEVICE_MAX_READ_IMAGE_ARGS, m_max_read_image_args);
  134.         query_param(id, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, m_max_write_image_args);
  135. #if defined(CL_VERSION_2_0)
  136.         query_param(id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, m_max_read_write_image_args);
  137. #endif
  138.         query_param(id, CL_DEVICE_IMAGE2D_MAX_WIDTH, m_image2d_max_width);
  139.         query_param(id, CL_DEVICE_IMAGE2D_MAX_HEIGHT, m_image2d_max_height);
  140.         query_param(id, CL_DEVICE_IMAGE3D_MAX_WIDTH, m_image3d_max_width);
  141.         query_param(id, CL_DEVICE_IMAGE3D_MAX_HEIGHT, m_image3d_max_height);
  142.         query_param(id, CL_DEVICE_IMAGE3D_MAX_DEPTH, m_image3d_max_depth);
  143. #if defined(CL_VERSION_1_2)
  144.         query_param(id, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, m_image_max_buffer_size);
  145.         query_param(id, CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, m_image_max_array_size);
  146. #endif
  147.         query_param(id, CL_DEVICE_MAX_SAMPLERS, m_max_samplers);
  148. #if defined(CL_VERSION_1_2)
  149.         query_param(id, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, m_image_pitch_alignment);
  150.         query_param(id, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, m_image_base_address_alignment);
  151. #endif
  152. #if defined(CL_VERSION_2_0)
  153.         query_param(id, CL_DEVICE_MAX_PIPE_ARGS, m_max_pipe_args);
  154.         query_param(id, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, m_pipe_max_active_reservations);
  155.         query_param(id, CL_DEVICE_PIPE_MAX_PACKET_SIZE, m_pipe_max_packet_size);
  156. #endif
  157.         query_param(id, CL_DEVICE_MAX_PARAMETER_SIZE, m_max_parameter_size);
  158.         query_param(id, CL_DEVICE_MEM_BASE_ADDR_ALIGN, m_mem_base_addr_align);
  159.         query_param(id, CL_DEVICE_SINGLE_FP_CONFIG, m_single_fp_config);
  160. #if defined(CL_VERSION_1_2)
  161.         query_param(id, CL_DEVICE_DOUBLE_FP_CONFIG, m_double_fp_config);
  162. #endif
  163.         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, m_global_mem_cache_type);
  164.         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, m_global_mem_cacheline_size);
  165.         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, m_global_mem_cache_size);
  166.         query_param(id, CL_DEVICE_GLOBAL_MEM_SIZE, m_global_mem_size);
  167.         query_param(id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, m_max_constant_buffer_size);
  168.         query_param(id, CL_DEVICE_MAX_CONSTANT_ARGS, m_max_constant_args);
  169. #if defined(CL_VERSION_2_0)
  170.         query_param(id, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, m_max_global_variable_size);
  171.         query_param(id, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, m_global_variable_preferred_total_size);
  172. #endif
  173.         query_param(id, CL_DEVICE_LOCAL_MEM_TYPE, m_local_mem_type);
  174.         query_param(id, CL_DEVICE_LOCAL_MEM_SIZE, m_local_mem_size);
  175.         query_param(id, CL_DEVICE_ERROR_CORRECTION_SUPPORT, m_error_correction_support);
  176. #if defined(CL_VERSION_1_1)
  177.         query_param(id, CL_DEVICE_HOST_UNIFIED_MEMORY, m_host_unified_memory);
  178. #endif
  179.         query_param(id, CL_DEVICE_PROFILING_TIMER_RESOLUTION, m_profiling_timer_resolution);
  180.         query_param(id, CL_DEVICE_ENDIAN_LITTLE, m_endian_little);
  181.         query_param(id, CL_DEVICE_AVAILABLE, m_available);
  182.         query_param(id, CL_DEVICE_COMPILER_AVAILABLE, m_compiler_available);
  183. #if defined(CL_VERSION_1_2)
  184.         query_param(id, CL_DEVICE_LINKER_AVAILABLE, m_linker_available);
  185. #endif
  186.         query_param(id, CL_DEVICE_EXECUTION_CAPABILITIES, m_execution_capabilities);
  187.         query_param(id, CL_DEVICE_QUEUE_PROPERTIES, m_queue_properties);
  188. #if defined(CL_VERSION_2_0)
  189.         query_param(id, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES, m_queue_on_host_properties);
  190.         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES, m_queue_on_device_properties);
  191.         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, m_queue_on_device_preferred_size);
  192.         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, m_queue_on_device_max_size);
  193.         query_param(id, CL_DEVICE_MAX_ON_DEVICE_QUEUES, m_max_on_device_queues);
  194.         query_param(id, CL_DEVICE_MAX_ON_DEVICE_EVENTS, m_max_on_device_events);
  195. #endif
  196. #if defined(CL_VERSION_1_2)
  197.         query_param(id, CL_DEVICE_BUILT_IN_KERNELS, m_built_in_kernels);
  198. #endif
  199.         query_param(id, CL_DEVICE_PLATFORM, m_platform);
  200.         query_param(id, CL_DEVICE_NAME, m_name);
  201.         query_param(id, CL_DEVICE_VENDOR, m_vendor);
  202.         query_param(id, CL_DRIVER_VERSION, m_driver_version);
  203.         query_param(id, CL_DEVICE_PROFILE, m_profile);
  204.         query_param(id, CL_DEVICE_VERSION, m_version);
  205. #if defined(CL_VERSION_1_1)
  206.         query_param(id, CL_DEVICE_OPENCL_C_VERSION, m_opencl_c_version);
  207. #endif
  208.         query_param(id, CL_DEVICE_EXTENSIONS, m_extensions);
  209. #if defined(CL_VERSION_1_2)
  210.         query_param(id, CL_DEVICE_PRINTF_BUFFER_SIZE, m_printf_buffer_size);
  211.         query_param(id, CL_DEVICE_PREFERRED_INTEROP_USER_SYNC, m_preferred_interop_user_sync);
  212.         query_param(id, CL_DEVICE_PARENT_DEVICE, m_parent_device);
  213.         query_param(id, CL_DEVICE_PARTITION_MAX_SUB_DEVICES, m_partition_max_sub_devices);
  214.         query_param(id, CL_DEVICE_PARTITION_PROPERTIES, m_partition_properties);
  215.         query_param(id, CL_DEVICE_PARTITION_AFFINITY_DOMAIN, m_partition_affinity_domain);
  216.         query_param(id, CL_DEVICE_PARTITION_TYPE, m_partition_type);
  217.         query_param(id, CL_DEVICE_REFERENCE_COUNT, m_reference_count);
  218. #endif
  219.         return CL_SUCCESS;
  220.     }
  221.  
  222.     std::string Name() { return m_name; }
  223.  
  224. private:
  225.     template<typename T>
  226.     cl_int query_param(cl_device_id id, cl_device_info param, T& value)
  227.     {
  228.         cl_int res;
  229.         size_t size = 0;
  230.  
  231.         res = clGetDeviceInfo(id, param, 0, 0, &size);
  232.         if (CL_SUCCESS != res && size != 0)
  233.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  234.  
  235.         if (0 == size)
  236.             return CL_SUCCESS;
  237.  
  238.         if (sizeof(T) != size)
  239.             throw std::runtime_error(std::string("clGetDeviceInfo: param size mismatch"));
  240.  
  241.         res = clGetDeviceInfo(id, param, size, &value, 0);
  242.         if (CL_SUCCESS != res)
  243.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  244.  
  245.         return CL_SUCCESS;
  246.     }
  247.  
  248.     template<typename T>
  249.     cl_int query_param(cl_device_id id, cl_device_info param, std::vector<T>& value)
  250.     {
  251.         cl_int res;
  252.         size_t size;
  253.  
  254.         res = clGetDeviceInfo(id, param, 0, 0, &size);
  255.         if (CL_SUCCESS != res)
  256.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  257.  
  258.         if (0 == size)
  259.             return CL_SUCCESS;
  260.  
  261.         value.resize(size / sizeof(T));
  262.  
  263.         res = clGetDeviceInfo(id, param, size, &value[0], 0);
  264.         if (CL_SUCCESS != res)
  265.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  266.  
  267.         return CL_SUCCESS;
  268.     }
  269.  
  270.     cl_int query_param(cl_device_id id, cl_device_info param, std::string& value)
  271.     {
  272.         cl_int res;
  273.         size_t size;
  274.  
  275.         res = clGetDeviceInfo(id, param, 0, 0, &size);
  276.         if (CL_SUCCESS != res)
  277.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  278.  
  279.         value.resize(size + 1);
  280.  
  281.         res = clGetDeviceInfo(id, param, size, &value[0], 0);
  282.         if (CL_SUCCESS != res)
  283.             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
  284.  
  285.         // just in case, ensure trailing zero for ASCIIZ string
  286.         value[size] = 0;
  287.  
  288.         return CL_SUCCESS;
  289.     }
  290.  
  291. private:
  292.     cl_device_type                            m_type;
  293.     cl_uint                                   m_vendor_id;
  294.     cl_uint                                   m_max_compute_units;
  295.     cl_uint                                   m_max_work_item_dimensions;
  296.     std::vector<size_t>                       m_max_work_item_sizes;
  297.     size_t                                    m_max_work_group_size;
  298.     cl_uint                                   m_preferred_vector_width_char;
  299.     cl_uint                                   m_preferred_vector_width_short;
  300.     cl_uint                                   m_preferred_vector_width_int;
  301.     cl_uint                                   m_preferred_vector_width_long;
  302.     cl_uint                                   m_preferred_vector_width_float;
  303.     cl_uint                                   m_preferred_vector_width_double;
  304. #if defined(CL_VERSION_1_1)
  305.     cl_uint                                   m_preferred_vector_width_half;
  306.     cl_uint                                   m_native_vector_width_char;
  307.     cl_uint                                   m_native_vector_width_short;
  308.     cl_uint                                   m_native_vector_width_int;
  309.     cl_uint                                   m_native_vector_width_long;
  310.     cl_uint                                   m_native_vector_width_float;
  311.     cl_uint                                   m_native_vector_width_double;
  312.     cl_uint                                   m_native_vector_width_half;
  313. #endif
  314.     cl_uint                                   m_max_clock_frequency;
  315.     cl_uint                                   m_address_bits;
  316.     cl_ulong                                  m_max_mem_alloc_size;
  317.     cl_bool                                   m_image_support;
  318.     cl_uint                                   m_max_read_image_args;
  319.     cl_uint                                   m_max_write_image_args;
  320. #if defined(CL_VERSION_2_0)
  321.     cl_uint                                   m_max_read_write_image_args;
  322. #endif
  323.     size_t                                    m_image2d_max_width;
  324.     size_t                                    m_image2d_max_height;
  325.     size_t                                    m_image3d_max_width;
  326.     size_t                                    m_image3d_max_height;
  327.     size_t                                    m_image3d_max_depth;
  328. #if defined(CL_VERSION_1_2)
  329.     size_t                                    m_image_max_buffer_size;
  330.     size_t                                    m_image_max_array_size;
  331. #endif
  332.     cl_uint                                   m_max_samplers;
  333. #if defined(CL_VERSION_1_2)
  334.     cl_uint                                   m_image_pitch_alignment;
  335.     cl_uint                                   m_image_base_address_alignment;
  336. #endif
  337. #if defined(CL_VERSION_2_0)
  338.     cl_uint                                   m_max_pipe_args;
  339.     cl_uint                                   m_pipe_max_active_reservations;
  340.     cl_uint                                   m_pipe_max_packet_size;
  341. #endif
  342.     size_t                                    m_max_parameter_size;
  343.     cl_uint                                   m_mem_base_addr_align;
  344.     cl_device_fp_config                       m_single_fp_config;
  345. #if defined(CL_VERSION_1_2)
  346.     cl_device_fp_config                       m_double_fp_config;
  347. #endif
  348.     cl_device_mem_cache_type                  m_global_mem_cache_type;
  349.     cl_uint                                   m_global_mem_cacheline_size;
  350.     cl_ulong                                  m_global_mem_cache_size;
  351.     cl_ulong                                  m_global_mem_size;
  352.     cl_ulong                                  m_max_constant_buffer_size;
  353.     cl_uint                                   m_max_constant_args;
  354. #if defined(CL_VERSION_2_0)
  355.     size_t                                    m_max_global_variable_size;
  356.     size_t                                    m_global_variable_preferred_total_size;
  357. #endif
  358.     cl_device_local_mem_type                  m_local_mem_type;
  359.     cl_ulong                                  m_local_mem_size;
  360.     cl_bool                                   m_error_correction_support;
  361. #if defined(CL_VERSION_1_1)
  362.     cl_bool                                   m_host_unified_memory;
  363. #endif
  364.     size_t                                    m_profiling_timer_resolution;
  365.     cl_bool                                   m_endian_little;
  366.     cl_bool                                   m_available;
  367.     cl_bool                                   m_compiler_available;
  368. #if defined(CL_VERSION_1_2)
  369.     cl_bool                                   m_linker_available;
  370. #endif
  371.     cl_device_exec_capabilities               m_execution_capabilities;
  372.     cl_command_queue_properties               m_queue_properties;
  373. #if defined(CL_VERSION_2_0)
  374.     cl_command_queue_properties               m_queue_on_host_properties;
  375.     cl_command_queue_properties               m_queue_on_device_properties;
  376.     cl_uint                                   m_queue_on_device_preferred_size;
  377.     cl_uint                                   m_queue_on_device_max_size;
  378.     cl_uint                                   m_max_on_device_queues;
  379.     cl_uint                                   m_max_on_device_events;
  380. #endif
  381. #if defined(CL_VERSION_1_2)
  382.     std::string                               m_built_in_kernels;
  383. #endif
  384.     cl_platform_id                            m_platform;
  385.     std::string                               m_name;
  386.     std::string                               m_vendor;
  387.     std::string                               m_driver_version;
  388.     std::string                               m_profile;
  389.     std::string                               m_version;
  390. #if defined(CL_VERSION_1_1)
  391.     std::string                               m_opencl_c_version;
  392. #endif
  393.     std::string                               m_extensions;
  394. #if defined(CL_VERSION_1_2)
  395.     size_t                                    m_printf_buffer_size;
  396.     cl_bool                                   m_preferred_interop_user_sync;
  397.     cl_device_id                              m_parent_device;
  398.     cl_uint                                   m_partition_max_sub_devices;
  399.     std::vector<cl_device_partition_property> m_partition_properties;
  400.     cl_device_affinity_domain                 m_partition_affinity_domain;
  401.     std::vector<cl_device_partition_property> m_partition_type;
  402.     cl_uint                                   m_reference_count;
  403. #endif
  404. };
  405.  
  406. } // namespace opencl
  407.  
  408.  
  409. class App
  410. {
  411. public:
  412.     App(CommandLineParser& cmd);
  413.     ~App();
  414.  
  415.     int initOpenCL();
  416.     int initVideoSource();
  417.  
  418.     int process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* cl_buffer);
  419.     int process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u);
  420.     int process_cl_image_with_opencv(cl_mem image, cv::UMat& u);
  421.  
  422.     int run();
  423.  
  424.     bool isRunning() { return m_running; }
  425.     bool doProcess() { return m_process; }
  426.     bool useBuffer() { return m_use_buffer; }
  427.  
  428.     void setRunning(bool running)      { m_running = running; }
  429.     void setDoProcess(bool process)    { m_process = process; }
  430.     void setUseBuffer(bool use_buffer) { m_use_buffer = use_buffer; }
  431.  
  432. protected:
  433.     bool nextFrame(cv::Mat& frame) { return m_cap.read(frame); }
  434.     void handleKey(char key);
  435.     void timerStart();
  436.     void timerEnd();
  437.     std::string timeStr() const;
  438.     std::string message() const;
  439.  
  440. private:
  441.     bool                        m_running;
  442.     bool                        m_process;
  443.     bool                        m_use_buffer;
  444.  
  445.     int64                       m_t0;
  446.     int64                       m_t1;
  447.     float                       m_time;
  448.     float                       m_frequency;
  449.  
  450.     string                      m_file_name;
  451.     int                         m_camera_id;
  452.     cv::VideoCapture            m_cap;
  453.     cv::Mat                     m_frame;
  454.     cv::Mat                     m_frameGray;
  455.  
  456.     opencl::PlatformInfo        m_platformInfo;
  457.     opencl::DeviceInfo          m_deviceInfo;
  458.     std::vector<cl_platform_id> m_platform_ids;
  459.     cl_context                  m_context;
  460.     cl_device_id                m_device_id;
  461.     cl_command_queue            m_queue;
  462.     cl_program                  m_program;
  463.     cl_kernel                   m_kernelBuf;
  464.     cl_kernel                   m_kernelImg;
  465.     cl_mem                      m_img_src; // used as src in case processing of cl image
  466.     cl_mem                      m_mem_obj;
  467.     cl_event                    m_event;
  468. };
  469.  
  470.  
  471. App::App(CommandLineParser& cmd)
  472. {
  473.     cout << "\nPress ESC to exit\n" << endl;
  474.     cout << "\n      'p' to toggle ON/OFF processing\n" << endl;
  475.     cout << "\n       SPACE to switch between OpenCL buffer/image\n" << endl;
  476.  
  477.     m_camera_id  = cmd.get<int>("camera");
  478.     m_file_name  = cmd.get<string>("video");
  479.  
  480.     m_running    = false;
  481.     m_process    = false;
  482.     m_use_buffer = false;
  483.  
  484.     m_t0         = 0;
  485.     m_t1         = 0;
  486.     m_time       = 0.0;
  487.     m_frequency  = (float)cv::getTickFrequency();
  488.  
  489.     m_context    = 0;
  490.     m_device_id  = 0;
  491.     m_queue      = 0;
  492.     m_program    = 0;
  493.     m_kernelBuf  = 0;
  494.     m_kernelImg  = 0;
  495.     m_img_src    = 0;
  496.     m_mem_obj    = 0;
  497.     m_event      = 0;
  498. } // ctor
  499.  
  500.  
  501. App::~App()
  502. {
  503.     if (m_queue)
  504.     {
  505.         clFinish(m_queue);
  506.         clReleaseCommandQueue(m_queue);
  507.         m_queue = 0;
  508.     }
  509.  
  510.     if (m_program)
  511.     {
  512.         clReleaseProgram(m_program);
  513.         m_program = 0;
  514.     }
  515.  
  516.     if (m_img_src)
  517.     {
  518.         clReleaseMemObject(m_img_src);
  519.         m_img_src = 0;
  520.     }
  521.  
  522.     if (m_mem_obj)
  523.     {
  524.         clReleaseMemObject(m_mem_obj);
  525.         m_mem_obj = 0;
  526.     }
  527.  
  528.     if (m_event)
  529.     {
  530.         clReleaseEvent(m_event);
  531.     }
  532.  
  533.     if (m_kernelBuf)
  534.     {
  535.         clReleaseKernel(m_kernelBuf);
  536.         m_kernelBuf = 0;
  537.     }
  538.  
  539.     if (m_kernelImg)
  540.     {
  541.         clReleaseKernel(m_kernelImg);
  542.         m_kernelImg = 0;
  543.     }
  544.  
  545.     if (m_device_id)
  546.     {
  547.         clReleaseDevice(m_device_id);
  548.         m_device_id = 0;
  549.     }
  550.  
  551.     if (m_context)
  552.     {
  553.         clReleaseContext(m_context);
  554.         m_context = 0;
  555.     }
  556. } // dtor
  557.  
  558.  
  559. int App::initOpenCL()
  560. {
  561.     cl_int res = CL_SUCCESS;
  562.     cl_uint num_entries = 0;
  563.  
  564.     res = clGetPlatformIDs(0, 0, &num_entries);
  565.     if (CL_SUCCESS != res)
  566.         return -1;
  567.  
  568.     m_platform_ids.resize(num_entries);
  569.  
  570.     res = clGetPlatformIDs(num_entries, &m_platform_ids[0], 0);
  571.     if (CL_SUCCESS != res)
  572.         return -1;
  573.  
  574.     unsigned int i;
  575.  
  576.     // create context from first platform with GPU device
  577.     for (i = 0; i < m_platform_ids.size(); i++)
  578.     {
  579.         cl_context_properties props[] =
  580.         {
  581.             CL_CONTEXT_PLATFORM,
  582.             (cl_context_properties)(m_platform_ids[i]),
  583.             0
  584.         };
  585.  
  586.         m_context = clCreateContextFromType(props, CL_DEVICE_TYPE_GPU, 0, 0, &res);
  587.         if (0 == m_context || CL_SUCCESS != res)
  588.             continue;
  589.  
  590.         res = clGetContextInfo(m_context, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &m_device_id, 0);
  591.         if (CL_SUCCESS != res)
  592.             return -1;
  593.  
  594.         m_queue = clCreateCommandQueue(m_context, m_device_id, 0, &res);
  595.         if (0 == m_queue || CL_SUCCESS != res)
  596.             return -1;
  597.  
  598.         const char* kernelSrc =
  599.             "__kernel "
  600.             "void bitwise_inv_buf_8uC1("
  601.             "    __global unsigned char* pSrcDst,"
  602.             "             int            srcDstStep,"
  603.             "             int            rows,"
  604.             "             int            cols)"
  605.             "{"
  606.             "    int x = get_global_id(0);"
  607.             "    int y = get_global_id(1);"
  608.             "    int idx = mad24(y, srcDstStep, x);"
  609.             "    pSrcDst[idx] = ~pSrcDst[idx];"
  610.             "}"
  611.             "__kernel "
  612.             "void bitwise_inv_img_8uC1("
  613.             "    read_only  image2d_t srcImg,"
  614.             "    write_only image2d_t dstImg)"
  615.             "{"
  616.             "    int x = get_global_id(0);"
  617.             "    int y = get_global_id(1);"
  618.             "    int2 coord = (int2)(x, y);"
  619.             "    uint4 val = read_imageui(srcImg, coord);"
  620.             "    val.x = (~val.x) & 0x000000FF;"
  621.             "    write_imageui(dstImg, coord, val);"
  622.             "}";
  623.         size_t len = strlen(kernelSrc);
  624.         m_program = clCreateProgramWithSource(m_context, 1, &kernelSrc, &len, &res);
  625.         if (0 == m_program || CL_SUCCESS != res)
  626.             return -1;
  627.  
  628.         res = clBuildProgram(m_program, 1, &m_device_id, 0, 0, 0);
  629.         if (CL_SUCCESS != res)
  630.             return -1;
  631.  
  632.         m_kernelBuf = clCreateKernel(m_program, "bitwise_inv_buf_8uC1", &res);
  633.         if (0 == m_kernelBuf || CL_SUCCESS != res)
  634.             return -1;
  635.  
  636.         m_kernelImg = clCreateKernel(m_program, "bitwise_inv_img_8uC1", &res);
  637.         if (0 == m_kernelImg || CL_SUCCESS != res)
  638.             return -1;
  639.  
  640.         m_platformInfo.QueryInfo(m_platform_ids[i]);
  641.         m_deviceInfo.QueryInfo(m_device_id);
  642.  
  643.         // attach OpenCL context to OpenCV
  644.         cv::ocl::attachContext(m_platformInfo.Name(), m_platform_ids[i], m_context, m_device_id);
  645.  
  646.         break;
  647.     }
  648.  
  649.     return m_context != 0 ? CL_SUCCESS : -1;
  650. } // initOpenCL()
  651.  
  652.  
  653. int App::initVideoSource()
  654. {
  655.     try
  656.     {
  657.         if (!m_file_name.empty() && m_camera_id == -1)
  658.         {
  659.             m_cap.open(m_file_name.c_str());
  660.             if (!m_cap.isOpened())
  661.                 throw std::runtime_error(std::string("can't open video file: " + m_file_name));
  662.         }
  663.         else if (m_camera_id != -1)
  664.         {
  665.             m_cap.open(m_camera_id);
  666.             if (!m_cap.isOpened())
  667.             {
  668.                 std::stringstream msg;
  669.                 msg << "can't open camera: " << m_camera_id;
  670.                 throw std::runtime_error(msg.str());
  671.             }
  672.         }
  673.         else
  674.             throw std::runtime_error(std::string("specify video source"));
  675.     }
  676.  
  677.     catch (std::exception e)
  678.     {
  679.         cerr << "ERROR: " << e.what() << std::endl;
  680.         return -1;
  681.     }
  682.  
  683.     return 0;
  684. } // initVideoSource()
  685.  
  686.  
  687. // this function is an example of "typical" OpenCL processing pipeline
  688. // It creates OpenCL buffer or image, depending on use_buffer flag,
  689. // from input media frame and process these data
  690. // (inverts each pixel value in half of frame) with OpenCL kernel
  691. int App::process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* mem_obj)
  692. {
  693.     cl_int res = CL_SUCCESS;
  694.  
  695.     CV_Assert(mem_obj);
  696.  
  697.     cl_kernel kernel = 0;
  698.     cl_mem mem = mem_obj[0];
  699.  
  700.     if (0 == mem || 0 == m_img_src)
  701.     {
  702.         // allocate/delete cl memory objects every frame for the simplicity.
  703.         // in real applicaton more efficient pipeline can be built.
  704.  
  705.         if (use_buffer)
  706.         {
  707.             cl_mem_flags flags = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR;
  708.  
  709.             mem = clCreateBuffer(m_context, flags, frame.total(), frame.ptr(), &res);
  710.             if (0 == mem || CL_SUCCESS != res)
  711.                 return -1;
  712.  
  713.             res = clSetKernelArg(m_kernelBuf, 0, sizeof(cl_mem), &mem);
  714.             if (CL_SUCCESS != res)
  715.                 return -1;
  716.  
  717.             res = clSetKernelArg(m_kernelBuf, 1, sizeof(int), &frame.step[0]);
  718.             if (CL_SUCCESS != res)
  719.                 return -1;
  720.  
  721.             res = clSetKernelArg(m_kernelBuf, 2, sizeof(int), &frame.rows);
  722.             if (CL_SUCCESS != res)
  723.                 return -1;
  724.  
  725.             int cols2 = frame.cols / 2;
  726.             res = clSetKernelArg(m_kernelBuf, 3, sizeof(int), &cols2);
  727.             if (CL_SUCCESS != res)
  728.                 return -1;
  729.  
  730.             kernel = m_kernelBuf;
  731.         }
  732.         else
  733.         {
  734.             cl_mem_flags flags_src = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR;
  735.  
  736.             cl_image_format fmt;
  737.             fmt.image_channel_order     = CL_R;
  738.             fmt.image_channel_data_type = CL_UNSIGNED_INT8;
  739.  
  740.             cl_image_desc desc_src;
  741.             desc_src.image_type        = CL_MEM_OBJECT_IMAGE2D;
  742.             desc_src.image_width       = frame.cols;
  743.             desc_src.image_height      = frame.rows;
  744.             desc_src.image_depth       = 0;
  745.             desc_src.image_array_size  = 0;
  746.             desc_src.image_row_pitch   = frame.step[0];
  747.             desc_src.image_slice_pitch = 0;
  748.             desc_src.num_mip_levels    = 0;
  749.             desc_src.num_samples       = 0;
  750.             desc_src.buffer            = 0;
  751.             m_img_src = clCreateImage(m_context, flags_src, &fmt, &desc_src, frame.ptr(), &res);
  752.             if (0 == m_img_src || CL_SUCCESS != res)
  753.                 return -1;
  754.  
  755.             cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR;
  756.  
  757.             cl_image_desc desc_dst;
  758.             desc_dst.image_type        = CL_MEM_OBJECT_IMAGE2D;
  759.             desc_dst.image_width       = frame.cols;
  760.             desc_dst.image_height      = frame.rows;
  761.             desc_dst.image_depth       = 0;
  762.             desc_dst.image_array_size  = 0;
  763.             desc_dst.image_row_pitch   = 0;
  764.             desc_dst.image_slice_pitch = 0;
  765.             desc_dst.num_mip_levels    = 0;
  766.             desc_dst.num_samples       = 0;
  767.             desc_dst.buffer            = 0;
  768.             mem = clCreateImage(m_context, flags_dst, &fmt, &desc_dst, 0, &res);
  769.             if (0 == mem || CL_SUCCESS != res)
  770.                 return -1;
  771.  
  772.             size_t origin[] = { 0, 0, 0 };
  773.             size_t region[] = { (size_t)frame.cols, (size_t)frame.rows, 1 };
  774.             res = clEnqueueCopyImage(m_queue, m_img_src, mem, origin, origin, region, 0, 0, &m_event);
  775.             if (CL_SUCCESS != res)
  776.                 return -1;
  777.  
  778.             res = clWaitForEvents(1, &m_event);
  779.             if (CL_SUCCESS != res)
  780.                 return -1;
  781.  
  782.             res = clSetKernelArg(m_kernelImg, 0, sizeof(cl_mem), &m_img_src);
  783.             if (CL_SUCCESS != res)
  784.                 return -1;
  785.  
  786.             res = clSetKernelArg(m_kernelImg, 1, sizeof(cl_mem), &mem);
  787.             if (CL_SUCCESS != res)
  788.                 return -1;
  789.  
  790.             kernel = m_kernelImg;
  791.         }
  792.     }
  793.  
  794.     m_event = clCreateUserEvent(m_context, &res);
  795.     if (0 == m_event || CL_SUCCESS != res)
  796.         return -1;
  797.  
  798.     // process left half of frame in OpenCL
  799.     size_t size[] = { (size_t)frame.cols / 2, (size_t)frame.rows };
  800.     res = clEnqueueNDRangeKernel(m_queue, kernel, 2, 0, size, 0, 0, 0, &m_event);
  801.     if (CL_SUCCESS != res)
  802.         return -1;
  803.  
  804.     res = clWaitForEvents(1, &m_event);
  805.     if (CL_SUCCESS != res)
  806.         return - 1;
  807.  
  808.     mem_obj[0] = mem;
  809.  
  810.     return  0;
  811. }
  812.  
  813.  
  814. // this function is an example of interoperability between OpenCL buffer
  815. // and OpenCV UMat objects. It converts (without copying data) OpenCL buffer
  816. // to OpenCV UMat and then do blur on these data
  817. int App::process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u)
  818. {
  819.     cv::ocl::convertFromBuffer(buffer, step, rows, cols, type, u);
  820.  
  821.     // process right half of frame in OpenCV
  822.     cv::Point pt(u.cols / 2, 0);
  823.     cv::Size  sz(u.cols / 2, u.rows);
  824.     cv::Rect roi(pt, sz);
  825.     cv::UMat uroi(u, roi);
  826.     cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3));
  827.  
  828.     if (buffer)
  829.         clReleaseMemObject(buffer);
  830.     m_mem_obj = 0;
  831.  
  832.     return 0;
  833. }
  834.  
  835.  
  836. // this function is an example of interoperability between OpenCL image
  837. // and OpenCV UMat objects. It converts OpenCL image
  838. // to OpenCV UMat and then do blur on these data
  839. int App::process_cl_image_with_opencv(cl_mem image, cv::UMat& u)
  840. {
  841.     cv::ocl::convertFromImage(image, u);
  842.  
  843.     // process right half of frame in OpenCV
  844.     cv::Point pt(u.cols / 2, 0);
  845.     cv::Size  sz(u.cols / 2, u.rows);
  846.     cv::Rect roi(pt, sz);
  847.     cv::UMat uroi(u, roi);
  848.     cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3));
  849.  
  850.     if (image)
  851.         clReleaseMemObject(image);
  852.     m_mem_obj = 0;
  853.  
  854.     if (m_img_src)
  855.         clReleaseMemObject(m_img_src);
  856.     m_img_src = 0;
  857.  
  858.     return 0;
  859. }
  860.  
  861.  
  862. int App::run()
  863. {
  864.     if (0 != initOpenCL())
  865.         return -1;
  866.  
  867.     if (0 != initVideoSource())
  868.         return -1;
  869.  
  870.     Mat img_to_show;
  871.  
  872.     // set running state until ESC pressed
  873.     setRunning(true);
  874.     // set process flag to show some data processing
  875.     // can be toggled on/off by 'p' button
  876.     setDoProcess(true);
  877.     // set use buffer flag,
  878.     // when it is set to true, will demo interop opencl buffer and cv::Umat,
  879.     // otherwise demo interop opencl image and cv::UMat
  880.     // can be switched on/of by SPACE button
  881.     setUseBuffer(true);
  882.  
  883.     // Iterate over all frames
  884.     while (isRunning() && nextFrame(m_frame))
  885.     {
  886.         cv::cvtColor(m_frame, m_frameGray, COLOR_BGR2GRAY);
  887.  
  888.         UMat uframe;
  889.  
  890.         // work
  891.         timerStart();
  892.  
  893.         if (doProcess())
  894.         {
  895.             process_frame_with_open_cl(m_frameGray, useBuffer(), &m_mem_obj);
  896.  
  897.             if (useBuffer())
  898.                 process_cl_buffer_with_opencv(
  899.                     m_mem_obj, m_frameGray.step[0], m_frameGray.rows, m_frameGray.cols, m_frameGray.type(), uframe);
  900.             else
  901.                 process_cl_image_with_opencv(m_mem_obj, uframe);
  902.         }
  903.         else
  904.         {
  905.             m_frameGray.copyTo(uframe);
  906.         }
  907.  
  908.         timerEnd();
  909.  
  910.         uframe.copyTo(img_to_show);
  911.  
  912.         putText(img_to_show, "Version : " + m_platformInfo.Version(), Point(5, 30), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
  913.         putText(img_to_show, "Name : " + m_platformInfo.Name(), Point(5, 60), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
  914.         putText(img_to_show, "Device : " + m_deviceInfo.Name(), Point(5, 90), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
  915.         cv::String memtype = useBuffer() ? "buffer" : "image";
  916.         putText(img_to_show, "interop with OpenCL " + memtype, Point(5, 120), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
  917.         putText(img_to_show, "Time : " + timeStr() + " msec", Point(5, 150), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
  918.  
  919.         imshow("opencl_interop", img_to_show);
  920.  
  921.         handleKey((char)waitKey(3));
  922.     }
  923.  
  924.     return 0;
  925. }
  926.  
  927.  
  928. void App::handleKey(char key)
  929. {
  930.     switch (key)
  931.     {
  932.     case 27:
  933.         setRunning(false);
  934.         break;
  935.  
  936.     case ' ':
  937.         setUseBuffer(!useBuffer());
  938.         break;
  939.  
  940.     case 'p':
  941.     case 'P':
  942.         setDoProcess( !doProcess() );
  943.         break;
  944.  
  945.     default:
  946.         break;
  947.     }
  948. }
  949.  
  950.  
  951. inline void App::timerStart()
  952. {
  953.     m_t0 = getTickCount();
  954. }
  955.  
  956.  
  957. inline void App::timerEnd()
  958. {
  959.     m_t1 = getTickCount();
  960.     int64 delta = m_t1 - m_t0;
  961.     m_time = (delta / m_frequency) * 1000; // units msec
  962. }
  963.  
  964.  
  965. inline string App::timeStr() const
  966. {
  967.     stringstream ss;
  968.     ss << std::fixed << std::setprecision(1) << m_time;
  969.     return ss.str();
  970. }
  971.  
  972.  
  973. int main(int argc, char** argv)
  974. {
  975.     const char* keys =
  976.         "{ help h ?    |          | print help message }"
  977.         "{ camera c    | -1       | use camera as input }"
  978.         "{ video  v    |          | use video as input }";
  979.  
  980.     CommandLineParser cmd(argc, argv, keys);
  981.     if (cmd.has("help"))
  982.     {
  983.         cmd.printMessage();
  984.         return EXIT_SUCCESS;
  985.     }
  986.  
  987.     App app(cmd);
  988.  
  989.     try
  990.     {
  991.         app.run();
  992.     }
  993.  
  994.     catch (const cv::Exception& e)
  995.     {
  996.         cout << "error: " << e.what() << endl;
  997.         return 1;
  998.     }
  999.  
  1000.     catch (const std::exception& e)
  1001.     {
  1002.         cout << "error: " << e.what() << endl;
  1003.         return 1;
  1004.     }
  1005.  
  1006.     catch (...)
  1007.     {
  1008.         cout << "unknown exception" << endl;
  1009.         return 1;
  1010.     }
  1011.  
  1012.     return EXIT_SUCCESS;
  1013. } // main()
4.2.2 Makefile
  1. TARGET = main
  2. CXX = g++
  3. CFLAGS += -I/usr/include -I/usr/local/include/opencv -I/usr/local/include/opencv2 -L/usr/lib  -L/usr/local/lib -L/lib -std=c++98
  4.  
  5. CFLAGS +=  -lopencv_core -lopencv_objdetect -lopencv_highgui -lopencv_videoio -lopencv_imgcodecs -lopencv_imgproc -lOpenCL -lpthread -lrt
  6.  
  7. all:
  8.     @$(CXX)  $(TARGET).cpp -o $(TARGET) $(CFLAGS)
  9. clean:
  10.     rm -rf  $(TARGET)
4.2.3 编译运行
  1. root@NanoPC-T6:/opt/opencl-project/opencv-ocl# make 
  2. root@NanoPC-T6:/opt/opencl-project/opencv-ocl# ./main -c

如下图所示：

参考文章

[1] RK3588实战：调用npu加速，yolov5识别图像、ffmpeg发送到rtmp服务器

[2] 嵌入式AI应用开发实战指南—基于LubanCat-RK系列板卡

[3] RK3588边缘计算

[4] OpenCL学习笔记(四)手动编译开发库(ubuntu+gcc+rk3588)

[5] 如何在RK3399中使用opencl并安装QT开发

[6] Arm Mali GPU OpenCL Developer Guide

[7] 什么是OpenCL

[8] 高性能计算

[9] OpenCL练习(一)：使用OpenCL+OpenCV进行RGB转灰度图

[10] https://opencv.org/opencl

[11] https://github.com/opencv/opencv/wiki/OpenCL-optimizations

[12] OpenCV OpenCL support

[13] 一、Opencv-OCL编程基础

声明：本文内容由网友自发贡献，不代表【wpsshop博客】立场，版权归原作者所有，本站不承担相应法律责任。如您发现有侵权的内容，请联系我们。转载请注明出处：https://www.wpsshop.cn/w/我家小花儿/article/detail/535448
推荐阅读
相关标签

Copyright © 2003-2013 www.wpsshop.cn 版权所有，并保留所有权利。

  

闽ICP备14008679号