前言

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/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中,有4中为随机算法,4种真随机算法

伪随机算法     对应的状态

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  生成4个服从泊松分布的int

curand_normal  生成服从正态分布的float

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

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

curand_normal_double  生成服从正态分布的double

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

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

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

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

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

curand_log_normal2_double 生成2个服从对数正态分布的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;
}

更多推荐

cuda随机数生成