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日

相关文章

  • caffe多标签输入和多任务学习

    拷贝convert_imageset,生成新工程convert_imageset_multi_label   修改源码 std::ifstream infile(argv[2]); std::vector<std::pair<std::string, std::vector<float>> > lines; std::st…

    2023年4月8日
    00
  • Caffe初试(二)windows下的cafee训练和测试mnist数据集

        一、mnist数据集 mnist是一个手写数字数据库,由Google实验室的Corinna Cortes和纽约大学柯朗研究院的Yann LeCun等人建立,它有60000个训练样本集和10000个测试样本集。mnist数据库官方网址为:http://yann.lecun.com/exdb/mnist/ 。可直接下载四个解压文件,分别对应:训练集样本、…

    Caffe 2023年4月7日
    00
  • Conv in caffe

    本文用于提醒自己caffe中的卷积,包括前向传播与反向传播。        Caffe利用im2col将输入图片(包括多个图片,即通道)转换成矩阵,从而使得卷积运算转换成矩阵运算。im2col函数可以参考matlab中的im2col。 转换图如下: 在caffe中卷积过程如下: 从上一篇博客可以知道全连接层的求导。随后的卷积层求导无非是逐次求导的过程: 其中…

    2023年4月8日
    00
  • mxnet与tensorflow的卷积实现细节比较

      mxnet的卷积 kernel = 3  pad=1边界补充0后,不管stride是否1还是2,imgw = 奇数或者偶数, 都是从图像位置(0,0)开始卷积   tensorlfow的卷积 kernel = 3 pad=‘SAME’边界补充0后, imgw = 偶数           stride=1, 是从图像位置(0,0)开始卷积 stride=…

    卷积神经网络 2023年4月8日
    00
  • 【目标检测】Fast RCNN算法详解

    转自:https://blog.csdn.net/shenxiaolu1984/article/details/51036677 Girshick, Ross. “Fast r-cnn.” Proceedings of the IEEE International Conference on Computer Vision. 2015. 继2014年的RCN…

    2023年4月8日
    00
  • 编译caffe报错:_ZN5boost16exception_detail10bad_alloc_D2Ev

    具体报错信息很长的。 text._ZN5boost16exception_detail10bad_alloc_D2Ev[_ZN5boost16exception_detail10bad_alloc_D5Ev] of .build_release/src/caffe/data_reader.o 报了这个错误好奇怪的。查了好久都没解决。 摸索半天终于解决了。好爽…

    Caffe 2023年4月7日
    00
  • Keras深度学习之卷积神经网络(CNN)

    一、总结 一句话总结: 卷积就是特征提取,后面可接全连接层来分析这些特征     二、Keras深度学习之卷积神经网络(CNN) 转自或参考:Keras深度学习之卷积神经网络(CNN)https://www.cnblogs.com/wj-1314/articles/9621901.html Keras–基于python的深度学习框架        Keras…

    2023年4月7日
    00
  • 解决Jupyter notebook[import tensorflow as tf]报错

     参考: https://blog.csdn.net/caicai_zju/article/details/70245099

    tensorflow 2023年4月6日
    00
合作推广
合作推广
分享本页
返回顶部