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

1次阅读

没有评论

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

移动设备的 GPU（比如 ARM Mali 系列）正在成为处理密集型任务的关键组件。与传统的 CPU 相比，GPU 擅长并行计算，特别适合以下场景：

计算机视觉：图像处理、物体识别等任务可以分解为大量相似的并行操作
机器学习推理：神经网络中的矩阵运算能充分利用 GPU 的并行计算能力
实时滤镜和特效：视频处理需要高吞吐量的像素操作

在移动设备上使用 GPU 计算不仅能提升性能，还能显著降低功耗，这对电池供电的设备尤为重要。

ARM Mali GPU 采用了一些与桌面 GPU 不同的设计理念：

Tile-Based 渲染：
将屏幕分成小块（tile）分别渲染
减少内存带宽需求，特别适合移动设备
需要特别注意内存访问模式
内存层级结构：
全局内存（Global Memory）：所有工作项可见，访问延迟高
局部内存（Local Memory）：工作组内共享，访问速度快
寄存器（Private Memory）：每个工作项私有，速度最快
节能设计：
更精细的功耗管理
动态调整频率和电压
对热设计更敏感

OpenCL 内核是在 GPU 上执行的函数，基本结构如下：

__kernel void matrix_mult(__global const float* A,
                          __global const float* B,
                          __global float* C,
                          const int width) {
    // 获取当前工作项的全局 ID
    int row = get_global_id(0);
    int col = get_global_id(1);

    // 执行矩阵乘法
    float sum = 0.0f;
    for(int k = 0; k < width; ++k) {sum += A[row * width + k] * B[k * width + col];
    }

    // 写入结果
    C[row * width + col] = sum;
}

关键点说明：

__kernel修饰符表示这是一个 OpenCL 内核函数
__global限定符表示指针指向全局内存
get_global_id()获取当前工作项的索引

OpenCL 的执行模型基于以下概念：

工作项（Work-item）：
最基本的执行单元
类似于 CPU 上的线程
通过 get_global_id() 获取唯一标识
工作组（Work-group）：
工作项的集合
组内工作项可以同步和共享局部内存
通过 get_local_id() 获取组内标识
NDRange：
定义整个计算空间的维度
可以是 1D、2D 或 3D
需要合理选择工作组大小

#include <CL/cl.h>
#include <vector>
#include <iostream>

int main() {
    // 初始化 OpenCL 环境
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;

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

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

    // 3. 准备矩阵数据
    const int size = 1024;
    std::vector<float> A(size * size, 1.0f);
    std::vector<float> B(size * size, 2.0f);
    std::vector<float> C(size * size, 0.0f);

    // 4. 创建内存对象
    cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                sizeof(float) * A.size(), A.data(), NULL);
    cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                sizeof(float) * B.size(), B.data(), NULL);
    cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                                sizeof(float) * C.size(), NULL, NULL);

    // 5. 创建程序对象
    const char* source = "..."; // 内核代码字符串
    program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
    clBuildProgram(program, 1, &device, NULL, NULL, NULL);

    // 6. 创建内核对象
    kernel = clCreateKernel(program, "matrix_mult", NULL);

    // 7. 设置内核参数
    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), &size);

    // 8. 执行内核
    size_t global[2] = {size, size};
    clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, NULL, 0, NULL, NULL);

    // 9. 读取结果
    clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0,
                       sizeof(float) * C.size(), C.data(), 0, NULL, NULL);

    // 10. 清理资源
    clReleaseMemObject(bufA);
    clReleaseMemObject(bufB);
    clReleaseMemObject(bufC);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);

    return 0;
}

局部内存优化：

__kernel void optimized_matmul(__global const float* A,
                              __global const float* B,
                              __global float* C,
                              const int width) {
    // 使用局部内存缓存数据块
    __local float tileA[16][16];
    __local float tileB[16][16];

    int row = get_global_id(0);
    int col = get_global_id(1);
    int localRow = get_local_id(0);
    int localCol = get_local_id(1);

    float sum = 0.0f;

    // 分块计算
    for(int t = 0; t < width/16; ++t) {
        // 从全局内存加载数据到局部内存
        tileA[localRow][localCol] = A[row * width + (t * 16 + localCol)];
        tileB[localRow][localCol] = B[(t * 16 + localRow) * width + col];

        // 确保所有工作项完成数据加载
        barrier(CLK_LOCAL_MEM_FENCE);

        // 计算当前块
        for(int k = 0; k < 16; ++k) {sum += tileA[localRow][k] * tileB[k][localCol];
        }

        // 等待所有工作项完成计算
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    C[row * width + col] = sum;
}