第4章 Exynos GPU架构与编程

第4章：Exynos GPU架构与编程：Mali GPU管线、OpenCL基础、在Exynos上运行第一个GPU Kernel

好，咱们今天来聊聊Exynos里最带劲的部分——GPU。说实话，很多做嵌入式开发的朋友，一听到GPU就觉得那是做图形渲染的，跟我搞算法的没啥关系。其实不然。你想想看，现在的手机芯片里，GPU占的面积比CPU还大，算力更是碾压CPU。放着这么好的资源不用，是不是有点浪费？

我个人习惯把Exynos的Mali GPU看作一个「并行计算黑盒子」。你给它一堆数据，它哗啦一下全给你算完。但前提是，你得懂它的脾气。今天我们就从Mali的管线讲起，再到OpenCL怎么跟它打交道，最后亲手跑一个Kernel。嗯，走起。

4.1 Mali GPU管线：别被「图形」两个字骗了

Mali GPU虽然是做图形出身，但它的计算单元本质上是一堆Shader Core。我刚开始接触Mali时，总觉得它跟桌面GPU差不多，后来发现坑不少。

Mali的管线大致分这么几段：

顶点处理：处理几何数据，比如三角形的位置、颜色。在计算场景下，这部分可以被我们绕过去。
片元处理：处理像素级别的计算。OpenCL的Work-item其实就映射到这里。
执行引擎：真正的算力核心。Mali G76有多个执行引擎，每个引擎可以同时处理多个线程。

我在项目中遇到过一个问题：同样的算法，在Mali G52上跑得飞快，换到G76反而慢了。后来一查，是缓存命中率的问题。Mali的L2缓存是所有Shader Core共享的，如果你让多个Core同时访问同一块数据，就会产生竞争。说白了，就是「抢饭吃」。

关键点：Mali GPU的计算能力，取决于你如何组织数据。尽量让每个Work-group访问连续的内存地址，这样缓存友好。

4.2 OpenCL基础：给GPU下指令的语言

OpenCL，全称是Open Computing Language。你可以把它理解成一种「让CPU指挥GPU干活」的语言。它不像CUDA那样是NVIDIA独家的，OpenCL是开放的，ARM、Qualcomm、Imagination都支持。

OpenCL的核心概念其实不多，我总结成四个：

概念	通俗解释	我踩过的坑
Platform	就是你的硬件平台，比如Exynos 9820	同一个芯片可能有多个Platform，记得选对
Device	具体的计算设备，比如Mali GPU	别选成CPU了，否则白忙活
Context	运行环境，管理内存和命令队列	Context创建失败，八成是驱动没装好
Command Queue	你给GPU发的指令列表	队列顺序很重要，乱序执行会出bug

你可能会问：「为什么要有这么多层？」其实就是为了跨平台。你写的OpenCL代码，理论上可以在任何支持OpenCL的设备上跑。当然，实际优化还得针对具体硬件。

小技巧：在Exynos上开发，建议先调用clGetPlatformIDs打印出所有Platform，确认你用的是ARM的Mali平台。我曾经因为没检查，结果跑到了CPU上，性能差了一个数量级。

4.3 在Exynos上运行第一个GPU Kernel

好，理论说完了，咱们来点实际的。我带你写一个最简单的Kernel：两个向量相加。别小看这个例子，它包含了OpenCL的所有核心流程。

先看Kernel代码（.cl文件）：

__kernel void vec_add(__global const float *a,
                      __global const float *b,
                      __global float *c)
{
    int id = get_global_id(0);
    c[id] = a[id] + b[id];
}

这段代码很简单吧？每个Work-item处理一个元素。但注意，get_global_id(0)返回的是全局ID，不是局部ID。我刚开始写的时候，经常搞混get_local_id和get_global_id，结果数据全乱了。

然后是主机端代码（C语言）：

// 1. 获取平台和设备
cl_platform_id platform;
cl_device_id device;
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);

// 2. 创建上下文和命令队列
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);

// 3. 创建内存对象
cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, N * sizeof(float), NULL, NULL);
cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, N * sizeof(float), NULL, NULL);
cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, N * sizeof(float), NULL, NULL);

// 4. 将数据从CPU拷贝到GPU
clEnqueueWriteBuffer(queue, bufA, CL_TRUE, 0, N * sizeof(float), a, 0, NULL, NULL);
clEnqueueWriteBuffer(queue, bufB, CL_TRUE, 0, N * sizeof(float), b, 0, NULL, NULL);

// 5. 编译Kernel并设置参数
cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, NULL);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(program, "vec_add");
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);

// 6. 执行Kernel
size_t global_work_size = N;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);

// 7. 读取结果
clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, N * sizeof(float), c, 0, NULL, NULL);

注意：第5步的clBuildProgram很容易失败。我建议你加上错误检查：

cl_int err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if (err != CL_SUCCESS) {
    // 打印编译日志
    size_t log_size;
    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
    char *log = (char*)malloc(log_size);
    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
    printf("Build log: %s\n", log);
    free(log);
}

我曾经在Exynos 9810上遇到一个坑：Kernel里用了double类型，但Mali G72不支持double，编译直接报错。后来改成float才搞定。

4.4 性能调优：别让GPU闲着

第一个Kernel跑通了，恭喜你！但你可能发现，性能并没有想象中那么快。为什么？因为数据传输占了大部分时间。

我建议你记住这个原则：尽量减少CPU和GPU之间的数据拷贝。具体怎么做？

使用CL_MEM_ALLOC_HOST_PTR标志，让OpenCL在主机可访问的内存区域分配缓冲区。
用clEnqueueMapBuffer代替clEnqueueReadBuffer，避免显式拷贝。
如果可能，把多个小Kernel合并成一个，减少启动开销。

嗯，说到启动开销，Mali GPU的Kernel启动延迟大概在几十微秒级别。如果你的Kernel执行时间只有几微秒，那大部分时间都浪费在启动上了。这时候，你就得考虑把多个任务打包成一个Kernel。

实战建议：在Exynos上做GPU计算，我一般遵循「三步走」：

先跑通功能，不管性能。
用clGetEventProfilingInfo测量每个步骤的时间。
找到瓶颈，针对性地优化（通常是数据传输或Work-group大小）。

好了，今天的内容就到这。你回去可以试试把向量加法的例子跑起来，然后改一改Work-group大小，看看性能变化。下一章我们会深入Mali的缓存架构，讲讲怎么写出「缓存友好」的Kernel。到时候见！

📚 课程目录