RK3588 GPU Mali-G610 MP4 OpenCL开发

一、环境准备

确认RK3588平台操作系统及GPU驱动正常安装，支持Mali-G610 MP4。
安装OpenCL运行时环境（支持OpenCL 2.0），可通过命令clinfo验证。
确保开发环境支持OpenCL编译与调试，如GCC、LLVM或厂商SDK。

sudo apt install clinfo
clinfo

二、OpenCL开发基本流程

查询平台和设备信息：
```
clGetPlatformIDs();
clGetDeviceIDs();
```

创建OpenCL上下文和命令队列：

clCreateContext();
clCreateCommandQueueWithProperties();

编译OpenCL内核程序：

clCreateProgramWithSource();
clBuildProgram();
clCreateKernel();

创建内存缓冲区，进行数据传输：
```
clCreateBuffer();
clEnqueueWriteBuffer();
```
设置合理的全局和局部工作项大小并执行内核：
```
clEnqueueNDRangeKernel();
```

三、工作组（Work-Group）设置要点

核心概念定义

NDRange（N维范围）：定义了全局计算域的大小（global size）。
Global Size（全局工作项数）：所有工作项的总数，通常等于图像的尺寸，例如 width × height。
Local Size（局部工作组尺寸）：每个工作组中工作项的数量（可选参数）。必须整除 global size。
Work-item（工作项）：每个工作项可以处理图像中的一个像素或一个数据单元。
Work-group（工作组）：OpenCL 将工作项组织成工作组，便于共享本地内存和同步。

设置要点

工作组大小影响GPU资源调度和执行效率，注意：

Local Work Size（局部工作组大小）通常设置为2的幂次方，如16×16，最大不要超过设备允许的CL_DEVICE_MAX_WORK_GROUP_SIZE，一般为256或512。
一个Shader Core（Mali-G610 MP4有4个核心）会执行一个工作组，因此局部工作组大小应合理控制在单核心可承载范围内，避免过大导致资源争用或执行效率下降。
全局工作项（Global Work Size）可远大于核心数，用于覆盖整个数据规模。

示例设置：

size_t global_work_size[2] = {1920, 1280};
size_t local_work_size[2] = {16, 16};

四、Mali-G610 MP4 GPU内存架构简介

层级	名称 / 特征	访问速度	容量范围	说明
寄存器	每个 Shader Core 独立	🚀最快（1cycle）	十几个~上百个 / Core	用于 warp/thread 间变量寄存
L0 Cache	特定类型如纹理缓存（可选）	很快	通常16~64KB	Shader Core 访问纹理时命中
L1 Cache	各 Shader Core 私有	快	通常32~64KB	仅该核心访问，有助于局部性
L2 Cache	所有 Shader Core 共享	中等	256KB ~ 1MB（G610）	GPU 内部各核心协同通信关键
系统主存	SoC DDR / LPDDR（外部 DRAM）	最慢（~100ns）	几 GB ~ 32GB	与 CPU/NPU/ISP 共享

OpenCL内存模型对应关系

OpenCL内存类型	硬件映射	特点
`__private`	寄存器	线程私有，访问速度最快，数量有限
`__local`	L1 Cache / 片上存储	同一工作组共享，延迟低，容量有限
`__global`	系统主存（通过L2缓存）	共享系统DRAM，延迟高，带宽受限
`__constant`	L2 Cache专用区域	只读缓存区域，适合存储不变数据
`image2d_t`等	L0纹理缓存	专用缓存优化纹理采样访问

五、性能优化建议

尽量减少对__global内存的访问次数，优先使用__local内存做缓存。
避免内核中复杂分支，采用数学替代分支操作。
使用向量数据类型（如float4）提升内存访问效率。
合理设置局部工作组大小，避免超过硬件支持的最大工作组大小。
利用Mali-G610的L1和L2缓存机制，优化内存访问模式。

六、示例代码

简单的图像像素反转内核

__kernel void invert(__global uchar *image, const int size) {
    int gid = get_global_id(0);
    if (gid < size) {
        image[gid] = 255 - image[gid];
    }
}

七、查询最大工作组大小示例代码

size_t maxWorkGroupSize = 0;
clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &maxWorkGroupSize, nullptr);
printf("Max work group size: %zu\n", maxWorkGroupSize);

RK3588 GPU Mali-G610 MP4 OpenCL开发指导

一、环境准备

二、OpenCL开发基本流程

三、工作组（Work-Group）设置要点

核心概念定义

设置要点

四、Mali-G610 MP4 GPU内存架构简介

OpenCL内存模型对应关系

五、性能优化建议

六、示例代码

简单的图像像素反转内核

七、查询最大工作组大小示例代码

八、常用工具

总结