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

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 人在当地工作的特别之处？

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

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

Answer 1

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

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

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

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

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

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

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

但是：

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

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