OpenCL 上一维工作负载的 global_work_size 最高并行度的公式是什么？

Question

相关信息来自

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(...)
}

是否有某种公式可以“按硬件配置”计算出 1D 工作负载的理想 global_work_size 和 local_work_size ？有什么经验法则吗？

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 的正确平衡，以实现 1D 工作负载的最大并行性

Answer 1

此处本地工作组大小应为 32 或 32 的倍数，否则您将无法充分使用硬件，并且每个 CU 的某些核心将在您的设备上保持空闲状态。例如，在 Nvidia Pascal GPU 上使用大小为 16 的工作组将使一半的 CUDA 核心随时闲置，并且您只能获得一半的计算吞吐量。跨所有设备的最佳兼容、通用的工作组大小是 64，因为 AMD GPU 至少需要 64 或其倍数。

如果您的 OpenCL 内核不使用本地内存，则工作组大小 32、64、128 和 256 之间的性能差异通常可以忽略不计。

1D 全局范围应远大于 GPU 核心数，以实现硬件完全饱和。确切的值并不重要，但它必须是工作组大小的干净倍数，通常是通过将预期的全局范围四舍五入到下一个倍数。全球范围达到数百万甚至数十亿的情况并不罕见。 GPU 可以很好地处理这个问题。

示例：具有 320 个核心的 GPU 一次计算 320 个线程，如果全局范围为 512，则 GPU 需要内核的 2 次迭代，而第二次迭代仅使其部分饱和。最重要的是来自 PCIe 总线的约 10us 调度延迟。性能将是 < 512/(2*320) = 80% of peak. But if you choose global range >> GPU 核心数量，例如高 1000 倍，与几毫秒的内核执行时间相比，不完整的最后迭代和调度延迟将变得可以忽略不计，并且性能接近峰值的 100%。