在我的内核中，有必要对一个小的查找表进行大量随机访问(仅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。换句话说，正在使用查找表的所有位，因此不能将其打包成较小的表示形式。
最佳答案
我能想到的最快的解决方案是首先不要使用数组:而是使用单个变量，并使用某种访问函数来访问它们，就像它们是数组一样。 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数据。
关于opencl - 在OpenCL内核中实现小型查找表的最佳方法是什么，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/7840720/