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

Question

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

6

在我的内核中，需要对一个只包含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内存空间中可获得最佳性能。然而，由于查找表的访问方式是随机的，因此仍会存在大量性能损失。通过删除查找表代码（例如替换为简单算术运算），虽然内核会提供错误答案，但性能会提高三倍以上。

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

【编辑】注意，X的最大值为7，但Y的最大值可达2 ^ 32-1。换句话说，查找表的所有比特位都被使用了，因此不能将其打包成更小的表示形式。

- Dustin

只是为了确保我理解正确，LUT []和X []对于每个单独的工作项都是唯一的吗？ - Adam S.

你尝试过在查找表中使用__constant内存吗？GPU通常为常量内存实现单独的缓存和内存访问路径，以加速共享查找表等操作。 - user57368

我曾尝试将global_table作为__constant传递，但由于某种原因，这并没有提高性能。每个内核都在独立的global_table、X和Y部分上操作（取决于线程ID）。 - Dustin

如果是这种情况，__constant 变量很可能没有帮助，因为在我看过的大多数 GPU 上，它基本上只是全局内存的一个小缓存。如果每个线程读取不同的值，那么缓存就无法发挥作用。 - Adam S.

英特尔的 OpenCL 指南讨论了在 LUTs 中使用“共享本地内存”。https://software.intel.com/en-us/articles/accelerate-performance-using-opencl-with-intel-hd-graphics#_Toc357619019 在一般情况下可能会有用。 - Peter Cordes

2个回答

1

正如rtollert所述，将LUT[]放置在寄存器中还是全局内存中取决于实现。通常情况下，内核中的数组是不推荐使用的，但由于它很小，很难说它会被放置在哪里。假设LUT[]被放置在寄存器中，我认为它比简单算术运算花费更长时间的原因不是因为它是随机访问的，而是因为每个工作项需要额外进行8次（编辑：显然更多）全局读取X来计算LUT索引。根据省略的内容，您是否可以像这样做：Y[i*gsi+gid] = global_table[someIndex + X[i*gsi+gid]]];？

- Adam S.

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

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- rtollert · Accepted Answer

我能想到的最快解决方案是首先不使用数组：改为使用单独的变量，并使用某种访问函数来访问它们，就像它们是一个数组一样。我IRC（至少对于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数据。