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

在我的内核中，需要对一个小查找表（只有 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 功能，该功能可以为非常小的内存块提供高效的随机访问？是否有使用向量类型的有效解决方案？

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

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

即使您需要一个巨大的 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 目标甚至不have私有内存，您在那里声明的任何内容都会进入 __global。使用寄存器存储是一个更大的胜利。

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