在我的内核中,有必要对一个小的查找表进行大量的随机访问(只有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]];
}
}
Run Code Online (Sandbox Code Playgroud)
由于体积小,我通过将表保存在__private内存空间中获得最佳性能.但是,由于访问查找表的随机性,仍然存在很大的性能损失.删除查找表代码(例如,用简单的算术运算代替),虽然内核会提供错误的答案,但性能提高了3倍以上.
有没有更好的办法?我是否忽略了一些OpenCL功能,它为非常小的内存块提供了有效的随机访问?使用矢量类型可以有效的解决方案吗?
[编辑]注意,X的最大值是7,但Y的最大值是2 ^ 32-1.换句话说,正在使用查找表的所有位,因此不能将其打包成较小的表示.
opencl ×1