opencl - 我应该如何查看全局和本地工作规模

Question

出于爱好目的，我使用 OpenCL 已经有一段时间了。我想知道是否有人可以解释我应该如何查看全局和本地工作空间。我已经玩了一段时间了，但我似乎无法绕开它。

我有这段代码，内核的全局工作大小为8，本地工作大小为4

    __kernel void foo(__global int *bar)
    {    
        bar[get_global_id(0)] = get_local_id(0);
    }

bar 中的结果如下所示：

{0, 1, 2, 3, 0, 1, 2, 3, 4}

由于我使用的工作尺寸，我知道为什么会发生这种情况。但我似乎无法理解我应该如何看待这个问题。

这是否意味着有 4 个线程在本地工作，8 个在全局工作，所以我总共有 4 * 8 个线程在运行？如果是这样，是什么让这 4 个在当地工作的人特别？

或者这是否意味着内核的主体只有两个计数器？一个来自本地，一个来自全球，但这有什么意义呢？

我知道我可能有点含糊，我的问题可能看起来很愚蠢。但我不知道如何更优化地使用它以及我应该如何看待它？

标签： opencl

Answer 1

• 全局大小是工作项的总数。
• 工作组细分了这个总工作量，本地大小定义了全局大小内每个组的大小。

因此，对于 8 的全局工作大小和 4 的局部大小，每个都在一维中，您将有 2 个组。每个线程的您get_global_id(0)将有所不同：0…7。get_local_id(0)将为每个组中的 4 个不同线程返回 0…3。这就是您在输出的索引 0 到 7 中看到的内容。

这也意味着如果您的全局工作大小为 8，那么bar您的内核将只设置前 8 项。因此，除此之外的任何内容（4输出中索引 8 处的值）都是未定义的。

这是否意味着有 4 个线程在本地工作，8 个在全局工作，所以我总共有 4 * 8 个线程在运行？如果是这样，是什么让这 4 个在当地工作的人特别？

你想多了。一共有8个线程。它们被细分为 2 组，每组 4 个线程。这些组中线程的“本地”是它们共享对同一local内存的访问。不在同一组中的线程只能通过全局内存“通信”。

使用本地内存可以极大地提高某些工作负载的效率：

• 它非常快。
• 工作组中的线程可以使用屏障来确保它们处于锁步状态，即它们可以相互等待以保证另一个线程已写入特定的local内存位置。（不同组中的线程不能互相等待。）

但：

• 本地内存很小（通常为几 KiB） - 在一组中使用所有内存通常会进一步降低效率。
• 本地内存必须由内核内部的数据填充，内核完成后其内容会丢失。（OpenCL 2 中的设备调度内核除外）
• 由于硬件限制，组大小有严格的限制。

因此，如果您不使用本地内存，工作组和本地工作大小基本上与您无关。