OpenCL：__constant 与 __local？

Question

假设我有一个大数组值（仍然小于 64 kB），在内核中经常非常读取，但不写入。然而，它可以从外部改变。该数组有两组值，让我们称它们为左右。 所以问题是，将大数组作为 __global 并将其写入 __local left 和 __local right 数组是否更快？还是将其作为一个常量 __constant large 并处理内核中的访问？例如：

__kernel void f(__global large, __local left, __local right, __global x, __global y) {
    for(int i; i < size; i++) {
        left[i] = large[i];
        right[i] = large[i + offset];
    }
    ...
    x = foo * left[idx];
    y = bar * right[idx];
}

对比：

__kernel void f(__constant large, __global x, __global y) {
    ...
    x = foo * large[idx];
    y = bar * large[idx * offset];
}

（索引有点复杂，但可以用宏来制作，例如） 我读到常量内存存在于全局空间中，所以它应该更慢吗？ 它将在 Nvidia 卡中运行。

Answer 1

首先，在第二种情况下，您应该通过某种方式使结果可用于您的 CPU。我假设您在计算后复制回globalspace。

我认为这取决于您在内核中所做的工作。例如，如果您的内核计算量很大（每个线程的计算量很大），那么第一个选项可能会付出代价。为什么？

• 您花了一些时间将数据从global large 空间复制到local 空间left 和right - 可以接受
• 你在本地空间的数据上做了很多计算 - OK
• 您花了一些时间从local left 和right 复制回global large。 - 可以接受。

但是，如果你的内核相对较轻，即每个线程都会做一些小的计算，那么

• 您对constant 空间上的数据进行了一些计算。这很可能意味着您不需要经常访问它。
• 您将中间结果存储在本地空间中。
• 您花了一些时间从local 空间复制回global 空间。 - 可以接受。

总而言之，对于大内核，第一个选项更好。第二个是小内核。

附：还有一点需要注意的是，如果你有多个内核一个接一个地在large 上工作，那么肯定会选择第一个选项。因为这样您就可以将数据保存在全局内存空间中，而不必每次启动内核时都进行复制。

编辑：既然你说它经常被非常访问，那么我认为你应该选择第一个选项。