CUDA 的随机数算法 API

这篇具有很好参考价值的文章主要介绍了CUDA 的随机数算法 API。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

参考自 Nvidia cuRand 官方 API 文档

一、具体使用场景

如下是是在 dropout 优化中手写的 uniform_random 的 Kernel:

#include <cuda_runtime.h>
#include <curand_kernel.h>

__device__ inline float cinn_nvgpu_uniform_random_fp32(int seed){
  curandStatePhilox4_32_10_t state;
  int idx = threadIdx.x + blockIdx.x * blockDim.x;
  curand_init(seed, idx, 1, &state);
  return curand_uniform(&state);
}

二、API 解析

我们首先来看 curand_init 函数的签名和语义:

__device__ void
curand_init(unsigned long long seed,
            unsigned long long subsequence,
            unsigned long long offset,
            curandStatePhilox4_32_10_t *state)

给定相同的seed、sequence、offset 参数下,curand_init 会保证产生相同的其实状态 state。另外此函数会在调用 2^67 ⋅ sequence + offset次 cu_rand API 之后「重置」为起始状态。关于 sequence 和 offset 的如何生效的机制,参考 StackOverflow。

注意:在 seed 相同、sequence 不同时,一般不会产生有统计学上关联的结果。原文:Sequences generated with the same seed and different sequence numbers will not have statistically correlated values.

另外在CUDA 并行产生随机数实践上,也有一些经验上的建议:

  • 若保证高质量的伪随机数,建议使用不同的 seed
  • 若是并行在一个「实验」里,建议指定不同的sequence参数,且最好「单调递增」
  • 若果线程里的config都是一样,即 state 一样,则可以把「随机状态量」放到 global memory里,以减少 setup 开销

参考原文:For the highest quality parallel pseudorandom number generation, each experiment should be assigned a unique seed. Within an experiment, each thread of computation should be assigned a unique sequence number. If an experiment spans multiple kernel launches, it is recommended that threads between kernel launches be given the same seed, and sequence numbers be assigned in a monotonically increasing way. If the same configuration of threads is launched, random state can be preserved in global memory between launches to avoid state setup time.

然后我们看下Nvidia 主要提供了哪些常用的随机数生成API:

__device__ float
curand_uniform (curandState_t *state);   // <----  It may return from 0.0 to 1.0, where 1.0 is included and 0.0 is excluded.

__device__ float
curand_normal (curandState_t *state);    // <-----  returns a single normally distributed float with mean 0.0 and standard deviation 1.0.

__device__ float
curand_log_normal (curandState_t *state, float mean, float stddev); // <----- returns a single log-normally distributed float based on a normal distribution with the given mean and standard deviation.

// 如下是上述 3 个API 的 double 版本
__device__ double
curand_uniform_double (curandState_t *state);

__device__ double
curand_normal_double (curandState_t *state);

__device__ double
curand_log_normal_double (curandState_t *state, double mean, double stddev);

上面的 device API 在每次调用时,只会生成一个 float/double 的随机数。Nvidia 同样提供了一次可以生成 2个或4个 device API:

__device__ uint4
curand4 (curandStatePhilox4_32_10_t *state);

__device__ float4
curand_uniform4 (curandStatePhilox4_32_10_t *state);

__device__ float4
curand_normal4 (curandStatePhilox4_32_10_t *state);

__device__ float4
curand_log_normal4 (curandStatePhilox4_32_10_t *state, float mean, float stddev);

从上面的函数接口以及 Nvidia 的文档来看,在初始化某种类型的 state 状态量后,每次调用类似 curand() 的 API 后,state 都会自动进行 offset 偏移。

因此,Nvidia 官网上也提供了单独对 state 进行位移的 API,其效果等价于调用多次无返回值的 curand() API,且性能更好:

__device__ void
skipahead(unsigned long long n, curandState_t *state); // <----- == calls n*curand()

__device__ void
skipahead_sequence(unsigned long long n, curandState_t *state); // <----- == calls n*2^67 curand()

三、性能分析

Nvidia 的官网明确指出了存在的性能问题,给开发者实现高性能 Kernel 提供了充分的经验指导:

  • curand_init()要比curand()和curand_uniform()慢!
  • curand_init()在 offset 比较大时性能也会比小 offset 差!
  • save/load操作 state 比每次重复创建起始 state 性能要快很多 !

原文如下:Calls to curand_init() are slower than calls to curand() or curand_uniform(). Large offsets to curand_init() take more time than smaller offsets. It is much faster to save and restore random generator state than to recalculate the starting state repeatedly.

对于上述第三点,Nvidia 建议可以将 state 存放到 global memory 中,如下是一个样例代码:

__global__ void example(curandState *global_state)
{
    curandState local_state;
    local_state = global_state[threadIdx.x];
    for(int i = 0; i < 10000; i++) {
        unsigned int x = curand(&local_state);
        ...
    }
    global_state[threadIdx.x] = local_state;
}

从另一个维度来讲,相对于A产生随机数操作的API,初始化 state 会占用更多的「寄存器」和 local memory 资源。因此 nvidia 建议将 curand_initcurand() API 拆分放到不同的 Kernel 中,可以获得最大的性能收益;

原文:Initialization of the random generator state generally requires more registers and local memory than random number generation. It may be beneficial to separate calls to curand_init() and curand() into separate kernels for maximum performance.
State setup can be an expensive operation. One way to speed up the setup is to use different seeds for each thread and a constant sequence number of 0. This can be especially helpful if many generators need to be created. While faster to set up, this method provides less guarantees about the mathematical properties of the generated sequences. If there happens to be a bad interaction between the hash function that initializes the generator state from the seed and the periodicity of the generators, there might be threads with highly correlated outputs for some seed values. We do not know of any problem values; if they do exist they are likely to be rare.

Nvidia 提供的 API 样例代码如下:文章来源地址https://www.toymoban.com/news/detail-434818.html

#include <stdio.h>
#include <stdlib.h>

#include <cuda_runtime.h>
#include <curand_kernel.h>

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

__global__ void setup_kernel(curandState *state)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    /* Each thread gets same seed, a different sequence
       number, no offset */
    curand_init(1234, id, 0, &state[id]);
}

__global__ void generate_uniform_kernel(curandStatePhilox4_32_10_t *state,
                                int n,
                                unsigned int *result)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int count = 0;
    float x;
    /* Copy state to local memory for efficiency */
    curandStatePhilox4_32_10_t localState = state[id];
    /* Generate pseudo-random uniforms */
    for(int i = 0; i < n; i++) {
        x = curand_uniform(&localState);
        /* Check if > .5 */
        if(x > .5) {
            count++;
        }
    }
    /* Copy state back to global memory */
    state[id] = localState;
    /* Store results */
    result[id] += count;
}

到了这里,关于CUDA 的随机数算法 API的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • Hadoop中的加密解密机制——伪随机数生成算法介绍

    作者:禅与计算机程序设计艺术 Hadoop作为当下最流行的大数据处理平台,提供了丰富的功能支持,如海量数据的存储、分析与计算。其中一个重要的环节就是数据安全问题。无论是存储集群还是计算集群,都需要提供数据保护措施来确保数据的完整性和可用性。今天我将介绍

    2024年02月06日
    浏览(58)
  • C# 随机数生成 Mersenne Twister 马特赛特旋转演算法 梅森旋转算法

    NuGet安装MathNet.Numerics 引用: using MathNet.Numerics.Random; Mersenne Twister算法是一种伪随机数发生器,名字来源于其周期长度通常取自Mersenne质数。它是由Makoto Matsumoto(松本)和Takuji Nishimura(西村)于1997年开发的,基于有限二进制字段上的矩阵线性再生。它的主要作用是生成伪随机数

    2024年02月09日
    浏览(35)
  • iOS MT19937随机数生成,结合AES-CBC加密算法实现。

    按处理顺序说明: 1. 生成随机数序列字符串函数 生成方法MT19937,初始种子seed,利用C++库方法,生成: 2. 对第一部中的随机数序列字符串进行sha256加密,得到64字节的一个数据流函数。  3. AES-CBC加密解密方法 /*     CCCrypt方法提供了CBC 和 ECB 两种AES加密模式,     如果不传

    2024年04月09日
    浏览(52)
  • 【Python】蒙特卡洛模拟 | PRNG 伪随机数发生器 | 马特赛特旋转算法 | LCG 线性同余算法 | Python Random 模块

          猛戳订阅!  👉 《一起玩蛇》🐍 💭 写在前面: 本篇博客将介绍经典的伪随机数生成算法,我们将  重点讲解 LCG(线性同余发生器) 算法与马特赛特旋转算法,在此基础上顺带介绍 Python 的 random 模块。   本篇博客还带有练习,无聊到喷水的练习,咳咳…… 学完前

    2024年02月04日
    浏览(48)
  • JS - 生成随机数的方法汇总(不同范围、类型的随机数)

    (1)使用 random() 方法可以返回一个介于 0 ~ 1 之间的伪随机数(包括 0,不包括 1)。 (2)下面是一个测试样例 (1)这种最简单,因为和 random 的特点保持一致。只需使用如下公式即可: (2)比如下面生成 [10,15) 范围内的随机浮点数。 因为 random 的特点,要取得这几个区间

    2023年04月08日
    浏览(52)
  • Unity 中的随机数的基础常用的随机数生成方法

    在 Unity 中,可以使用 Random 类来生成随机数。以下是一些常用的随机数生成方法: Random.Range(min, max):生成一个在[min, max)范围内的随机整数。 Random.value:生成一个在[0, 1)范围内的随机浮点数。 Random.insideUnitCircle:生成一个在单位圆内的随机二维向量。 Random.insideUnitSphere:生成

    2024年02月20日
    浏览(55)
  • Hutool 生成随机数和随机字符串

    官方文档: https://www.hutool.cn/docs/#/core/工具类/随机工具-RandomUtil 整理完毕,完结撒花~

    2024年02月16日
    浏览(58)
  • MySQL、Oracle 生成随机ID、随机数、随机字符串

    UUID():是由128位的数字组成的全局唯一标识符。每次都生成一个新的随机数。 它通常以32个十六进制数的形式表示,分为5个部分,以连字符分隔。 UUID的长度是36个字符,包括32个十六进制数字和4个连字符。 UUID的标准格式是由 8-4-4-4-12 个十六进制数字组成的,其中每个部分的

    2024年01月16日
    浏览(57)
  • java生成随机数

       bound 必须是正数。 以下代码生成的是 0 到 30 的随机数。 生成区间的随机数:[最小值,最大值] 学的不是技术,更是梦想!!!

    2024年02月07日
    浏览(66)
  • Flutter 生成随机数

    如何让随机数变化? 我们尝试过的都知道,当你创建出来一个随机数后,调用他他的值是随机的,但是,这时候他的值就会固定住,不管怎么样都是随机出来的那个数,如果想要他每次都不一样的话,我们就想要使用刷新来让他变化了。 我们可以使用这样的方法来使他每次不一

    2024年02月13日
    浏览(47)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包