Mali Midgard系列GPU点积运算支持情况及最优OpenCL内核咨询

点积运算支持情况 #

Mali T-624属于ARM Midgard架构T600系列GPU，原生支持硬件级点积运算，完全兼容OpenCL标准定义的点积内置函数，无需软件模拟实现。
Midgard架构的ALU流水线针对单精度浮点、8/16位整型的点积场景做了专门优化，4分量点积运算可在单周期内完成乘加合并计算，相比手动拆分的逐分量乘加代码，指令开销可降低75%左右。

最优效率OpenCL内核编写优化要点 #

• 优先调用内置点积函数：直接使用dot()内置函数完成计算，不要手动编写分量相乘后累加的代码，编译器会自动将dot()映射为专属硬件点积指令，避免冗余的指令调度开销。
• 保证访存地址对齐：输入缓冲区指针建议添加__attribute__((aligned(16)))修饰，匹配4分量向量的16字节宽度要求，消除非对齐访存带来的额外延迟。
• 使用向量化数据类型：优先采用float4、int4等4分量向量类型存储输入数据，完全匹配Midgard架构的向量寄存器宽度，单次访存即可加载4个分量数据，充分利用硬件点积的并行计算能力。
• 合并访存减少带宽占用：如果需要批量计算大量点积，建议每个工作项负责计算多组点积，将零散的随机访存合并为连续的块访存，降低显存带宽瓶颈对性能的影响。
• 明确指针别名与工作组大小：给指针添加restrict关键字明确无地址重叠，使用__attribute__((reqd_work_group_size(256, 1, 1)))指定工作组大小匹配Midgard架构的波前宽度，帮助编译器做更激进的优化。

参考实现代码 #

__kernel __attribute__((reqd_work_group_size(256, 1, 1)))
void batch_dot_product(__global const float4* restrict a,
                       __global const float4* restrict b,
                       __global float* restrict output,
                       const int total_vec_num) {
    int global_id = get_global_id(0);
    float acc = 0.0f;
    // 每个工作项计算16组点积，合并访存提升带宽利用率
    int start_offset = global_id * 16;
    int end_offset = min(start_offset + 16, total_vec_num);
    for (int i = start_offset; i < end_offset; i++) {
        acc += dot(a[i], b[i]);
    }
    output[global_id] = acc;
}

内容的提问来源于stack exchange，提问作者marios