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

请注意，某些 OpenCL 目标甚至没有私有内存，您在那里声明的任何内容都只会进入 __global。使用寄存器存储是一个更大的胜利。

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