CUDA：随机数

本文主要参考文档：http://docs.nvidia.com/cuda/curand/introduction.html#introduction
示例代码请见：https://github.com/Yannnnnnnnnnnn/cuda_random
首先声明一下，由于我完全不懂伪随机中的具体算法，也不懂准随机是什么。所以此处并不介绍这两种算法的区别，并且为了简单起见，我只介绍伪随机及其用法。
本博客的主要重点是大概介绍一下CUDA中随机数怎么用，并且大概描述下怎么做性能更高。

---

总体介绍

CUDA总共提供两种类型的随机数生成，一种是位于host（CPU）端的随机数，一种是位于device（GPU）端的随机数。
- 对于host端的随机数生成，需要/include/curand.h头文件，还需要curand依赖库，（如何添加请看：http://blog.csdn.net/u012348774/article/details/78901205）。不过host端随机数不一定要在host端生成，也可以指定在device端生成。当指定在host端生成随机数时，程序调用CPU完成所有的随机数，并将所有的结果保存在host端；而当制定在device端生成随机数时，程序的调用由CPU完成，程序的执行却是在device端，最终生成的随机数也保存着在device的全局内存中。
- 对于device端的随机数生成，则需要/include/curand_kernel.h。此时随机数生成也是使用GPU生成，由于此时的代码可以直接编写在device或者global函数中，生成的随机数可以直接被核函数调用。

---

host端随机数生成

host端device模式的随机数生成大致分为以下几个步骤：
- 使用curandCreateGenerator()选择一个合适的随机数生成方式，初始化随机数生成器；
- 设置随机数生成器的一些参数，例如种子点、偏移量等等；值得注意的是，此处的seed如果是一样的，那么每次生成的随机数都会是一样的。因此要确保生成的随机数不重复，最好设置独一无二的seed，例如当前时间。
- 使用cudaMalloc()在device端声明内存；
- 使用curandGenerate() 生成随机数并保存用于使用。

host端host模式的随机数生成大致分为以下几个步骤：
- 使用curandCreateGeneratorHost()选择一个合适的随机数生成方式，初始化随机数生成器；
- 设置随机数生成器的一些参数，例如种子点、偏移量等等；值得注意的是，此处的seed如果是一样的，那么每次生成的随机数都会是一样的。因此要确保生成的随机数不重复，最好设置独一无二的seed，例如当前时间。
- 使用malloc()在host端声明内存；
- 使用curandGenerate() 生成随机数并保存用于使用。

生成函数
上文只提到了一个随机数生成函数curandGenerate，其实cuda还支持好几种随机数生成方式，其中部分如下（更详细的介绍请参考生成函数）。同时还要注意，一个伪随机序列的长度是有限的，即不改变随机数生成器大量调用会得到重复的随机序列。

//生成伪随机序列，返回的随机数是32-bit的无符号整形
curandGenerate( curandGenerator_t generator, unsigned int *outputPtr, size_t num)

//生成64bit的伪随机序列
curandGenerateLongLong( curandGenerator_t generator, unsigned long long *outputPtr, size_t num)

//生产0-1之间的单精度浮点型伪随机序列
curandGenerateUniform(
    curandGenerator_t generator, 
    float *outputPtr, size_t num)

//按照指定的均值和方差生成伪随机序列
curandGenerateNormal( curandGenerator_t generator, float *outputPtr, size_t n, float mean, float stddev)

性能
除开各种算法之间的速度差异，通常情况下调用随机数生成器的次数越少，程序的效率也越高。因此通常做法都是调用一次随机数生成器，然后生成一定数量的随机数。而不是一个随机数生成器只生成一个随机数。

---

device端随机数生成

相比而言，device端就简单多了，只需要调用两个类型的函数：

1. curand_init，初始化随机数生成器（MTGP32 generator、Philox_4x32_10 generator不介绍）；
2. 使用curand函数生成随机数。

curand_init
为了更好的理解，有必要对该函数做进一步介绍。

__device__ void curand_init ( unsigned long long seed, unsigned long long sequence, unsigned long long offset, curandState_t *state)

毋庸置疑，相同的seed、sequence总是会生成相同的伪随机数，并且在调用2 67 ⋅ sequence + offset次curand函数后，也会生成相同的伪随机数。
另外补充，不同seed生成的伪随机序列通常是不想关的，但是也存在相关的可能。同一个seed不同sequence生成的伪随机序列一定不相关。
针对并行计算，CUDA给出了一些建议：每个核函数可以使用同一个seed，但是最好每个线程都是不同的sequence。每个核函数一个seed和sequence，会导致程序效率异常低下。

生成函数
device端的随机数也支持好几个不同类型的生成函数，其中我觉得比较重要的有以下几个（更多请参考生成函数）：

//返回整数
__device__ unsigned int
curand (curandState_t *state)

//返回包括[0,1]之间的伪随机数单精度浮点序列，包括0和1
__device__ float 
curand_uniform (curandState_t *state)

//返回均值为0，标准差为1的伪随机数单精度浮点序列
__device__ float 
curand_normal (curandState_t *state)

//双精度
__device__ double 
curand_uniform_double (curandState_t *state)

//双精度
__device__ double 
curand_normal_double (curandState_t *state)

性能
前文已经提到，curand_init()的速度非常慢，因此不建议在每个线程中调用该函数；并且curand_init中的offset参数越大其速度也越慢。
因此给出的建议是最好不要每个线程只生成一个随机数，这样会导致程序的效率非常低下。
在我提供的demo中，random.cu是每个线程一个curand_init()，random_gen.cu是一系列线程调用一个curand_init()，大家可以感受下速度差异。