前言

host端当然可以用cstdlib中的rand()函数生成随机数，但设备端如何使用这些随机数？每调用一次rand()，就通过cudaMemcpy传递给显存吗？显然不是，这会消耗太多I/O时间；先调用n次，一次性传到device中吗？虽然可行，但并不是最优解。能否用一种方法，让主机仅仅传给设备一个信号，使得多个随机数在device端被生成，这样就能避免过多的I/O带来的耗时问题；或者调用一个设备端函数，使得设备端可以根据需要动态生成随机数。于是curand库和curand_kernel库进入了我们的视野。

参考文献

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() 类型

//生成伪随机序列，返回的随机数是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端随机数生成

步骤

使用curand的deviceAPI生成随机数主要需要三个步骤

a．创建一个随机算法状态的对象，如curandState state

b．对状态进行初始化，使用curand_init（）函数

c．生成随机数，使用curand等

参考文献

a步：创建随机算法状态对象

在deviceAPI中，有４中为随机算法，４种真随机算法

伪随机算法　　　　　对应的状态

Mtgp32 curandStateMtgp32_t

MRG32k3a curandStateMRG32k3a_t

Philox4_32_10 curandStatePhilox4_32_10_t

XORWOW curandStateWORWOW_t

真随机算法　　　　　对应的状态

scrambled Sobol64 curandStateScrambledSobol64_t

Sobol64　 curandStateSobol64_t

scrambled Sobol32 curandStateScrambledSobol32_t

Sobol32　 curandStateSobol32_t

如要创建算法XORWOW状态对象：curandStateWORWOW_t state

b步：curand_init（）

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

c步：生成随机数

生成随机数除了可以使用不同的算法（由参数决定）以外，还能生成不同分布、不同浮点类型的随机数，这根据调用的生成随机数的API决定。

主要有以下几种API：

curand 生成unsigned int型整数，没有特殊分布

curand_uniform 生成服从均匀分布的float

curand_uniform4 生成4个服从均匀分布的float

curand_uniform_double　 生成服从均匀分布的double

curand_uniform2_double 生成2个服从均匀分布的double

curand_poisson　 生成服从泊松分布的int

curand_poisson4　 生成４个服从泊松分布的int

curand_normal　 生成服从正态分布的float

curand_normal2　 生成２个服从正态分布的float

curand_normal4　 生成４个服从正态分布的float

curand_normal_double　 生成服从正态分布的double

curand_normal2_double　 生成２个服从正态分布的double

curand_log_normal　 生成服从对数正态分布的float

curand_log_normal2 生成2个服从对数正态分布的float

curand_log_normal4 生成４个服从对数正态分布的float

curand_log_normal_double 生成服从对数正态分布的double

curand_log_normal2_double 生成２个服从对数正态分布的double

curand_mtgp32_single　 生成mtgp32算法中服从均匀分布的float

curand_mtgp32_single_specific 生成mtgp32算法中服从特殊泊松分布的float

curand_mtgp32_specific 生成mtgp32算法中的32bits的int

如要生成一个int型的随机数，采用XORWOW算法：curand(curandStateXORWOW_t* state)

unsigned int x = curand(&localState);

float x = curand_uniform(&localState);

float2 x = curand_normal2(&localState);

double x = curand_uniform_double(&localState);

double x = curand_normal_double(&localState);

一个例子

1. 创建cuda_runtime的项目文件后，

2. 转到"属性"->" CUDA"，然后将" CUDA Toolkit自定义目录"设置为CUDA工具包目录。
   对我来说是：C：\ Program Files \ nVIDIA GPU Computing Toolkit \ CUDA \ v10.0

3. 将以下代码复制进去

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "curand_kernel.h"

#include <stdio.h>

#define BLOCKS 32
#define THREAD_NUM 6

__global__ void setup_kernel(curandState *state, unsigned long seed)
{
	int tid = blockIdx.x *blockDim.x + threadIdx.x; //获取线程号0~blocks*THREAD_NUM-1  grid划分成1维，block划分为1维
	curand_init(seed, tid, 0, &state[tid]);// initialize the state
}

__global__ void use(curandState *globalState)
{
	unsigned int j;
	int tid = blockIdx.x *blockDim.x + threadIdx.x; //获取线程号0~blocks*THREAD_NUM-1  grid划分成1维，block划分为1维
	curandState localState = globalState[tid];
	j = (curand(&localState));
	printf("%u\n", j);
}

int main()
{
	curandState* devStates;  //创建一个随机算法状态的对象
	cudaMalloc(&devStates, BLOCKS * THREAD_NUM * sizeof(curandState));
	srand(time(0));

	setup_kernel << <BLOCKS, THREAD_NUM >> > (devStates, rand()); // blocks number is POP. thread number is P
	use << < BLOCKS, THREAD_NUM >> > (devStates);
    
    return 0;
}