【问题标题】:Computing Maximum Concurrent Workgroups计算最大并发工作组
【发布时间】:2015-08-11 19:00:29
【问题描述】:

我想知道是否有一种标准方法可以以编程方式确定可在 GPU 上运行的最大并发工作组数。

例如,在具有 5 个计算单元(或 SM)的 NVIDIA 卡上,每个计算单元最多可以有 8 个工作组(或块),因此可以同时运行的最大工作组数为 40。

因为我可以使用clGetDeviceInfo 找到计算单元的数量,所以我只需要一个计算单元上可以运行的最大工作组数。

谢谢!

【问题讨论】:

    标签: opencl gpu gpgpu


    【解决方案1】:

    每个执行单元/SM 的最大组数受硬件资源限制。让我以 Intel Gen8 GPU 为例。每个子片包含 16 个屏障寄存器。因此,最多可以同时运行 16 个工作组。

    此外,每个子片可用的共享本地内存量 (64KB)。例如,如果一个工作组需要 32KB 的共享本地内存,那么无论工作组大小如何,这些工作组中只有 2 个可以同时运行。

    【讨论】:

      【解决方案2】:

      我通常使用计算单元的数量作为工作组的数量。我喜欢扩大组的大小以使硬件饱和,而不是强制 gpu '同时'安排许多工作组。

      我不知道如何在不查看供应商规范的情况下确定最大组数。

      【讨论】:

      • 我无法从任何供应商文档(AMD、Intel)中找到最大工作组数,我认为没有像最大工作组数这样的东西,工作组数仅受以下限制工作项的数量,所以如果你的工作项是 1024,那么1024/CL_DEVICE_MAX_WORK_GROUP_SIZE 确定没有工作组。
      • 情况并非如此。当您使用 clEnqueueNDRangeKernel 时,您会给出一个 local_work_size 参数。这是工作组的大小(和维度)。您还可以指定 global_work_size,它是全局大小和维度。假设一维布局 工作组的数量是 global_work_size/local_work_size。最初的问题是询问如何找出 gpu 上“正在运行”的并发工作组的限制。理论上可能没有这样的限制,但在实践中,缓存和寄存器交换会显着影响超过一些组的性能。在其他时候这是不可能的。
      • 如果你使用这种方法,它不会给你保证设备上的完全占用。另外我不想限制可以从主机启动的工作组的数量。
      • @ocluser 这不能保证有效。某些 NVIDIA 卡(Fermi 架构)允许每个工作组最多有 1024 个工作项,但允许在一个计算单元上运行 1536 个线程。因此,如果我有 512 个,那么您的方法将产生每个计算单元 2 个工作组,而实际上它是 3 个。
      猜你喜欢
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 1970-01-01
      • 2015-07-31
      相关资源
      最近更新 更多