OpenCL：__constant对__local？

Question

假设我有大量的值(仍然小于64 kB)，这些值通常在内核中被读取，但不是写到的。然而，它可以从外部改变。数组有两组值，让我们调用它们的左和右。所以问题是，将大数组作为__global并将其写入__local左数组和右数组是更快吗?还是将其作为常量__constant大并处理内核中的访问呢？例如：

__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];
}

(索引有点复杂，但可以使用宏，例如)我读到常量内存存在于全局空间，那么应该慢一点吗？它会在一张英伟达卡里运行。

Answer 1

首先，在第二种情况下，您应该有某种方法将结果提供给您的CPU。我假设您在计算后复制回global space。

我认为这取决于您在内核中做了什么。例如，如果内核计算量很大(每个线程的计算量很大)，那么第一个选项可能会支付。为什么？

• 您需要花费一些时间将数据从global large空间复制到local空间left和right --可以接受。
• 你对本地空间上的数据做了大量的计算--好的
• 您需要花一些时间从local left和right复制回global large。-可以接受。

但是，如果您的内核相对较轻，即每个线程将进行一些小的计算，那么

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

总之，对于大内核来说，第一种选择更好。对于第二颗小粒来说。

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

编辑:既然您已经说过它经常被访问，所以我认为您应该选择第一个选项。