cuRAND :: CUDA Toolkit Documentation

2. 主机API概述

要使用主机API，用户代码应包含库头文件curand.h并动态链接cuRAND库。该库使用CUDA运行时，因此在使用静态cuRAND库时，用户也需要链接CUDA运行时。

随机数由生成器产生。cuRAND中的生成器封装了生成伪随机或准随机数序列所需的所有内部状态。通常的操作顺序如下：

1. 使用curandCreateGenerator()创建一个所需类型的新生成器(参见Generator Types)。

2. 设置生成器选项（参见Generator Options）；例如，使用curandSetPseudoRandomGeneratorSeed()来设置种子。

3. 使用cudaMalloc()在设备上分配内存。

4. 使用curandGenerate()或其他生成函数生成随机数。

5. 使用结果。

6. 如果需要，可以通过多次调用curandGenerate()来生成更多随机数。

7. 使用curandDestroyGenerator()进行清理。

要在主机CPU上生成随机数，在上述第一步中调用curandCreateGeneratorHost()，在第三步中分配一个主机内存缓冲区来接收结果。无论您是在设备上还是在主机CPU上生成随机数，所有其他调用的工作方式都相同。

同时创建多个生成器是合法的。每个生成器封装了一个独立的状态，并且与其他所有生成器互不干扰。每个生成器产生的数字序列都是确定性的。给定相同的设置参数，每次运行程序时都会生成相同的序列。在设备上生成随机数将产生与在主机CPU上生成相同的序列。

请注意，上述步骤4中的curandGenerate()会启动一个内核并异步返回。如果您在另一个流中启动另一个内核，且该内核需要使用curandGenerate()的结果，您必须调用cudaThreadSynchronize()或使用流管理/事件管理例程，以确保在启动新内核之前随机数生成内核已完成执行。

请注意，将主机内存指针传递给在设备上运行的生成器是无效的，同样将设备内存指针传递给在CPU上运行的生成器也是无效的。这些情况下的行为是未定义的。

2.1. 生成器类型

随机数生成器通过向curandCreateGenerator()传递类型来创建。cuRAND中有九种随机数生成器，分为两大类。CURAND_RNG_PSEUDO_XORWOW、CURAND_RNG_PSEUDO_MRG32K3A、CURAND_RNG_PSEUDO_MTGP32、CURAND_RNG_PSEUDO_PHILOX4_32_10和CURAND_RNG_PSEUDO_MT19937是伪随机数生成器。CURAND_RNG_PSEUDO_XORWOW使用XORWOW算法实现，属于xor-shift家族的伪随机数生成器。CURAND_RNG_PSEUDO_MRG32K3A是组合多重递归家族的伪随机数生成器成员。CURAND_RNG_PSEUDO_MT19937和CURAND_RNG_PSEUDO_MTGP32属于梅森旋转家族的伪随机数生成器。CURAND_RNG_PSEUDO_MTGP32的参数针对GPU操作进行了定制。CURAND_RNG_PSEUDO_MT19937的参数与CPU版本相同，但排序方式不同。CURAND_RNG_PSEUDO_MT19937仅支持HOST API，且只能在sm_35或更高架构上使用。CURAND_RNG_PHILOX4_32_10属于Philox家族，是D E Shaw Research在SC11会议上提出的三种非加密基于计数器的随机数生成器之一。基本SOBOL准随机数生成器有4种变体，所有变体都能生成高达20,000维的序列。CURAND_RNG_QUASI_SOBOL32、CURAND_RNG_QUASI_SCRAMBLED_SOBOL32、CURAND_RNG_QUASI_SOBOL64和CURAND_RNG_QUASI_SCRAMBLED_SOBOL64是准随机数生成器类型。CURAND_RNG_QUASI_SOBOL32是生成32位序列的Sobol生成器。CURAND_RNG_QUASI_SCRAMBLED_SOBOL32是生成32位序列的加扰Sobol生成器。CURAND_RNG_QUASI_SOBOL64是生成64位序列的Sobol生成器。CURAND_RNG_QUASI_SCRAMBLED_SOBOL64是生成64位序列的加扰Sobol生成器。

2.2. 生成器选项

创建后，可以使用通用选项seed、offset和order来定义随机数生成器。

2.2.1. 种子

seed参数是一个64位整数，用于初始化伪随机数生成器的起始状态。相同的seed总是产生相同的结果序列。

2.2.2. 偏移量

offset参数用于在序列中跳过前面的部分。如果offset = 100，生成的第一个随机数将是序列中的第100个。这允许同一程序的多次运行可以从同一序列继续生成结果而不会重叠。请注意，跳过功能不适用于CURAND_RNG_PSEUDO_MTGP32和CURAND_RNG_PSEUDO_MT19937生成器。

2.2.3. 排序

order参数用于选择结果在全局内存中的排序方式。它还会直接影响cuRAND生成函数的性能。

伪随机序列有五种排序选择：CURAND_ORDERING_PSEUDO_DEFAULT、CURAND_ORDERING_PSEUDO_LEGACY、CURAND_ORDERING_PSEUDO_BEST、CURAND_ORDERING_PSEUDO_SEEDED和CURAND_ORDERING_PSEUDO_DYNAMIC。准随机数有一种排序选择CURAND_ORDERING_QUASI_DEFAULT。伪随机数生成器的默认排序是CURAND_ORDERING_PSEUDO_DEFAULT，而准随机数生成器的默认排序是CURAND_ORDERING_QUASI_DEFAULT。

两种伪随机排序方式CURAND_ORDERING_PSEUDO_DEFAULT和CURAND_ORDERING_PSEUDO_BEST对所有伪随机生成器产生相同的输出顺序，除了MT19937生成器——对于MT19937，CURAND_ORDERING_PSEUDO_DEFAULT等同于CURAND_ORDERING_PSEUDO_LEGACY。对于MT19937，CURAND_ORDERING_PSEUDO_BEST可能在不同型号的GPU上产生不同输出，且不能与通过curandCreateGeneratorHost()创建的主机生成器一起使用。未来cuRAND版本可能会更改CURAND_ORDERING_PSEUDO_BEST关联的排序方式以提升性能或结果质量。但始终保证使用CURAND_ORDERING_PSEUDO_BEST获得的排序是确定性的，且在每次程序运行时保持一致。而使用CURAND_ORDERING_PSEUDO_LEGACY获得的排序保证在所有cuRAND版本中保持不变。

CURAND_ORDERING_PSEUDO_DYNAMIC排序方式不能与通过curandCreateGeneratorHost()创建的主机生成器一起使用，目前仅支持以下伪随机生成器：CURAND_RNG_PSEUDO_XORWOW、CURAND_RNG_PSEUDO_PHILOX4_32_10、CURAND_RNG_PSEUDO_MRG32K3A和CURAND_RNG_PSEUDO_MTGP32。当选择CURAND_ORDERING_PSEUDO_DYNAMIC排序方式时，cuRAND会尝试最大化GPU利用率以获得最佳性能。使用CURAND_ORDERING_PSEUDO_DYNAMIC获得的排序在不同GPU上可能不同。不保证：在所有cuRAND版本中保持一致，以及对所有分布都相同。但保证是确定性的。

每种生成器类型在排序参数上的行为差异概述如下：

XORWOW伪随机数生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序与CURAND_ORDERING_PSEUDO_BEST相同。
- CURAND_ORDERING_PSEUDO_BEST
  
  在当前版本中，CURAND_ORDERING_PSEUDO_BEST的输出顺序与CURAND_ORDERING_PSEUDO_LEGACY相同。
- CURAND_ORDERING_PSEUDO_LEGACY
  
  全局内存中偏移量 $n$ 处的结果来自原始XORWOW序列中的位置
  $(n mod 4096) \cdot 2^{67} + ⌊ n / 4096 ⌋$
- CURAND_ORDERING_PSEUDO_DYNAMIC
  
  CURAND_ORDERING_PSEUDO_DYNAMIC的输出顺序在不同GPU上可能有所不同。
- CURAND_ORDERING_PSEUDO_SEEDED
  
  全局内存中偏移量 $n$ 处的结果来自XORWOW序列中 $n / 4096 ⌋$ 位置的值，该序列由用户种子和数值 $n mod 4096$ 组合生成。换句话说，4096个线程各自使用不同的种子。这种种子生成方法减少了状态设置时间，但对于某些用户种子值可能导致伪随机输出的统计特性较弱。
MRG32k3a伪随机数生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序与CURAND_ORDERING_PSEUDO_BEST相同。
- CURAND_ORDERING_PSEUDO_BEST
  
  全局内存中偏移量 $n$ 处的结果来自原始MRG32k3a序列中的位置
  $(n mod 81920) \cdot 2^{76} + ⌊ n / 81920 ⌋$
  （请注意，MRG32k3a后续样本之间的步长与XORWOW不同）
- CURAND_ORDERING_PSEUDO_LEGACY
  
  全局内存中偏移量 $n$ 处的结果来自原始MRG32k3a序列中的位置
  $(n mod 4096) \cdot 2^{76} + ⌊ n / 4096 ⌋$
  （请注意MRG32k3a的后续样本之间的步长与XORWOW不同）
- CURAND_ORDERING_PSEUDO_DYNAMIC
  
  CURAND_ORDERING_PSEUDO_DYNAMIC的输出顺序在不同GPU上可能有所不同。
MTGP32伪随机数生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序与CURAND_ORDERING_PSEUDO_BEST相同。
- CURAND_ORDERING_PSEUDO_BEST
  
  MTGP32生成器实际上基于基础算法的不同参数集生成了192个不同的序列。设 $S (p)$ 为参数集 $p$ 对应的序列。
  
  全局内存中偏移量 $n$ 处的结果来自序列 $n mod 256$ 位置上的值，该序列为
  $S (⌊ n / 256 ⌋ mod 192)$
  换句话说，来自 $S (0)$ 的256个样本后跟着来自 $S (1)$ 的256个样本，依此类推，直到 $S (191)$ 。这个模式会重复，因此接下来的256个样本来自 $S (0)$ ，随后是来自 $S (1)$ 的256个样本，以此类推。
- CURAND_ORDERING_PSEUDO_LEGACY
  
  MTGP32生成器实际上基于基础算法的不同参数集生成了64个不同的序列。设 $S (p)$ 表示参数集 $p$ 对应的序列。
  
  全局内存中偏移量 $n$ 处的结果来自序列 $n mod 256$ 的位置
  $S (⌊ n / 256 ⌋ mod 64)$
  换句话说，先是来自 $S (0)$ 的256个样本，接着是来自 $S (1)$ 的256个样本，依此类推，直到 $S (63)$ 。这个模式会重复，因此接下来的256个样本又来自 $S (0)$ ，然后是来自 $S (1)$ 的256个样本，以此类推。
- CURAND_ORDERING_PSEUDO_DYNAMIC
  
  CURAND_ORDERING_PSEUDO_DYNAMIC的输出顺序在不同GPU上可能不同。在这种排序方式下，MTGP32可以使用与原始MTGP32实现不同的预计算参数。
MT19937伪随机数生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序与CURAND_ORDERING_PSEUDO_LEGACY相同。
- CURAND_ORDERING_PSEUDO_LEGACY
  
  排序方式主要基于标准的MT19937 CPU实现。输出由8192个独立的生成器产生。每个生成器生成原始序列的连续子序列。每个子序列的长度为 $2^{1000}$ 。随机数以8个为一组生成，因此前8个元素来自第一个子序列，接下来的8个元素来自第二个子序列，依此类推。为了提高性能，结果采用了不同于原始方式的排列。排序方式与所使用的硬件无关。更多信息请参阅[18]。
- CURAND_ORDERING_PSEUDO_BEST
  
  CURAND_ORDERING_PSEUDO_BEST的输出排序方式会根据GPU中SM(流式多处理器)的数量来优化性能。随机数的生成方式与CURAND_ORDERING_PSEUDO_LEGACY相同，但生成器的数量可能不同以实现更好的性能。使用这种排序方式生成种子的速度会快很多。
  
  排序方式CURAND_ORDERING_PSEUDO_BEST仅支持GPU cuRAND随机数生成器，不能用于通过curandCreateGeneratorHost()创建的宿主生成器。
Philox_4x32_10 伪随机数生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT的输出顺序与CURAND_ORDERING_PSEUDO_BEST相同。
- CURAND_ORDERING_PSEUDO_BEST
  
  在当前版本中，CURAND_ORDERING_PSEUDO_BEST的输出顺序与CURAND_ORDERING_PSEUDO_LEGACY相同。
- CURAND_ORDERING_PSEUDO_LEGACY
  
  Philox_4x32_10生成器中的每个线程会根据基础算法的不同参数集生成不同的序列。在主机API中存在65536种不同的序列。每个序列的四个值之后会接着下一个序列的四个值。
- CURAND_ORDERING_PSEUDO_DYNAMIC
  
  CURAND_ORDERING_PSEUDO_DYNAMIC的输出顺序在不同GPU上可能有所不同。
32位和64位SOBOL及加扰SOBOL准随机数生成器
- CURAND_ORDERING_QUASI_DEFAULT
  
  当在 $d$ 维度中生成 $n$ 个结果时，输出将包含 $n / d$ 个来自维度1的结果，接着是 $n / d$ 个来自维度2的结果，依此类推直到维度 $d$ 。只能生成维度大小的整数倍结果。维度参数 $d$ 通过curandSetQuasiRandomGeneratorDimensions()设置，默认值为1。

2.3. 返回值

所有cuRAND主机库函数调用都会返回一个curandStatus_t类型的值。成功执行且无错误的调用将返回CURAND_STATUS_SUCCESS。如果发生错误，则会根据具体错误返回其他值。由于CUDA允许内核与CPU代码异步执行，因此在调用库函数时可能会检测到非cuRAND内核中的错误。这种情况下将返回CURAND_STATUS_PREEXISTING_ERROR。

2.4. 生成函数

curandStatus_t 
curandGenerate(
    curandGenerator_t generator, 
    unsigned int *outputPtr, size_t num)
    
curandStatus_t 
curandGenerateLongLong(
    curandGenerator_t generator, 
    unsigned long long *outputPtr, size_t num)

curandGenerate() 函数用于为 XORWOW、MRG32k3a、MTGP32、MT19937、Philox_4x32_10 和 SOBOL32 生成器生成伪随机或准随机比特输出。每个输出元素是一个32位无符号整数，其中所有比特都是随机的。对于 SOBOL64 生成器，每个输出元素是一个64位无符号长整型，其中所有比特都是随机的。curandGenerate() 对 SOBOL64 生成器会返回错误。请使用 curandGenerateLongLong() 来通过 SOBOL64 生成器生成64位整数。

curandStatus_t 
curandGenerateUniform(
    curandGenerator_t generator, 
    float *outputPtr, size_t num)

curandGenerateUniform() 函数用于生成0.0到1.0之间均匀分布的浮点数值，其中不包含0.0但包含1.0。

curandStatus_t 
curandGenerateNormal(
    curandGenerator_t generator, 
    float *outputPtr, size_t n, 
    float mean, float stddev)

curandGenerateNormal() 函数用于生成符合给定均值和标准差的正态分布浮点数值。

curandStatus_t 
curandGenerateLogNormal(
    curandGenerator_t generator, 
    float *outputPtr, size_t n, 
    float mean, float stddev)

curandGenerateLogNormal() 函数用于基于给定均值和标准差的正态分布生成对数正态分布的浮点数值。

curandStatus_t 
curandGeneratePoisson(
    curandGenerator_t generator, 
    unsigned int *outputPtr, size_t n, 
    double lambda)

curandGeneratePoisson() 函数用于基于给定lambda参数的泊松分布生成泊松分布的整数值。

curandStatus_t
curandGenerateUniformDouble(
    curandGenerator_t generator, 
    double *outputPtr, size_t num)

curandGenerateUniformDouble() 函数生成双精度均匀分布的随机数。

curandStatus_t
curandGenerateNormalDouble(
    curandGenerator_t generator,
    double *outputPtr, size_t n, 
    double mean, double stddev)

curandGenerateNormalDouble() 使用给定的均值和标准差生成双精度正态分布结果。双精度结果只能在计算能力1.3或更高的设备以及主机上生成。

curandStatus_t
curandGenerateLogNormalDouble(
    curandGenerator_t generator,
    double *outputPtr, size_t n, 
    double mean, double stddev)

curandGenerateLogNormalDouble() 基于给定均值和标准差的正态分布，生成双精度对数正态分布结果。

对于准随机生成，返回结果的数量必须是生成器维度的倍数。

生成函数可以在同一个生成器上多次调用，以生成连续的结果块。对于伪随机生成器，多次调用生成函数将产生与一次大尺寸调用相同的结果。对于准随机生成器，由于内存中维度的排序方式，多次较短调用不会在内存中产生与一次较大调用相同的结果；然而生成的 $n$ 维向量将是相同的。

双精度结果只能在计算能力1.3或更高的设备以及主机上生成。

2.5. 主机API示例


/*
 * This program uses the host CURAND API to generate 100
 * pseudorandom floats.
 */
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand.h>

#define CUDA_CALL(x) do { if((x)!=cudaSuccess) { \
    printf("Error at %s:%d\n",__FILE__,__LINE__);\
    return EXIT_FAILURE;}} while(0)
#define CURAND_CALL(x) do { if((x)!=CURAND_STATUS_SUCCESS) { \
    printf("Error at %s:%d\n",__FILE__,__LINE__);\
    return EXIT_FAILURE;}} while(0)

int main(int argc, char *argv[])
{
    size_t n = 100;
    size_t i;
    curandGenerator_t gen;
    float *devData, *hostData;

    /* Allocate n floats on host */
    hostData = (float *)calloc(n, sizeof(float));

    /* Allocate n floats on device */
    CUDA_CALL(cudaMalloc((void **)&devData, n*sizeof(float)));

    /* Create pseudo-random number generator */
    CURAND_CALL(curandCreateGenerator(&gen,
                CURAND_RNG_PSEUDO_DEFAULT));

    /* Set seed */
    CURAND_CALL(curandSetPseudoRandomGeneratorSeed(gen,
                1234ULL));

    /* Generate n floats on device */
    CURAND_CALL(curandGenerateUniform(gen, devData, n));

    /* Copy device memory to host */
    CUDA_CALL(cudaMemcpy(hostData, devData, n * sizeof(float),
        cudaMemcpyDeviceToHost));

    /* Show result */
    for(i = 0; i < n; i++) {
        printf("%1.4f ", hostData[i]);
    }
    printf("\n");

    /* Cleanup */
    CURAND_CALL(curandDestroyGenerator(gen));
    CUDA_CALL(cudaFree(devData));
    free(hostData);
    return EXIT_SUCCESS;
}

2.6. 静态库支持

从6.5版本开始，cuRAND库在Linux和Mac系统上也以静态库形式提供，文件名为libcurand_static.a。Windows系统不支持静态库。静态cuRAND库依赖于一个通用的线程抽象层库，在Linux和Mac上称为libcuos.a，在Windows上称为cuos.lib。

例如，在Linux系统上，要使用cuRAND动态库编译一个小型应用程序，可以使用以下命令：

nvcc myCurandApp.c  -lcurand  -o myCurandApp

若要针对静态cuRAND库进行编译，则必须使用以下命令：

     
nvcc myCurandApp.c  -lcurand_static   -lculibos -o myCurandApp

也可以使用原生的Host C++编译器。根据主机操作系统的不同，在链接行可能需要一些额外的库，如pthread或dl。在Linux上建议使用以下命令：

        
g++ myCurandApp.c  -lcurand_static   -lculibos -lcudart_static -lpthread -ldl -I <cuda-toolkit-path>/include -L <cuda-toolkit-path>/lib64 -o myCurandApp

请注意，在后一种情况下，不需要库cuda。CUDA运行时会根据需要尝试显式打开cuda库。对于未安装CUDA驱动的系统，这允许应用程序优雅地处理此问题，并在仅支持CPU路径可用时继续运行。

2.7. 性能说明

通常情况下，通过生成尽可能大的随机数块，您将从cuRAND库获得最佳性能。生成大量随机数的少量调用比生成少量随机数的多次调用更高效。默认的伪随机生成器XORWOW在首次调用时需要一些时间进行初始化设置（采用默认排序方式时）。后续生成调用则不需要这种设置。要避免这种初始化时间，请使用CURAND_ORDERING_PSEUDO_SEEDED排序方式。

MTGP32梅森旋转算法与线程和块数量密切相关。MTGP32的状态结构实际上包含来自特定参数集确定的序列的256个连续样本的状态。64个块中的每一个使用不同的参数集，256个线程中的每一个从状态生成一个样本并更新状态。因此，MTGP32最高效的使用方式是生成16384个样本的倍数。

MT19937算法的性能取决于单次调用期间生成的样本数量。在生成超过2GB数据时可达到峰值性能，但仅生成80MB数据时也能实现80%的峰值性能。具体参考请见[18]。

Philox_4x32_10算法与线程和块数量密切相关。每个线程在同一时间计算4个随机数，因此最有效使用Philox_4x32_10的方式是生成线程数量的4倍数值。

为了获得cuRAND主机API的最佳性能，建议用户使用CURAND_ORDERING_PSEUDO_BEST或CURAND_ORDERING_PSEUDO_DYNAMIC排序方式。

2.8. 线程安全

只要不同的主机线程使用不同的生成器，生成器不是MT19937（CURAND_RNG_PSEUDO_MT19937），并且输出不重叠，cuRAND主机API就是线程安全的。

请注意，当与MT19937生成器（CURAND_RNG_PSEUDO_MT19937）一起使用时，cuRAND主机API不是线程安全的。