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 样例代码如下:

#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;
}

原文链接:https://www.cnblogs.com/CocoML/p/17376020.html

本站文章如无特殊说明,均为本站原创,如若转载,请注明出处:CUDA 的随机数算法 API - Python技术站

(0)
上一篇 2023年5月6日
下一篇 2023年5月6日

相关文章

  • linux下caffe的命令运行脚本

    参考:https://www.cnblogs.com/denny402/p/5076285.html 首先编译: make -j8 make pycaffe   注:下面的–solver=…. 等价于 -solver …. ########################## -solver:必选参数 ###################set …

    Caffe 2023年4月8日
    00
  • 生成对抗网络GAN资料汇总(附pdf下载)

    如今,GAN技术已延伸至AI各方向,如图像处理、自然语言处理和语音等。深度学习三巨头之一 Yann LeCun :“这是过去十年间机器学习领域最让人激动的点子!”。2018 GAN 与人造胚胎等入选《麻省理工科技评论》全球十大突破性技术。 训练不稳定、难控制,并不影响学界对其追捧;这也反过来促进GAN发展和完善。短视频火热,多家互联网公司(腾讯百度等)专门设…

    GAN生成对抗网络 2023年4月7日
    00
  • pytorch的visdom启动不了、蓝屏

    pytorch的visdom启动不了、蓝屏     问题描述:我是在ubuntu16。04上用python3.5安装的visdom。可是启动是蓝屏:在网上找了很久的解决方案:有三篇博文:      https://blog.csdn.net/qq_22194315/article/details/78827185 https://blog.csdn.net/…

    PyTorch 2023年4月8日
    00
  • 吴恩达深度学习 course4 卷积神经网络

    1.computer vision cv是深度学习的一个重要方向,cv一般而言包括:图像识别,目标检测,神经风格转换 传统的神经网络所存在的问题:图片的输入维度比较大,具体如下图所示,这就造成了权重w的维度比较大,那么他所占用的内存也会比较大,计算w的计算量也会很大 所以我们会引入卷积神经网络 2.Edge detection example 神经网络由浅层…

    2023年4月8日
    00
  • Pytorch 中 tensor的维度拼接

    torch.stack() 和 torch.cat() 都可以按照指定的维度进行拼接,但是两者也有区别,torch.satck() 是增加新的维度进行堆叠,即其维度拼接后会增加一个维度;而torch.cat() 是在原维度上进行堆叠,即其维度拼接后的维度个数和原来一致。具体说明如下: torch.stack(input,dim) input: 待拼接的张量序…

    PyTorch 2023年4月8日
    00
  • 【tensorflow】tf.keras + 神经网络类class 6 步搭建神经网络

    tf.keras + Sequential() 可以搭建出上层输入就是下层输出的顺序网络结构,但是无法写出一些带有跳连的非顺序网络结构。 这时候可以选择用类 class 搭建神经网络结构,即使用 class 类封装一个网络结构:   … class MyModel(Model):   def __init__(self):     super(MyMod…

    Keras 2023年4月8日
    00
  • caffe 生成检测框并绘图

      Step 1 使用训练好的模型检测图片: build/examples/ssd/ssd_detect.bin models/VGGNet/VOC0712/SSD_300x300/deploy.prototxt /home/daisida/ssd/models/VGGNet/VOC0712/SSD_300x300/VGG_VOC0712_SSD_300x3…

    Caffe 2023年4月8日
    00
  • 基于Hadoop的机器学习开源项目

    Mahout 是机器学习和数据挖掘的一个分布式框架,区别于其他的开源数据挖掘软件,它是基于hadoop之上的; 所以Mahout的优势就是hadoop的优势。官网http://mahout.apache.org/ 上说的Scalable就是指hadoop的可扩展性。 Mahout用map-reduce实现了部分数据挖掘算法,解决了并行挖掘的问题。这里说的“解…

    机器学习 2023年4月10日
    00
合作推广
合作推广
分享本页
返回顶部