调整OpenCL工作组大小未提升内存合并性能，为何反而变慢？

云原生阿岚

2026-6-17

问题描述

我向GPU上传了一个包含一系列“records（记录）”的缓冲区。记录数量在50到10000之间，所有记录长度相同（通常约40000个float值）。随后运行核函数处理缓冲区中每个记录内指定“regions（区域）”的值，共约200个区域，宽度为10-100个值，且所有记录共用该区域列表。

首次尝试的核函数使用全局工作大小[num records, num regions]、本地工作大小为NULL，每个线程处理一个区域的所有值，代码如下：

for (int i = regionStart; i <= regionEnd; i++)
{
    buff[i] = do_something(buffer[i]);
}

（核函数通过get_global_id(0)和get_global_id(1)，结合传入的区域位置查找表计算regionStart和regionEnd）

为提升并行性与内存合并性能，我将全局工作大小改为[num records, num regions, 64]，本地工作大小设为[1, 1, 64]，预期每个区域由64个线程处理，核函数代码更新为：

int i = get_local_id(2); // 0-64
while (i + regionStart <= regionEnd)
{
    buff[i] = do_something(buffer[i]);
    i += get_local_size(2);
}

但核函数执行时间反而比原版本慢数倍。运行环境为Windows，显卡为Radeon Pro WX7100，请问我忽略或误解了什么？

问题分析与解决方案

新方案不仅没实现内存合并，反而彻底破坏了原有的连续访问优势：

原代码中单个线程连续遍历区域内的float值，内存访问是连续的，GPU内存控制器可以通过批量事务高效读写，带宽利用率极高。
新代码里，每个线程访问的是regionStart + get_local_id(2)、regionStart + get_local_id(2)+64这类间隔256字节（64个float）的地址，属于分散式随机访问。GCN架构对这种访问模式的处理效率极低，会触发大量独立内存事务，直接导致内存带宽利用率暴跌，这是性能骤降的核心原因。

Radeon Pro WX7100基于GCN 4架构，单计算单元（CU）含64个流处理器（SP），要求线程组尺寸为64的倍数以最大化SIMD利用率，但你的方案存在两个致命问题：

区域宽度仅10-100，用64个线程处理时，大量线程在第一次循环后就闲置（比如宽度10时，仅前10个线程有工作，剩下54个全程空转），造成严重的线程资源浪费。
原方案的线程总数是num_records * 200，这个规模已经足够填满WX7100的硬件线程池（共1280个SP），GPU能保持高利用率；新方案线程数虽多，但有效工作的线程占比极低，整体利用率反而大幅下降。

要真正提升性能，需从内存访问和线程映射两方面调整：

恢复连续内存访问：让相邻线程访问相邻内存地址。比如将全局工作大小设为[num_records, total_processed_elements]，每个线程处理单个元素，确保同一线程组内的线程访问同一条记录中的连续元素。
适配区域尺寸的线程粒度：针对10-100的区域宽度，选择32或64作为线程组尺寸，让线程组处理区域内的连续块（比如一个线程组处理64个连续元素，不足部分用循环补全），避免分散访问。
避免过度拆分线程：当单个线程仅处理1-2个float时，线程调度和同步的开销会超过计算收益，反而拖慢整体性能，需保证每个线程有足够的计算量。

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