如何通过置换实现OpenCL/CUDA索引的随机化？

我在OpenCL中处理尺寸为整数vdimx、vdimy、vdimz的float体数据，Kernel开头通常编写如下代码：

int i = get_global_id(0);
int j = get_global_id(1);
int k = get_global_id(2);

OpenCL运行时分配的i、j、k在同一时刻数值通常非常接近，比如当vdimx=vdimy=vdimz=1024时，最先运行的索引大概率是(0,0,0)、(0,0,1)、(0,0,2)这类连续值，而非随机分散的索引。但在CT投影器这类应用中，邻近体素索引会投影到同一像素，导致多GPU线程写入同一内存位置，拖慢原子操作的速度。

我想通过随机化体素位置解决这个问题，计划修改代码为：

int i = get_global_id(0);
int j = get_global_id(1);
int k = get_global_id(2);
i = perm(i);
j = perm(j);
k = perm(k);

核心问题解答 #

1. #define perm(i) ((i * 10) % vdimx)能否对任意vdimx实现一一映射？ #

不能。这种线性同余方式要实现一一映射（即置换），必须满足乘数与模数vdimx互质——也就是两者的最大公约数gcd(乘数, vdimx) = 1。比如当vdimx是偶数时，10和vdimx的gcd至少是2，此时不同的i会映射到同一个结果，无法实现一一对应。

2. 使用不整除vdimx的质数进行乘法，能否得到置换？ #

不一定。关键不是质数是否整除vdimx，而是质数和vdimx是否互质。比如vdimx=15时，质数3能整除15，gcd(3,15)=3≠1，用3作为乘数会出现重复映射；但用质数7时，gcd(7,15)=1，就能实现一一映射。换句话说，只要乘数是与vdimx互质的整数（不管是不是质数），都能通过(i * multiplier) % vdimx实现置换。

3. 其他高效的索引随机化技术 #

• 线性同余置换优化：除乘法外结合加法，比如perm(i) = ((i * a + b) % vdimx + vdimx) % vdimx，其中a与vdimx互质，b为任意整数。这种方式能进一步打乱索引分布，同时保持一一映射。
• 轻量哈希置换：使用低开销的哈希函数，比如perm(i) = (i * 1103515245 + 12345) >> 16 % vdimx（类似glibc的rand实现），需确保结果在[0, vdimx-1]范围内且无冲突，适合对随机性要求更高的场景。
• 预计算置换表：若vdimx、vdimy、vdimz固定，可在主机端用Fisher-Yates洗牌算法生成0到vdimx-1的随机排列数组，传入OpenCL Kernel后，perm(i)直接查表即可。这种方式随机性最好，Kernel端计算开销极低，但需额外内存存储置换表，适合维度固定且内存充足的场景。
• 位运算置换：针对2的幂次维度（如1024=2^10），可使用位混洗操作，比如perm(i) = (i << 5) | (i >> 7)（针对10位的i），利用位操作的高效性打乱索引，且天然是一一映射（位操作可逆）。

内容的提问来源于stack exchange，提问作者VojtaK