c++ - 如何使用 CUDA 生成随机排列

Question

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

一个顺序版本将是 Fisher-Yates 洗牌。

例子：

令 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__代码（受标准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].

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

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

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

Answer 2

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

http://thrust.github.com。

首先，创建一个长度为 s_L xn 且重复 sn 次的向量 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 xn 的向量均匀分布的随机变量

  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 应该反映改组。

Answer 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());

然后，每次你想洗牌 d_cards 调用这对：

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%）。

欢迎任何更正或建议。