如何使用 CUDA 生成随机排列
Posted
技术标签:
【中文标题】如何使用 CUDA 生成随机排列【英文标题】:How to generate random permutations with 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 可能感兴趣
【参考方案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%)。
欢迎任何更正或建议。
【讨论】:
以上是关于如何使用 CUDA 生成随机排列的主要内容,如果未能解决你的问题,请参考以下文章