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

Question

在我的内核中，有必要对一个小的查找表(只有 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。换句话说，查找表的所有位都被使用了，所以它不能被打包成一个更小的表示。

Answer 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 绑定(bind)的。 (当然，假设您有足够的备用寄存器用于 arr 变量。)

请注意，一些 OpenCL 目标甚至没有私有(private)内存，并且您在此处声明的任何内容都将转到 __global。在那里使用寄存器存储是一个更大的胜利。

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