cuRAND :: CUDA 工具包文档

2. 主机 API 概述

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

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

1. 使用 curandCreateGenerator() 创建所需类型的新生成器（参见生成器类型）。

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

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

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

5. 使用结果。

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

7. 使用 curandDestroyGenerator() 清理。

要在主机 CPU 上生成随机数，在上面的步骤 1 中调用 curandCreateGeneratorHost()，在步骤 3 中，分配一个主机内存缓冲区以接收结果。所有其他调用工作方式相同，无论您是在设备上还是在主机 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 算法实现，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 仅支持主机 API，并且只能在 sm_35 或更高架构上使用。CURAND_RNG_PHILOX4_32_10 是 Philox 家族的成员，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. 生成器选项

创建后，可以使用通用选项种子、偏移量和顺序来定义随机数生成器。

2.2.1. 种子

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

2.2.2. 偏移量

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

2.2.3. 顺序

顺序参数用于选择结果在全局内存中的排序方式。它也直接影响 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。

对于所有伪随机生成器，除了 MT19937，CURAND_ORDERING_PSEUDO_DEFAULT 和 CURAND_ORDERING_PSEUDO_BEST 这两个伪随机排序产生相同的输出排序，对于 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$ 的全局内存中的结果来自位置
  $(n mod 4096) \cdot 2^{67} + ⌊ n / 4096 ⌋$
  在原始 XORWOW 序列中。
- CURAND_ORDERING_PSEUDO_DYNAMIC
  
  CURAND_ORDERING_PSEUDO_DYNAMIC 的输出排序在不同的 GPU 上可能不同。
- CURAND_ORDERING_PSEUDO_SEEDED
  
  偏移量为 $n$ 的全局内存中的结果来自位置 $n / 4096 ⌋$ 在以用户种子和数字组合播种的 XORWOW 序列中 $n mod 4096$ 。换句话说，4096 个线程中的每个线程都使用不同的种子。这种播种方法减少了状态设置时间，但对于某些用户种子值，可能会导致伪随机输出的统计弱点。
MRG32k3a 伪随机生成器
- CURAND_ORDERING_PSEUDO_DEFAULT
  
  在当前版本中，CURAND_ORDERING_PSEUDO_DEFAULT 的输出排序与 CURAND_ORDERING_PSEUDO_BEST 相同。
- CURAND_ORDERING_PSEUDO_BEST
  
  偏移量为 $n$ 的全局内存中的结果来自位置
  $(n mod 81920) \cdot 2^{76} + ⌊ n / 81920 ⌋$
  在原始 MRG32k3a 序列中。（请注意，MRG32k3a 的后续样本之间的步幅与 XORWOW 不同）
- CURAND_ORDERING_PSEUDO_LEGACY
  
  偏移量为 $n$ 的全局内存中的结果来自位置
  $(n mod 4096) \cdot 2^{76} + ⌊ n / 4096 ⌋$
  在原始 MRG32k3a 序列中。（请注意，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 个元素来自第二个子序列，依此类推。结果的排列方式与原始排列方式不同，以实现更高的性能。排序与您使用的硬件无关。有关更多信息，请参见[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
  
  当生成 $n$ 结果在 $d$ 维度中时，输出将由来自维度 1 的 $n / d$ 个结果、来自维度 2 的 $n / d$ 个结果以及直到维度 $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 库依赖于一个名为 libcuos.a（在 Linux 和 Mac 上）和 cuos.lib（在 Windows 上）的通用线程抽象层库。

例如，在 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 不是线程安全的。