工作组中部分封闭的工作人员

Question

让我们考虑一下，在 OpenCL 中有多个工作组和多个工作人员时。如果我们在工作组中拥有与 GPU 上的“核心”一样多的工作人员，则 GPU 将按顺序工作，其中工作组中的每个工作人员并行工作（对吗？）。完成一个工作组后，将执行下一个工作组。如果我们在工作组中的工作人员数量少于 GPU 上的“核心”，据我了解，GPU 将并行执行多个工作组，这会导致多个工作人员并行执行（对吗？）。在这种情况下，执行这段代码会发生什么？

__kernel void vector_add(__global const int *A) 
{ 
    // Get the index of the current element to be processed
    int i = get_global_id(0);
    if(A[i] == 0)
    {
        return; //stop like the half of all workers in a workgroup
    }
    // Do some time consuming calculation
    ...
}

这段代码会导致“分支”，但工作组中的某些工作人员会直接停止（return）其他人进行一些耗时的计算。我们可以称之为“分支”吗？一个大问题：returned 的“核心”会做什么？他们是否等到工作组中的每个工人都完成工作？或者因为他们中的很多人同时 returning，所以他们会跳到下一个工作组执行吗？

行为是否特定于供应商？或者在 OpenCL 中是否正确定义了这种情况？

Answer 1

如果您在内核中有分支并且在工作组中一些工作人员执行分支 A 和一些分支 B，则所有工作人员都必须计算两个分支并分别丢弃未使用的分支结果。这会对执行时间产生负面影响，这也是应尽可能避免在 GPU 上进行分支的原因。在您的 return 分支为空的示例中，如果工作组中只有一名工作人员必须执行 time consuming calculation，所有其他工作人员都必须等待，从而阻塞其他工作组的硬件资源。如果工作组很小，并且您很幸运所有线程都执行 return 分支，那么该特定工作组的执行速度非常快。

物理 GPU“核心”与工作组大小之间的匹配与计算结果无关，但会在一定程度上影响性能。 Workgroup 大小应该是 32 的倍数（GPU 将其“核心”细分为 32 个一组，即所谓的 warps）。因此，如果工作组大小为 16，则一半的 GPU 将始终处于空闲状态。另一方面，如果工作组大小非常大（例如 1024）并且您在内核中有分支，那么所有工作人员都执行相同分支并且最终出现上述情况的可能性较小。

如果您需要通过 local 内存在工作组之间进行通信，有时工作组大小需要权衡。更大的工作组允许更多的本地通信，但增加了“双分支”的可能性。如果您不使用 local 内存，您可以自由调整工作组大小以获得最佳性能（通常为 64-256）。

理想情况下，您希望用数百万个线程使 GPU 饱和，从而没有空闲的“核心”并获得最佳性能。

工作组中部分封闭的工作人员

partly closed worker in workgroup

opencl