[翻译]CURAND Libaray--Host API--(1)
原文来自:cuda curand toolkit document
Translated by xingoo
如果有错误请联系:xinghl90@gmail.com
2Host API简述
使用host api,用户需要在头文件的部分包含 curand.h,并且添加CURAND的动态链接库,即在LINKER的INPUT里面添加”curand.lib;”。这个文档是基于CUDA runtime的,所以用户的代码也应该是在runtime时调用的。而driver API是不支持CURAND的。
触发器生成随机数,CURAND同过内部的封装产生伪随机数列或者真随机数列。具体的步骤如下:
1 使用函数curandCreateGenerator创建一个新的目标类型(参考 触发器类型)触发器
2 设置触发器的选项(参考 触发器选项),比如,使用curandSetPseudoRandomGeneratorSeed() 来设置种子。
3 使用cudaMalloc()申请设备上的存储空间
4 通过curandGenerate()或者其他的生成函数 生成随机数
5 使用生成结果
6 如果希望生成更多的随机数,可以多次使用curandGenerate().
7 通过curandDestroyGenertator(),释放地址空间
如果希望在cpu的host端生成随机数,那么第一步应该调用curandCreateGeneratorHost(),;在第三步,要申请主机的一段地址空间来接收生成结果。其他步骤,在host端还是device端的操作都是相同的。
同一时间段创建几个触发器是可以的,每个触发器的封装都独立声明。每个触发器自己确定一组生成序列。如果每次运行的时候,设置相同的参数,生成随机序列也都是相同的。在device端生成的序列与在host端生成的也是相同的。
注意步骤4里面,curandGenerate()调用一段内核函数,并且异步返回。所以,如果你在不同的流里面运行了另一个内核函数,而这个内核函数又要使用curandGenerate()的结果,那么就必须调用cudaThreadSynchronized() 来实现同步,以确保生成随机序列的内核函数在被调用前结束执行。(也可以通过下面这种方法解决:use the stream management/event management routines 实在不知道这个方法怎么翻译)
通过一个host端的指针来调用在device上运行的触发器,以及通过一个device端的指针来调用运行在host端的触发器,这种情况都属于未定义。
2.1触发器类型Generator Types
随机数触发器通过curandCreateGenerator()来传递类型创建。在CURAND中有7种随机数触发器类型,我们把他们分为两类。
2.1.1 伪随机数触发器
CURAND_RNG_PSEUDO_XORWOW,(基于异或移位算法XORWOW)
CURAND_RNG_PSEUDO_MRG32K3A,(基于Combined Multiplie Recursive)
CURAND_RNG_PSEUDO_MTGP32(基于梅森旋转法Mersenne Twister)
2.1.2 真随机数触发器
这4种触发器都是基于SOBOL序列的变种的真随机数序列触发器
CURAND_RNG_QUASI_SOBOL32(是一种32位序列的SOBOL触发器)
CURAND_RNG_QUASI_SCRAMBLED_SOBOL32(srambled不知道怎么翻译才好)
CURAND_RNG_QUASI_SOBOL64(是一种64位序列的SOBOL触发器)
CURAND_RNG_QUASI_SCRAMBLED_SOBOL64
2.2触发器选项 Generator Options
创建时,随机数触发器就能通过基本选项---seed种子,offset消耗和order顺序来定义。
2.2.1 seed 种子
种子是一个64位的整型数,用来初始化伪随机触发器的产生。相同的种子,能够产生相同的序列。
2.2.2 Offset消耗/补偿
这个选项用来跳过序列开始的一段随机数。如果offset=100,那么第一个随机数产生的时候,会从原来的随机数序列的第100个开始。这让多次运行同一段程序可以继续使用相同序列产生的随机数,而不会重叠。这个跳过的功能不支持CURAND_RNG_PSEUDO_MTGP32触发器。
2.2.4 Order顺序
这个选项用来选择在全局内存上生成的随机数如何排序。
有三种伪随机序列顺序可以选择:
CURAND_ORDERING_PSEUDO_DEFAULT(默认的)
CURAND_ORDERING_PSEUDO_BEST
CURAND_ORDERING_PSEUDO_SEEDED.
有一种真随机序列顺序可以选择:
CURAND_ORDERING_DEFAULT.
目前,CURAND_ORDERING_PSEUDO_DEFAULT 和 CURAND_ORDERING_PSEUDO_BEST对于所有的伪随机触发器产生的结果都是相同的。然而,未来的版本,CURAND可能会调整CURAND_ORDERING_PSEUDO_BEST,使之无论是性能还是质量上都得以改善。而CURAND_ORDERING_PSEUDO_DEFAULT对于任何的CURAND版本都会保持相同的产生结果。在目前的版本中,只有XORWOW触发器拥有多种产生顺序。
关于产生顺序,每种触发器的做法如下:
XORWOW伪随机触发器:
CURAND_ORDERING_PSEUDO_BEST
目前版本同CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序是相同的
CURAND_ORDERING_PSEUDO_DEFAULT
当offset=n时,随机序列在全局内存中从(n mpd 4096)-2 67+[n/4096]开始
CURAND_ORDERING_PSEUDO_SEEDED
当offset=n时,随机序列在全局内存中从n mod 4096开始。也就是说,每4096个线程使用不同的种子。这种种子的方式,虽然节省了运行时间,但对于某些种子的值,伪随机序列输出结果统计也显示出了一些不足。
MRG32K3A 伪随机触发器
CURAND_ORDERING_PSEUDO_BEST
跟下面的DEFAULT一样
CURAND_ORDERING_PSEUDO_DEFAULT
当offset=n时,公式为(n mod 4096) -276+[n/4096](注意这个触发器产生的跨度跟XORWOW产生的跨度是不一样的)
MTGP32 伪随机触发器
CURAND_ORDERING_PSEUDO_BEST
同样跟DEFAULT的顺序相同
CURAND_ORDERING_PSEUDO_DEFAULT
MTGP32触发器基于基本算法,通过设置不同参数,能够产生64种随机序列。使用S(p)来定义序列,p为参数。Offset=n时,公式为S([n/256]mod 64).也就是说,S(0)会有256个例子,然后是S(1)的256个,一直到S(63),一直重复。
Philox_4x32_10 伪随机触发器
CURAND_ORDERING_PSEUDO_BEST
跟下面的DEFAULT一样
CURAND_ORDERING_PSEUDO_DEFAULT
每个触发器的线程都会产生不同的随机序列。在host API中有8192种不同的序列。来自于一个序列的每四个值都依赖于下一个序列的四个值(Each four values from one sequence are followed by four values from next sequence)。
32位和64位的SOBOL以及快速SOBOL真随机触发器
CURAND_ORDERING_QUASI_DEFAULT
当在d个dimensions上面生成n个结果,第1个dimesion由n/d个结果组成,第2个也由n/d个组成,一直到第d个dimension.触发器只能生成dimension的倍数大小的序列。这个dimension的参数d使用curandSetQuasiRandomGeneratorDimensions(),并且默认为1.