在 OpenCL 内核中实现小型查找表的最佳方法是什么答案

【问题标题】：What is the best way to implement a small lookup table in an OpenCL Kernel在 OpenCL 内核中实现小型查找表的最佳方法是什么
【发布时间】：2011-10-21 06:01:27
【问题描述】：

在我的内核中，有必要对一个小的查找表（只有 8 个 32 位整数）进行大量随机访问。每个内核都有一个唯一的查找表。下面是内核的简化版本，用于说明如何使用查找表。

__kernel void some_kernel(  
    __global uint* global_table,
    __global uint* X,
    __global uint* Y) {

    size_t gsi = get_global_size(0);
    size_t gid = get_global_id(0);

    __private uint LUT[8]; // 8 words of of global_table is copied to LUT

    // Y is assigned a value from the lookup table based on the current value of X
    for (size_t i = 0; i < n; i++) {
        Y[i*gsi+gid] = LUT[X[i*gsi+gid]];
    }   
}

由于体积小，我通过将表保存在 __private 内存空间中获得了最佳性能。但是，由于访问查找表的随机性，仍然存在很大的性能损失。删除查找表代码（例如，替换为简单的算术运算）后，虽然内核会提供错误的答案，但性能提高了 3 倍以上。

有没有更好的方法？我是否忽略了一些为非常小的内存块提供有效随机访问的 OpenCL 功能？是否有使用向量类型的有效解决方案？

[edit] 注意，X 的最大值为 7，但 Y 的最大值为 2^32-1。换句话说，查找表的所有位都在使用，因此不能将其打包成更小的表示形式。

【问题讨论】：

只是为了确保我理解正确，LUT[] 和 X[] 对于每个工作项都是唯一的吗？
您是否尝试过将__constant 内存用于查找表？ GPU 通常为恒定内存实现单独的缓存和内存访问路径，以加快共享查找表之类的速度。
我确实尝试将 global_table 作为 __constant 传递，但由于某种原因，它对性能没有帮助。每个内核都在 global_table、X 和 Y 的独立部分上运行（取决于线程 ID）。
如果是这种情况，__constant 很可能无济于事，因为至少在我见过的大多数 GPU 上，它基本上是一个用于全局内存的小型缓存。如果每个线程读取不同的值，缓存将无济于事。
英特尔的 OpenCL 指南谈到了将“共享本地内存”用于 LUT。 software.intel.com/en-us/articles/… 一般可能有用。

标签： opencl

【解决方案1】：

我能想到的最快的解决方案是首先不使用数组：使用单个变量，并使用某种访问函数来访问它们，就好像它们是一个数组一样。 IIRC（至少对于 AMD 编译器，但我很确定 NVidia 也是如此）：通常，数组总是存储在内存中，而标量可能存储在寄存器中。（但我对这件事有点模糊——我可能错了！）

即使你需要一个巨大的 switch 语句：

uint4 arr0123, arr4567;
uint getLUT(int x) {
    switch (x) {
    case 0: return arr0123.r0;
    case 1: return arr0123.r1;
    case 2: return arr0123.r2;
    case 3: return arr0123.r3;
    case 4: return arr4567.r0;
    case 5: return arr4567.r1;
    case 6: return arr4567.r2;
    case 7: default: return arr4567.r3;
    }
}

...与 __private 数组相比，您可能仍然在性能上领先，因为假设 arr 变量都适合寄存器是纯 ALU 绑定的。（当然，假设您有足够的备用寄存器用于 arr 变量。）

注意，一些 OpenCL 目标甚至没有拥有私有内存，并且您在其中声明的任何内容都只会转到 __global。在那里使用寄存器存储是一个更大的胜利。

当然，这种 LUT 方法可能初始化较慢，因为您需要至少两次单独的内存读取才能从全局内存中复制 LUT 数据。

【讨论】：

我之前已经排除了这样的解决方案，因为我害怕引起warp发散，但我会试一试看看。为什么你使用 2 个 uint4 向量而不是单个 uint8，或者只使用 8 个变量？
这将性能提高了 15% ......不是我希望的 200% 到 300%，但每一点都有帮助。
我用 uint4s 因为我是个白痴。 :) 我想 uint8s 也可以。将值收集到单个逻辑变量中意味着您可以使用 vload8() 来初始化表，而不是（可能）进行八次单独的内存访问，每个变量一次。但你已经知道了。
实际上，几乎没有分歧。 Switch 语句通常实现为算术 jmps（计算 goto）。因为在这些案例中你没有做任何工作，所以你会在最多一条指令中出现分歧，这一点都不错。

【解决方案2】：

正如 rtollert 所说，由实现决定 LUT[] 是放在寄存器中还是放在全局内存中。通常内核中的数组是禁止的，但由于它很小，很难说它会放在哪里。假设 LUT[] 被放入寄存器中，我会说与简单的算术运算相比它需要很长时间的原因不是因为它是随机访问的，而是因为每个工作项都会产生额外的 8（编辑：显然更多）全局读取 X 以计算 LUT 索引。根据省略的内容，您可以执行类似 Y[i*gsi+gid] = global_table[someIndex + X[i*gsi+gid]]]; 之类的操作吗？

【讨论】：

来自全局内存的 8 次读取被合并，并且也在循环之外。由于 n 很大（通常约为 1024），因此这些读取复制 LUT 的开销有效地分摊了。