ARM GPU编程入门指南：从基础概念到高效并行计算实践

1次阅读

共计 2404 个字符，预计需要花费 7 分钟才能阅读完成。

近年来，移动设备上的计算需求爆炸式增长，尤其是在计算机视觉和机器学习推理领域。传统 CPU 处理这些任务时往往力不从心，而 ARM Mali GPU 凭借其并行计算能力，可以显著提升性能。比如在手机上实时运行人脸识别、图像风格迁移等应用，GPU 加速能带来 3 - 5 倍的性能提升，同时保持较低的功耗。

与桌面级 GPU 不同，ARM Mali 采用 Tile-Based 渲染架构，这种设计更适合移动设备的低功耗需求。主要差异体现在：

内存层级：Mali GPU 的缓存体系更精简，全局内存带宽有限，因此需要特别注意内存访问模式
执行单元：采用较小的计算单元集群，适合细粒度并行任务
功耗管理：有更精细的功耗控制机制，频繁的内存访问会显著增加功耗

OpenCL 内核是运行在 GPU 上的函数，需要使用 __kernel 修饰符。一个典型的矩阵乘法内核如下：

__kernel void matrix_mult(__global const float* A,
                          __global const float* B,
                          __global float* C,
                          const int N) {int i = get_global_id(0); // 工作项 ID
    int j = get_global_id(1);

    float sum = 0.0f;
    for (int k = 0; k < N; ++k) {sum += A[i*N + k] * B[k*N + j];
    }
    C[i*N + j] = sum;
}

OpenCL 使用两级并行模型：

工作项(Work-item)：最基本的执行单元，相当于 CUDA 中的线程
工作组(Work-group)：工作项的集合，共享局部内存

与 CPU 线程模型不同，GPU 工作项是轻量级的，通常需要启动数千个才能充分利用 GPU 资源。

以下是包含 host 端代码的完整示例，着重展示了内存对象处理：

// Host 端代码（省略错误检查）cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, N*N*sizeof(float), NULL, &err);
cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, N*N*sizeof(float), NULL, &err);
cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, N*N*sizeof(float), NULL, &err);

// 数据传输到设备
clEnqueueWriteBuffer(queue, bufA, CL_TRUE, 0, N*N*sizeof(float), A, 0, NULL, NULL);
clEnqueueWriteBuffer(queue, bufB, CL_TRUE, 0, N*N*sizeof(float), B, 0, NULL, NULL);

// 设置内核参数
clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);
clSetKernelArg(kernel, 3, sizeof(int), &N);

// 执行内核
size_t globalSize[2] = {N, N};
clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0, NULL, NULL);

// 回读结果
clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, N*N*sizeof(float), C, 0, NULL, NULL);

局部内存 (Local Memory) 是工作组内共享的高速内存，合理使用可以大幅减少全局内存访问：

__kernel void optimized_matmul(__global const float* A,
                              __global const float* B,
                              __global float* C,
                              const int N) {__local float Asub[BLOCK_SIZE][BLOCK_SIZE];
    __local float Bsub[BLOCK_SIZE][BLOCK_SIZE];

    // 从全局内存加载数据块到局部内存
    int bx = get_group_id(0);
    int by = get_group_id(1);
    ...
}

Mali GPU 的局部内存通常分为多个 bank，要确保工作组内不同工作项不要同时访问同一个 bank。解决方法包括：