随机数的生成在应用科学、密码学和金融应用领域中都有很广泛的用途。一个随机数生成器(RNG)是一个没有任何参数的函数f，但是每次调用时都能返回随机值序列的下一个值。我们可以把它看作一个指针，该指针可以遍历随机数构成的数组，如下图所示：
cuRAND库在基于CUDA库的拟随机数和伪随机数的生成。接下来主要介绍随机数生成的一些背景、cuRAND的配置以及使用cuRAND的两个例子。
拟 随 机 数 或 伪 随 机 数 的 选 择
对于以计算机为基础的随机数生成的学习，需要明白的是没有真正的随机数生成。因为计算机是一个有序系统(以保证函数正常运行)，所以从根本上来说就没有能描绘随机序列的概念，也就不能生成真正的随机数。某些硬件解决方案可以生成被认为是真正的随机数，但是许多库用来生成随机数的RNG算法也都有很好的算法结构和定义，它能让用户有了生成真正随机数的错觉。然而，不能说这是一个不好的特性，虽然它在某些情况下还是有用的。
例如，RNG是从一个种子数据开始生成随机数的，这个种子是生成随机数序列的初始值。我们可以把它当作上图所示的最开始的第一个值，其他所有制都是在它的基础上生成的。我们可以每次都给一个定义好的RNG算法提供相同的初始值，而每次都会得到相同的随机数序列，这对于测试应用程序很有用，可以重复使用相同的随机序列。
RNG可以分为两大类：伪随机数生成器和拟随机数生成器。两者各有各的用途并且都被cuRAND所支持。一个伪随机数生成器(PRNG)使用RNG算法生成随机数序列。在这个序列中，每个值都是有效范围内的任意值，并且都是RNG所用的数据存储类型。例如，当从一个整数类型的PRNG中取回一个值时，返回值为1的概率P(1)和返回值为2的概率P(2)或是3的概率P(3)以及所有的P(INT_MAX)都相等。这对于PRNG中返回的每一个值都成立。这就意味着，不能仅仅因为上一个取回值是2就认为下一个取值是2的概率就会变小。换句话说，对一个伪随机数序列，每次采样都是独立统计事件，对以后抽样的样本观察值不会有影响。但是，这在拟随机数生成器(QRNG)中是不成立的。一个QRNG会尽量均匀地填充输出类型的范围。因此，如果QRNG采样的前一个值为2，那么下一个值是2的概率P(2)实际上会减小。一个QRNG的序列采样不是独立统计的。
PRNG和QRNG在不同的应用程序中发挥着不同的作用。当需要真正的随机时，PRNG是更好的选择。对于一个使用密码生成的应用程序，PRNG是一个比QRNG更好的选择，因为若是使用PRNG的话，已经生成的密码信息对同一序列内其他密码的生成概率没有影响。另一方面，QRNG在探索不为人知的空间时是很有用的。QRNG能保证多维空间有更均匀的采样，也有可能发现固定采样间隔所没有发现的一些性质。例如，采用蒙特卡罗方法的某些应用程序就得益于QRNG的使用。
cuRAND 库 概 述
cuRAND库可以用于伪随机序列和拟随机序列的采样。它既有一个主机端API又有一个设备端API，这意味着它可以直接被主机端调用，也可以直接被内核代码调用。cuRAND库的许多概念对于这两种API来说是共享的，但当使用设备端API时增加了一些选项。
概念介绍
主机和设备cuRAND API的配置项有4个：用于生成随机序列的RNG算法、返回值遵循的一个分布、初始种子数值和随机数序列开始采样的一个偏移量。由于设备端API指定每个参数都要依照规范进行设置，如果用户没有给定初始值，那么主机端API会自动将其设置为默认值。
主机和设备端也都有其他CUDA库里的句柄概念。在主机端API中，句柄即所谓的随机数生成器。我们可以用curandCreateGenerator或curandSetGeneratorOffset。只需有一个生成器来访问主机端API。而在设备端API中，句柄则是指cuRAND的状态。设备状态有很多种类型，每一种对应设备端API支持不同的RNG。然而，状态对象的作用仍然是维护GPU上单线程cuRAND上下文的配置和状态。因此，通常需要分配很多的设备状态对象，每一个对应于不同的GPU线程。
在cuRAND主机和设备端API中第一个配置项是RNG算法。在主机端API中，使用如下配置：
curandStatus_t curandCreateGenerator(curandGenerator_t *generator, curandRngType_t rng_type);
在设备端API，这是通过调用一个特定的RNG初始化函数来配置的，该函数包含RNG特定的状态对象，该对象在设备端API中被当成一个cuRAND生成器：
__device__ void curand_init(unsigned long long seed,unsigned long long subsequence,unsigned long long offset,curandStateXORWOW_t *state);
__device__ void curand_init(unsigned long long seed,unsigned long long subsequence,unsigned long long offset,curandStateMRG32k3a_t *state);
选择不同的RNG算法会影响随机数生成技术，并可能影响生成序列的随机性和算法性能，主机端和设备端都支持多种随机数生成算法。例如，可以使用以下命令在设备端初始化RNG，使用的是XORWOW RNG算法：
__global__ void kernel(...){
curandStateXORWOW_t rand_state;
curand_init(0,0,0,&rand_state);
然后，使用rand_state作为生成随机数的句柄。接下来，不同分布方式的选择也会影响到由RNG生成的随机数在一定取值范围内的分布(注意，即使是浮点类型的值也是在一个离散的有限范围内的)。一个RNG算法和所采用的分布方式之间的关系可能不太明显。RNG算法可以被看成是一个能产生二进制位的随机序列的黑盒。这些二进制位对RNG算法来收并没有实际意义。在RNG之上添加一个指向cuRAND的特定返回类型，然后用这些二进制来生成一个能表示所选择分布的特征数值。主机和设备API都支持正态分布、均匀分布、对数正态分布和泊松分布。当通过调用特定的分布函数来生成随机数时，所需的分布也就确定了。在主机端API中若使用均匀分布需要使用curandGenerateUniform:
curandGenerator_t rand_state;
int d_rand_length = ...;
float *d_rand;
curandCreateGenerator(&rand_state,CURAND_RNG_PSEUDO_DEFAULT);
cudaMalloc((void **)&d_rand,sizeof(float) * d_rand_length);
curandGenerateUniform(rand_state,d_rand,d_rand_length);
//在设备端，按照如下方式使用curand_uniform来代替：
__global__ void kernel(...){
curandStateXORWOW_t rand_state;
curand_init(0,0,0,&rand_state);
float f = curand_uniform(&rand_state);
}
cuRAND中第三个配置项是种子值。种子的概念出现在PRNG和QRNG中，但在两者中用法不同。为cuRAND PRNG选择种子的值是64位的，由人随机指定，为PRNG接下来生成随机序列打下基础。不同种子会产生不同的随机序列。主机端API允许使用curandSetPseudoRandomGeneratorSeed为PRNG选取种子：
curandGenerator_t rand_state;
curandCreateGenerator(&rand_state,CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(rand_state,9872349ULL);
如果没有特殊情况，那么会使用默认的种子值；另一方面，对于每个线程的PRNG。设备端API都要设置好确定的种子值。如下调用curand_init来选取种子，第一个参数为创建RNG指定的起始种子：
__device__ void curand_init(unsigned long long seed,unsigned long long subsequence,unsigned long long offset,curandStateXORWOW_t *state);
唯一被主机和设备端API支持的QRNG是基于Sobol拟随机序列的。一个Sobol序列以方向向量作为种子。回忆前面，一个拟随机数生成器的优点是每个采样都不是一个独立的统计事件，一个QRNG会特意在取值范围内均匀地取数。我们可以把这些方向向量看成是探索n维空间的起始方向，随机数据位就是从这个n维空间产生的。因此，所谓的种子，就是人为指定的初始的随机性，即使有时名字不同，但意义是一样的。在cuRAND主机端API，只有用于QRNG的维数可以使用curandSetQuasiRandomGeneratorDimensions来设置：
curandGenerator_t rand_state;
curandCreateGenerator(&rand_state,CURAND_RNG_QUASI_SOBOL32);
curandSetQuasiRandomGeneratorDimensions(rand_state,2);
//而设备端API允许指定种子的方向向量
__global__ void kernel(curandDirectionVectors32_t *direction_vector,...){
curandStateSobol32_t rand_state;
curand_init(*direction_vector,0,&rand_state);
}
curandDirectionVectors32_t *h_vectors;
curandGetDirectionVectors32(&h_vectors,CURAND_DIRECTION_VECTORS_32_JOEKUO6);
cudaMemcpy(d_vectors,h_vectors,sizeof(curandDirectionVectors32_t),cudaMemcpyHostToDevice);
kernel<<<blocks,threads>>>(d_vectors,...);
最后，cuRAND中第四个配置项是随机数序列在起始点的偏移量。也就是说，由不同种子起始生成的随机数序列也不同。这个偏移量允许你跳转到当前序列的第i个随机值上。在主机API中，可以使用curandSetGeneratorOffset来进行设置：
curandGenerator_t rand_state;
curandCreateGenerator(&rand_state,CURAND_RNG_PSEUDO_DEFAULT);
curandSetGeneratorOffset(rand_state,0ULL);
//在cuRAND设备端API中，偏移量为curand_init函数的一个参数(与种子类似)：
__device__ void curand_init(unsigned long long seed,unsigned long long subsequence,unsigned long long offset,curandStateXORWOW_t *state);
主机和设备端API对比
对一个特定应用程序来说，需要对是否使用主机或设备的cuRand API做出决定。两个API提供了相似的功能：相似的RNG，相同的分布，以及相同的可配置性(参数不同)。但是对这些功能进行访问的方法却大有不同。接下来将根据应用程序的需求，对选择过程提供给指南。
如果我们唯一的目标是在主机端应用程序中生成高效、高质量的随机数，那么主机端API是最好的选择。CUDA方面的专家已经着手编写相关程序，并达到了此应用目的的最佳性能，这远比我们自己编写内核程序来调用设备端API要好用的多。如果考虑到之后GPU内核的消耗，由于使用主机端API在GPU上预生成随机数，因此这并没有什么优势。按提供者与使用者对随机性进行的划分可能会导致代码可读性下降，也可能会因cuRAND主机函数库和内核启动开销的升高而导致性能变差，同时会要求在内核执行之前，就要知道必要的随机数数目。
如果我们想要对随机数生成有更多的控制权，如果我们正使用一组由手写CUDA内核生成的随机数，特别是内核中所需要的随机数是动态变化的，那么设备端API是正确的选择。只需要编写少量程序来初始化和管理设备端RNG的状态，就可以在CUDA程序进行内部操作时获得更大的灵活性。
cuRAND 介 绍
接下来主要介绍两个cuRAND API的例子。第一个是使用主机和设备端cuRAND API来替换系统自带的rand函数。第二个是使用主机和设备端cuRAND API自带的rand调用来创建一个CUDA内核随机数。
替换rand()
可以从Wrox.com中下载replace-rand.cu进行学习，此程序通过使用主机和设备端cuRAND API的调用来为之后的主机耗能生成随机数。这个例子会在依次函数库调用中产生大量的随机数，然后在主机请求新的随机数时对它们进行访问。只有当现有的随机数都用过了，才回到cuRAND函数库调用产生更多简单的随机数。
使用主机端API的工作流在本章中是比较简单的，其步骤如下：
1.使用curandCreateGenerator创建一个由所需RNG配置的cuRAND生成器。对这样的生成器进行配置是可行的(例如，使用curandSetStream)，但对于主机端API则是可选择的；
2.使用cudaMalloc为cuRAND预分配设备内存，使其用来存储输出的随机数；
3.通过执行一个cuRAND库调用来生成随机数，例如，curandGenerateUniform。
4.如果主机端一定要有能耗的化，则使用cudaMemcpy从设备内存中取回生成的随机数。
以下代码段是使用cuRAND主机端API的工作流：
float cuda_host_rand()
{
static float *dRand = NULL;
static float *hRand = NULL;
curandGenerator_t randGen;
static int dRand_length = 1000000;
static int dRand_used = 1000000;
if (dRand == NULL)
{
/*
* If the cuRAND state hasn't been initialized yet, construct a cuRAND
* host generator and pre-allocate memory to store the generated random
* values in.
*/
curandCreateGenerator(&randGen,CURAND_RNG_PSEUDO_DEFAULT);
cudaMalloc((void **)&dRand, sizeof(float) * dRand_length);
hRand = (float *)malloc(sizeof(float) * dRand_length);
}
if (dRand_used == dRand_length)
{
/*
* If all pre-generated random numbers have been consumed, regenerate a
* new batch using curandGenerateUniform.
*/
curandGenerateUniform(randGen, dRand, dRand_length);
cudaMemcpy(hRand, dRand, sizeof(float) * dRand_length,
cudaMemcpyDeviceToHost);
dRand_used = 0;
}
// Return the next pre-generated random number
return hRand[dRand_used++];
}
对设备端API的处理稍微有些复杂，有如下几步：
1.在设备内存中为每个线程预分配一套cuRAND状态对象来管理其RNG(随机数生成器)的状态；
2.如果cuRAND产生的随机值被规定为复制到主机或必须存储至后期的内核，那么来存储它们的预分配设备内存可选；
3.在设备内存中使用CUDA内核调用初始化所有cuRAND状态对象的状态；
4.执行一个CUDA内核，该内核调用一个cuRAND设备端函数(如curand_uniform)，然后使用预分配的cuRAND状态对象生成随机数。这一步和上一步可以合并到一个核函数中，但需要注意的是，在获取之后要用随机数时，不用重新初始化状态对象；
5.如果在第2步预分配了用于取回随机数的设备内存，那么就要把随机数传回主机端。
下面是关于cuRAND设备端API的工作流代码段：
/*
* An implementation of rand() that uses the cuRAND device API.
*/
float cuda_device_rand()
{
static curandState *states = NULL;
static float *dRand = NULL;
static float *hRand = NULL;
static int dRand_length = 1000000;
static int dRand_used = 1000000;
int threads_per_block = 256;
int blocks_per_grid = 30;
if (dRand == NULL)
{
/*
* If the cuRAND state hasn't been initialized yet, pre-allocate memory
* to store the generated random values in as well as the cuRAND device
* state objects.
*/
cudaMalloc((void **)&dRand, sizeof(float) * dRand_length);
cudaMalloc((void **)&states, sizeof(curandState) *
threads_per_block * blocks_per_grid);
hRand = (float *)malloc(sizeof(float) * dRand_length);
// Initialize states on the device
initialize_state<<<blocks_per_grid, threads_per_block>>>(states);
}
if (dRand_used == dRand_length)
{
/*
* If all pre-generated random numbers have been consumed, regenerate a
* new batch.
*/
refill_randoms<<<blocks_per_grid, threads_per_block>>>(dRand,
dRand_length, states);
cudaMemcpy(hRand, dRand, sizeof(float) * dRand_length,
cudaMemcpyDeviceToHost);
dRand_used = 0;
}
// Return the next pre-generated random number
return hRand[dRand_used++];
}
cuda_host_rand采用cuRAND的主机端API预生成一个大批量的随机数，然后在连续调用cuda_host_rand过程中对这些随机数进行遍历。它会记下预生成的随机数的数量和已使用的随机数的数量，然后根据这些统计数量来确定什么时候产生新的批处理。cuda_device_rand使用cuRAND设备端API,通过在设备上预分配一系列cuRAND状态对象来执行相同的操作，使用这些对象来生成大批量的随机数，并且采用与cuda_host_rand相同的方式来管理数据。
为CUDA内核生成随机数
我们可以在Wrox.com中的rand-kernel.cu中下载示例代码来进行学习，这个例子是利用cuRAND主机和设备端API来生成供CUDA内核使用的随机数。
use_host_api在调用一个使用这些随机数的CUDA内核之前，使用cuRAND主机端API预生成N个随机数。注意，这需要来自主机程序和cuRAND的多个内核调用，以及专门分配来存储所生成的随机数的设备内存。下面代码段的意义是，使用curandGenerateUniform生成随机数并传给dRand,将这些值传给host_api_kernel，然后将host_api_kernel的最终输出从设备端(dOut)传到主机端(hOut):
curandGenerateUniform(randGen, dRand, N);
// Consume the values generated by curandGenerateUniform
host_api_kernel<<<blocks_per_grid, threads_per_block>>>(dRand, dOut, N);
// Retrieve outputs
cudaMemcpy(hOut, dOut, sizeof(float) * N, cudaMemcpyDeviceToHost);
use_device_api函数使用cuRAND设备端API在GPU上按需生成随机数。注意，只需要一个包含所有cuRAND初始化和执行的单一的内核调用，无需分配用于存储这些随机值的CUDA设备内存。CUDA内核能直接使用生成的任意随机数。下面代码段的意义是，内核device_api_kernel在预分配的cuRAND设备生成器上的执行，以及将内核的最终输出从设备传到主机。
device_api_kernel<<<blocks_per_grid, threads_per_block>>>(states, dOut, N);
// Retrieve the results
cudaMemcpy(hOut, dOut, sizeof(float) * N, cudaMemcpyDeviceToHost);
cuRAND 发 展 中 的 重 要 主 题
cuRAND是一个使用起来简单而灵活的API。基于cuRAND的对象中最重要的部分就是了解随机性的要求。例如，不同的随机数生成器或不同的分布选择对应用程序的性能、正确性和结果会有很大的影响。特别是对像蒙特卡罗模拟这样依赖于随机数的应用程序。还有一点比较重要的是，要确保我们已经正确配置了cuRAND环境以生成期望的随机数类型。
扫一扫
407
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
