如何使用 CUDA 生成随机排列答案

【问题标题】：How to generate random permutations with CUDA如何使用 CUDA 生成随机排列
【发布时间】：2012-09-21 03:56:22
【问题描述】：

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

Fisher-Yates shuffle 的顺序版本。

例子：

令 S={1, 2, ..., 7} 为源索引集。目标是并行生成 n 个随机排列。 n 个排列中的每一个都只包含每个源索引一次，例如{7, 6, ..., 1}。

【问题讨论】：

制作 X 个线程局部随机生成器并在每个生成器上运行 Fisher-Yates...？
这肯定行得通，但由于 SIMD 执行模型，这对于使用 CUDA 的实现来说也是最坏的情况。
你能解释一下最坏的情况吗？您的意思是不同的种子会使所有线程遵循不同的控制路径吗？你为什么这么认为？ F-Y shuffle 是一个简单的循环
使用thrust::permutation_iterator 怎么样？但是，它确实需要您编写自己的重新索引方案。
this 可能感兴趣

标签： c++ algorithm cuda thrust

【解决方案1】：

如果 s 的长度 = s_L，那么在推力中可以实现一个非常粗略的方法：

http://thrust.github.com.

首先，创建一个长度为 s_L x n 且重复 s n 次的向量 val。

创建一个向量 val_keys 将 n 个重复 s_L 次的唯一键与 val 的每个元素相关联，例如，

  val = {1,2,...,7,1,2,...,7,....,1,2,...7}
  val_keys = {0,0,0,0,0,0,0,1,1,1,1,1,1,2,2,2,...., n,n,n}

现在是有趣的部分。创建一个长度为 s_L x n 的均匀分布随机变量的向量

  U  = {0.24, 0.1, .... , 0.83}

然后你可以在 val,val_keys 上做 zip 迭代器并根据 U 对它们进行排序：

http://codeyarns.com/2011/04/04/thrust-zip_iterator/

val, val_keys 会到处都是，所以你必须使用推力::stable_sort_by_key() 将它们重新组合在一起，以确保如果 val[i] 和 val[j] 都属于 key[k ] 并且 val[i] 在随机排序之后在 val[j] 之前，那么在最终版本中 val[i] 仍然应该在 val[j] 之前。如果一切按计划进行，val_keys 应该看起来和以前一样，但 val 应该反映改组。

【讨论】：

【解决方案2】：

Fisher-Yates shuffle 可以并行化。例如，4 个并发工作人员只需要 3 次迭代即可对 8 个元素的向量进行混洗。在第一次迭代中，它们交换 01, 23, 45, 67;在第二次迭代 02, 13, 45, 67;在最后一次迭代 04, 15, 26, 37.

这可以很容易地实现为 CUDA __device__ 代码（受标准 min/max reduction 启发）：

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]。

请注意，上面的代码有以下假设：

排列的长度是 2 * BLOCK_SIZE，其中 BLOCK_SIZE 是 2 的幂。
2 * BLOCK_SIZE 整数适合__shared__ CUDA 设备的内存
BLOCK_SIZE 是 CUDA 块的有效大小（通常在 32 到 512 之间）

要生成多个排列，我建议使用不同的 CUDA 块。如果目标是排列 7 个元素（正如在原始问题中提到的那样），那么我相信在单线程中完成它会更快。

【讨论】：

【解决方案3】：

对于大型集合，在随机键向量上使用排序原语可能足以满足您的需求。首先，设置一些向量：

const int N = 65535;
thrust:device_vector<uint16_t> d_cards(N);
thrust:device_vector<uint16_t> d_keys(N);
thrust::sequence(d_cards.begin(), d_cards.end());

然后，每次你想洗牌时，调用一对：

thrust::tabulate(d_keys.begin(), d_keys.end(), PRNFunc(rand()*rand());
thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_cards.begin());
// d_cards now freshly shuffled

随机密钥是从使用种子（在主机代码中评估并在启动时复制到内核）和密钥编号（在线程创建时以表格形式传入的）的仿函数生成的：

struct PRNFunc
{
  uint32_t seed;
  PRNFunc(uint32_t s) { seed = s; }
  __device__ __host__ uint32_t operator()(uint32_t kn) const
  {
    thrust::minstd_rand randEng(seed);
    randEng.discard(kn);
    return randEnd();
  }
};

我发现如果我能弄清楚如何缓存thrust::sort_by_key 在内部执行的分配，性能可以提高（可能提高30%）。

欢迎任何更正或建议。

【讨论】：