OpenCL 1D工作负载最大并行度的global_work_size公式及最优配置问询

问题核心 #

针对1D工作负载，如何根据硬件配置计算能实现最大并行度的global_work_size和local_work_size？是否有可遵循的公式或经验法则？

硬件配置（来自clinfo） #

Max compute units                               10
Max work item dimensions                        3
Max work item sizes                             256x256x256
Max work group size                             256
Preferred work group size multiple (kernel)     32

暴力尝试方案 #

本质是按批次循环调度内核：

loop {
  clEnqueueNDRangeKernel(...)
}

当前尝试的代码方案 #

let max_compute_units = 10;
let global_work_size = 256 * max_compute_units;
let local_work_size = 256;
assert!(global_work_size % local_work_size == 0);
assert!(local_work_size % 32 == 0);

该方案仍有优化空间，尚未找到主机向GPU调度工作的最优平衡。

经验法则与计算方法 #

1. local_work_size的选择原则 #

• 硬性约束：
  • 不能超过硬件的Max work group size（此处为256）
  • 必须是Preferred work group size multiple的整数倍（此处为32）
  • 1D场景下不超过Max work item sizes第一维（此处为256）
• 性能导向：
  • 优先选择与硬件Wavefront/Warp大小匹配的倍数：比如NVIDIA Warp为32、AMD Wavefront为64，结合此处硬件推荐的32倍数，可测试32、64、128、256这些值。
  • 避免过大的local_work_size：如果核函数依赖共享内存，过大的工作组会占用更多共享内存资源，反而降低并行效率。

2. global_work_size的选择原则 #

• 基础要求：global_work_size必须是local_work_size的整数倍，避免零散工作项被低效调度。
• 充分利用硬件并行性：
  • 核心公式：global_work_size = k × Max compute units × (每个计算单元支持的最大并发工作项数 / local_work_size)
  • 其中k为1_{4的整数（用来隐藏内存访问延迟，数值越大，延迟隐藏效果越好）；每个计算单元支持的最大并发工作项数可通过硬件规格估算（比如多数GPU单CU支持1024}2048个并发工作项）。
  • 以当前硬件为例，若单CU支持1024个并发工作项，local_work_size选64，则每个CU可容纳16个工作组，此时global_work_size可设为2 × 10 × 16 × 64 = 20480（k取2）。
• 固定工作负载场景：若总任务量为N，global_work_size取大于等于N的最小local_work_size整数倍，同时尽量满足上述硬件填充要求。

3. 优化验证方向 #

• 结合核函数特性调整：计算密集型核函数优先保证每个CU的并发工作项数达上限；内存密集型核函数需增大global_work_size以隐藏内存延迟。
• 实测对比：固定local_work_size为32、64、128、256，分别测试不同倍数的global_work_size，对比执行时间找到最优组合。

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