当前位置：首页 > news >正文

【ELF2学习板】OpenCL程序测试

news 来源：原创 2025/5/13 22:44:26

引言

OpenCL简介

主要特点

编程模型

应用场景

测试程序

代码说明

构建编译环境

头文件

库文件

程序编译

测试结果

结语

引言

ELF2开发板采用的是RK3588处理器，它是瑞芯微推出的一款高性能 SoC。RK3588 集成了 ARM Mali-G610 MP4 GPU，这款 GPU 具备强大的并行计算能力。要充分发挥GPU的计算能力，可以使用 OpenCL库。OpenCL（Open Computing Language）是一个开放的、跨平台的并行编程标准，用于编写可在多种计算设备（如 CPU、GPU、FPGA 等）上运行的并行程序。瑞芯微为 RK3588 提供专门的 GPU 驱动，支持 OpenCL 标准。

OpenCL简介

主要特点

跨平台性：OpenCL 可以在不同厂商的硬件设备上运行，包括英特尔、AMD、英伟达等公司生产的 CPU 和 GPU，以及其他类型的加速器，如 FPGA。这使得开发者能够编写一次代码，然后在多种设备上部署和运行。
并行计算能力：它允许开发者充分利用计算设备中的大量计算核心，通过并行计算来加速程序的执行。对于大规模数据处理、科学计算、图形处理等领域，OpenCL 可以显著提高程序的性能。
灵活性：OpenCL 提供了丰富的编程模型，开发者可以根据具体的应用场景和硬件特性，灵活地控制并行计算的粒度和方式。可以通过调整工作组大小、全局工作大小等参数来优化程序性能。

编程模型

OpenCL 的编程模型主要涉及以下几个关键概念：

平台（Platform）：代表一个 OpenCL 实现的集合，通常对应于一个硬件厂商的软件栈。例如，英特尔的 OpenCL 平台、英伟达的 OpenCL 平台等。
设备（Device）：指具体的计算设备，如 CPU、GPU 或 FPGA。一个平台可以包含多个设备。
上下文（Context）：是 OpenCL 运行时的一个环境，它管理着所有的 OpenCL 对象，如内存对象、程序对象、内核对象等。
命令队列（Command Queue）：用于向设备发送命令，如内核执行命令、内存读写命令等。命令队列可以是有序的或无序的，开发者可以根据需要选择不同的队列类型。
内核（Kernel）：是 OpenCL 中并行执行的基本单元，它是一段用 OpenCL C 语言编写的函数，在设备上并行执行。
内存对象（Memory Object）：用于在主机和设备之间传输数据，包括缓冲区（Buffer）和图像（Image）两种类型。

应用场景

科学计算：在气象预报、分子动力学模拟、金融建模等领域，需要处理大量的数据和复杂的计算任务。OpenCL 可以利用 GPU 等设备的并行计算能力，加速这些计算任务的执行。
图形处理：虽然 OpenGL 和 DirectX 是专门用于图形渲染的 API，但 OpenCL 可以用于图形处理中的一些辅助计算任务，如物理模拟、图像处理等。
数据挖掘和机器学习：在处理大规模数据集时，OpenCL 可以加速数据挖掘算法和机器学习模型的训练过程。例如，加速神经网络的前向传播和反向传播计算。

测试程序

以下是用豆包生成的一个简单的 OpenCL 测试程序，该程序会对两个数组进行加法运算。此程序既可以验证功能的正确性，也能通过计时来体现性能提升的能力。

#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>
#include <time.h>#define ARRAY_SIZE 1048576 // 增大数据规模// OpenCL 内核代码
const char *kernelSource ="__kernel void vector_add(__global const float *a, __global const float *b, __global float *c) {\n""    int i = get_global_id(0);\n""    c[i] = a[i] + b[i];\n""}\n";// CPU 上的数组加法函数
void cpu_vector_add(float *a, float *b, float *c, int size) {for (int i = 0; i < size; i++) {c[i] = a[i] + b[i];}
}int main() {cl_platform_id platform_id = NULL;cl_device_id device_id = NULL;cl_uint ret_num_devices;cl_uint ret_num_platforms;cl_int ret;// 获取平台信息ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);if (ret != CL_SUCCESS) {printf("Error: Failed to get platform IDs!\n");return -1;}// 获取设备信息ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);if (ret == CL_DEVICE_NOT_FOUND) {ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_CPU, 1, &device_id, &ret_num_devices);}if (ret != CL_SUCCESS) {printf("Error: Failed to get device IDs!\n");return -1;}// 创建上下文cl_context context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);if (ret != CL_SUCCESS) {printf("Error: Failed to create context!\n");return -1;}// 创建命令队列cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret);if (ret != CL_SUCCESS) {printf("Error: Failed to create command queue!\n");return -1;}// 创建内存对象float *a = (float *)malloc(ARRAY_SIZE * sizeof(float));float *b = (float *)malloc(ARRAY_SIZE * sizeof(float));float *c_gpu = (float *)malloc(ARRAY_SIZE * sizeof(float));float *c_cpu = (float *)malloc(ARRAY_SIZE * sizeof(float));for (int i = 0; i < ARRAY_SIZE; i++) {a[i] = (float)i;b[i] = (float)(i * 2);}cl_mem a_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, ARRAY_SIZE * sizeof(float), NULL, &ret);cl_mem b_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, ARRAY_SIZE * sizeof(float), NULL, &ret);cl_mem c_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, ARRAY_SIZE * sizeof(float), NULL, &ret);if (ret != CL_SUCCESS) {printf("Error: Failed to create memory objects!\n");return -1;}// 将数据写入内存对象ret = clEnqueueWriteBuffer(command_queue, a_mem_obj, CL_TRUE, 0, ARRAY_SIZE * sizeof(float), a, 0, NULL, NULL);ret |= clEnqueueWriteBuffer(command_queue, b_mem_obj, CL_TRUE, 0, ARRAY_SIZE * sizeof(float), b, 0, NULL, NULL);if (ret != CL_SUCCESS) {printf("Error: Failed to write data to memory objects!\n");return -1;}// 创建程序对象cl_program program = clCreateProgramWithSource(context, 1, (const char **)&kernelSource, NULL, &ret);if (ret != CL_SUCCESS) {printf("Error: Failed to create program object!\n");return -1;}// 构建程序ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);if (ret != CL_SUCCESS) {size_t len;char buffer[2048];clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);printf("Error: Failed to build program:\n%s\n", buffer);return -1;}// 创建内核对象cl_kernel kernel = clCreateKernel(program, "vector_add", &ret);if (ret != CL_SUCCESS) {printf("Error: Failed to create kernel object!\n");return -1;}// 设置内核参数ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj);ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_mem_obj);ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_mem_obj);if (ret != CL_SUCCESS) {printf("Error: Failed to set kernel arguments!\n");return -1;}// 测量 GPU 执行时间clock_t start_gpu = clock();// 执行内核size_t global_item_size = ARRAY_SIZE;size_t local_item_size = 256; // 调整本地工作组大小ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);if (ret != CL_SUCCESS) {printf("Error: Failed to enqueue kernel!\n");return -1;}// 读取结果ret = clEnqueueReadBuffer(command_queue, c_mem_obj, CL_TRUE, 0, ARRAY_SIZE * sizeof(float), c_gpu, 0, NULL, NULL);if (ret != CL_SUCCESS) {printf("Error: Failed to read result!\n");return -1;}clock_t end_gpu = clock();double time_gpu = ((double)(end_gpu - start_gpu)) / CLOCKS_PER_SEC;// 测量 CPU 执行时间clock_t start_cpu = clock();cpu_vector_add(a, b, c_cpu, ARRAY_SIZE);clock_t end_cpu = clock();double time_cpu = ((double)(end_cpu - start_cpu)) / CLOCKS_PER_SEC;// 验证结果int correct = 1;for (int i = 0; i < ARRAY_SIZE; i++) {if (c_gpu[i] != c_cpu[i]) {correct = 0;break;}}if (correct) {printf("Results are correct!\n");} else {printf("Results are incorrect!\n");}// 输出性能对比结果printf("CPU execution time: %f seconds\n", time_cpu);printf("GPU execution time: %f seconds\n", time_gpu);if (time_gpu < time_cpu) {printf("GPU is %.2f times faster than CPU.\n", time_cpu / time_gpu);} else {printf("CPU is %.2f times faster than GPU.\n", time_gpu / time_cpu);}// 释放资源ret = clFlush(command_queue);ret = clFinish(command_queue);ret = clReleaseKernel(kernel);ret = clReleaseProgram(program);ret = clReleaseMemObject(a_mem_obj);ret = clReleaseMemObject(b_mem_obj);ret = clReleaseMemObject(c_mem_obj);ret = clReleaseCommandQueue(command_queue);ret = clReleaseContext(context);free(a);free(b);free(c_gpu);free(c_cpu);return 0;
}

代码说明

平台与设备信息获取：借助 clGetPlatformIDs 和 clGetDeviceIDs 来获取可用的 OpenCL 平台与设备。
上下文与命令队列创建：运用 clCreateContext 创建上下文，使用 clCreateCommandQueue 创建命令队列。
内存对象创建：通过 clCreateBuffer 创建内存对象，并且使用 clEnqueueWriteBuffer 把数据写入这些对象。
程序与内核创建：利用 clCreateProgramWithSource 创建程序对象，通过 clBuildProgram 构建程序，使用 clCreateKernel 创建内核对象。
内核参数设置与执行：使用 clSetKernelArg 设置内核参数，通过 clEnqueueNDRangeKernel 执行内核。
结果读取与验证：使用 clEnqueueReadBuffer 读取结果，然后验证计算结果的正确性。
资源释放：释放所有创建的 OpenCL 对象和分配的内存。

要充分发挥 GPU 的性能优势，需要增大数据规模，确保 GPU 处于良好的工作状态。

构建编译环境

如果使用的是Ubuntu桌面板，可以直接在ELF2开发板上进行本地编译。

对于Buildroot版本，正常的ELF2编译环境的构建应该基于厂商提供的SDK或者Linux系统源码。不过我偷了个懒，直接从OpenCL的官网下载了头文件，从开发板上拷贝出来所需要的库文件，搭建了一个简易的编译环境。

头文件

从GitHub - KhronosGroup/OpenCL-Headers at v2020.03.13下载需要的头文件，并保存在OpenCL-Headers-2020.03.13目录下。

库文件

把开发板上的/usr/lib/libmali.so.1.9.0文件拷贝到PC上，改名为libmali.so，放置在lib目录下。RK3588的OpenCL库文件就是这个libmali。

程序编译

建立一个src目录，将程序保存在该目录下，然后用如下命令进行编译：

aarch64-linux-gnu-gcc opencl_array_addition.c -o opencl_array_addition -lmali -I../OpenCL-Headers-2020.03.13 -L../lib/  -Wl,--allow-shlib-undefined

这里使用-Wl,--allow-shlib-undefined选项，忽略所有找不到函数的链接错误。因为Mali库还需要其他一些库的支持。

测试结果

程序运行的结果如下：

root@elf2-buildroot:~# ./opencl_array_addition
arm_release_ver: g13p0-01eac0, rk_so_ver: 10
Results are correct!
CPU execution time: 0.022247 seconds
GPU execution time: 0.006805 seconds
GPU is 3.27 times faster than CPU.

从这个结果可以看出，在数据量很大的情况下，GPU还是发挥出来其自身的并行优势的。