OpenCL下AMD GPU全局内存带宽仅达clpeak的4%，求优化建议

问题背景 #

在AMD RX 5700XT（gfx1010）GPU上使用OpenCL 2.0搭配ROCm驱动时，核心计算因不可避免的全局内存随机读写，仅能达到clpeak报告375GB/s全局内存带宽的约5%（~20GB/s），已排除CL_MEM_SVM_FINE_GRAIN_BUFFER导致的性能下降问题。复现测试内核如下：

kernel void testkernel(__global const ulong *loc,__global uchar *res,__global const uchar *val,__private ulong bsize)
{
    uint n= get_global_id(0);
    ulong i= n*bsize;
    ulong j= i+bsize;
    uchar z= 0;
    for (ulong k=i;k<j;++k)
    {
        ulong l= loc[k];
        ulong l2= loc[l];
        z+= val[l2];
    }
    res[n]= z;
}

测试参数：bsize=1000，loc/val数组填充随机值，通过clSVMAlloc分配内存，测试不同全局工作项数（10K/100K/700K）时带宽无明显提升。

当前性能受限的核心原因 #

1. 随机访问的存储体冲突与理想带宽差异
   clpeak的375GB/s是连续内存访问的理想带宽，而随机访问会触发大量存储体冲突（RX 5700XT有32个存储体），每个内存请求的延迟无法被有效隐藏，实际有效带宽会暴跌至理想值的几分之一甚至更低。

2. 内存请求串行化与延迟隐藏不足
   内核中每个线程的内循环是串行内存操作：先读loc[k]，再读loc[l]，最后读val[l2]，单线程内无法并行发起内存请求。GPU依赖大量活跃线程来掩盖单内存请求的延迟（RX 5700XT需要至少数千个活跃线程才能充分利用硬件），若活跃线程数不足，延迟无法被隐藏，整体性能会被内存延迟限制，而非带宽。

3. 细粒度数据的内存事务浪费
   val数组是uchar类型（1字节），而gfx10架构的全局内存事务最小单位是64字节。每次读取val[l2]都会触发64字节的内存事务，但仅使用其中1字节，造成约98%的带宽浪费，这是有效带宽极低的关键原因之一。

针对性优化建议 #

• 合并细粒度内存访问

  • 将val数组的uchar元素打包为ulong（8字节）或更大粒度的类型，单次内存读取可获取8个uchar值，大幅减少内存请求次数，提升有效带宽利用率。
  • 对loc数组的访问使用__global_prefetch指令提前发起预取，让内存操作与计算重叠，隐藏部分延迟。例如在循环中提前预取下一个loc[k]：

    __global_prefetch(loc[k+1]);
    ulong l= loc[k];
    

• 优化线程并发与任务拆分

  • 减小bsize，同时增加全局工作项数量，让活跃线程数达到GPU的推荐值（RX 5700XT建议至少2048个，最好8192以上），充分利用多线程调度掩盖内存延迟。
  • 避免单线程内的长串行循环，将工作拆分为更多并行小任务，让GPU的SIMD单元和内存控制器同时满负荷运行。

• 利用缓存提升命中率

  • 指定工作组大小为64（gfx10架构的SIMD宽度），通过__attribute__((reqd_work_group_size(64,1,1)))修饰内核，让工作组内线程共享L1缓存，提升随机访问的缓存命中率。
  • 若loc数组存在重复访问，将常用的loc值缓存到私有内存或局部内存中，减少全局内存访问次数。

• 内存分配方式验证

  • 尝试用普通的clCreateBuffer代替clSVMAlloc分配内存，对比两者性能差异，排除SVM带来的潜在额外开销。
  • 确保内存分配按64字节对齐，可通过指定CL_MEM_ALLOC_HOST_PTR或手动对齐数据，提升内存事务的执行效率。

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