如何使用CUDA生成随机排列

Question

我可以使用哪些并行算法从给定集合生成随机排列？特别是适合CUDA的论文提案或链接会有所帮助.

这种顺序版本将是Fisher-Yates shuffle.

例:

设S = {1,2,...,7}是源索引的集合.目标是并行生成n个随机排列.n个排列中的每一个恰好包含每个源索引一次,例如{7,6,...,1}.

Answer 1

Fisher-Yates shuffle可以并行化.例如,4个并发工作者只需要3次迭代来混洗8个元素的向量.在第一次迭代时,它们交换0 < - > 1,2 < - > 3,4 < - > 5,6 - 7; 在第二次迭代0 < - > 2,1 < - > 3,4 - 5,6 - 7; 并且在最后一次迭代0 - 4,1 - 5,2 - 6,3 - 7.

这可以很容易地实现为CUDA __device__代码(受标准最小/最大缩减的启发):

const int id  = threadIdx.x;
__shared__ int perm_shared[2 * BLOCK_SIZE];
perm_shared[2 * id]     = 2 * id;
perm_shared[2 * id + 1] = 2 * id + 1;
__syncthreads();

unsigned int shift = 1;
unsigned int pos = id * 2;  
while(shift <= BLOCK_SIZE)
{
    if (curand(&curand_state) & 1) swap(perm_shared, pos, pos + shift);
    shift = shift << 1;
    pos = (pos & ~shift) | ((pos & shift) >> 1);
    __syncthreads();
}

这里省略了curand初始化代码,并且方法swap(int *p, int i, int j)交换值p[i]和p[j].

请注意,上面的代码有以下假设:

1. 置换长度为2*BLOCK_SIZE,其中BLOCK_SIZE为2的幂.
2. 2*BLOCK_SIZE整数适合__shared__CUDA设备的内存
3. BLOCK_SIZE是CUDA块的有效大小(通常在32到512之间)

为了生成多个排列,我建议使用不同的CUDA块.如果目标是对7个元素进行排列(正如在原始问题中提到的那样),那么我相信在单个线程中执行它会更快.