random_generator.h

Go to the documentation of this file.

/*
 * Licensed to the Apache Software Foundation (ASF) under one
 * or more contributor license agreements.  See the NOTICE file
 * distributed with this work for additional information
 * regarding copyright ownership.  The ASF licenses this file
 * to you under the Apache License, Version 2.0 (the
 * "License"); you may not use this file except in compliance
 * with the License.  You may obtain a copy of the License at
 *
 *   http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing,
 * software distributed under the License is distributed on an
 * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
 * KIND, either express or implied.  See the License for the
 * specific language governing permissions and limitations
 * under the License.
 */
 
#ifndef MXNET_RANDOM_GENERATOR_H_
#define MXNET_RANDOM_GENERATOR_H_
 
#include <random>
#include <new>
#include "./base.h"
 
#if MXNET_USE_CUDA
#include <curand_kernel.h>
#include <math.h>
#endif  // MXNET_USE_CUDA
 
namespace mxnet {
namespace common {
namespace random {
 
template <typename Device, typename DType MSHADOW_DEFAULT_DTYPE>
class RandGenerator;
 
template <typename DType>
class RandGenerator<cpu, DType> {
 public:
  // at least how many random numbers should be generated by one CPU thread.
  static const int kMinNumRandomPerThread;
  // store how many global random states for CPU.
  static const int kNumRandomStates;
 
  // implementation class for random number generator
  // TODO(alexzai): move impl class to separate file - tracked in MXNET-948
  class Impl {
   public:
    typedef
        typename std::conditional<std::is_floating_point<DType>::value, DType, double>::type FType;
    explicit Impl(RandGenerator<cpu, DType>* gen, int state_idx)
        : engine_(gen->states_ + state_idx) {}
 
    Impl(const Impl&) = delete;
    Impl& operator=(const Impl&) = delete;
 
    MSHADOW_XINLINE int rand() {
      return engine_->operator()();
    }
 
    MSHADOW_XINLINE int64_t rand_int64() {
      return static_cast<int64_t>(engine_->operator()() << 31) + engine_->operator()();
    }
 
    MSHADOW_XINLINE FType uniform() {
      typedef typename std::conditional<std::is_integral<DType>::value,
                                        std::uniform_int_distribution<DType>,
                                        std::uniform_real_distribution<FType>>::type GType;
      GType dist_uniform;
      return dist_uniform(*engine_);
    }
 
    MSHADOW_XINLINE FType normal() {
      std::normal_distribution<FType> dist_normal;
      return dist_normal(*engine_);
    }
 
   private:
    std::mt19937* engine_;
  };  // class RandGenerator<cpu, DType>::Impl
 
  static void AllocState(RandGenerator<cpu, DType>* inst) {
    inst->states_ = new std::mt19937[kNumRandomStates];
  }
 
  static void FreeState(RandGenerator<cpu, DType>* inst) {
    delete[] inst->states_;
  }
 
  MSHADOW_XINLINE void Seed(mshadow::Stream<cpu>*, uint32_t seed) {
    for (int i = 0; i < kNumRandomStates; ++i)
      (states_ + i)->seed(seed + i);
  }
 
  // export global random states, used by c++ custom operator
  MSHADOW_XINLINE void* GetStates() {
    return static_cast<void*>(states_);
  }
 
 private:
  std::mt19937* states_;
};  // class RandGenerator<cpu, DType>
 
template <typename DType>
const int RandGenerator<cpu, DType>::kMinNumRandomPerThread = 64;
 
template <typename DType>
const int RandGenerator<cpu, DType>::kNumRandomStates = 1024;
 
#if MXNET_USE_CUDA
 
template <typename DType>
class RandGenerator<gpu, DType> {
 public:
  // at least how many random numbers should be generated by one GPU thread.
  static const int kMinNumRandomPerThread;
  // store how many global random states for GPU.
  static const int kNumRandomStates;
 
  // uniform number generation in Cuda made consistent with stl (include 0 but exclude 1)
  // by using 1.0-curand_uniform().
  // Needed as some samplers in sampler.h won't be able to deal with
  // one of the boundary cases.
  // TODO(alexzai): move impl class to separate file - tracked in MXNET-948
  class Impl {
   public:
    Impl& operator=(const Impl&) = delete;
    Impl(const Impl&)            = delete;
 
    // Copy state to local memory for efficiency.
    __device__ explicit Impl(RandGenerator<gpu, DType>* gen, int state_idx)
        : global_gen_(gen), global_state_idx_(state_idx), state_(*(gen->states_ + state_idx)) {}
 
    __device__ ~Impl() {
      // store the curand state back into global memory
      global_gen_->states_[global_state_idx_] = state_;
    }
 
    MSHADOW_FORCE_INLINE __device__ int rand() {
      return curand(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ int64_t rand_int64() {
      return static_cast<int64_t>(curand(&state_) << 31) + curand(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ float uniform() {
      return static_cast<float>(1.0) - curand_uniform(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ float normal() {
      return curand_normal(&state_);
    }
 
   private:
    RandGenerator<gpu, DType>* global_gen_;
    int global_state_idx_;
    curandStatePhilox4_32_10_t state_;
  };  // class RandGenerator<gpu, DType>::Impl
 
  static void AllocState(RandGenerator<gpu, DType>* inst);
 
  static void FreeState(RandGenerator<gpu, DType>* inst);
 
  void Seed(mshadow::Stream<gpu>* s, uint32_t seed);
 
  // export global random states, used by c++ custom operator
  void* GetStates();
 
 private:
  curandStatePhilox4_32_10_t* states_;
};  // class RandGenerator<gpu, DType>
 
template <>
class RandGenerator<gpu, double> {
 public:
  // uniform number generation in Cuda made consistent with stl (include 0 but exclude 1)
  // by using 1.0-curand_uniform().
  // Needed as some samplers in sampler.h won't be able to deal with
  // one of the boundary cases.
  // TODO(alexzai): move impl class to separate file - tracked in MXNET-948
  class Impl {
   public:
    Impl& operator=(const Impl&) = delete;
    Impl(const Impl&)            = delete;
 
    // Copy state to local memory for efficiency.
    __device__ explicit Impl(RandGenerator<gpu, double>* gen, int state_idx)
        : global_gen_(gen), global_state_idx_(state_idx), state_(*(gen->states_ + state_idx)) {}
 
    __device__ ~Impl() {
      // store the curand state back into global memory
      global_gen_->states_[global_state_idx_] = state_;
    }
 
    MSHADOW_FORCE_INLINE __device__ int rand() {
      return curand(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ int64_t rand_int64() {
      return static_cast<int64_t>(curand(&state_) << 31) + curand(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ double uniform() {
      return static_cast<float>(1.0) - curand_uniform_double(&state_);
    }
 
    MSHADOW_FORCE_INLINE __device__ double normal() {
      return curand_normal_double(&state_);
    }
 
   private:
    RandGenerator<gpu, double>* global_gen_;
    int global_state_idx_;
    curandStatePhilox4_32_10_t state_;
  };  // class RandGenerator<gpu, double>::Impl
 
 private:
  curandStatePhilox4_32_10_t* states_;
};  // class RandGenerator<gpu, double>
 
#endif  // MXNET_USE_CUDA
 
}  // namespace random
}  // namespace common
}  // namespace mxnet
#endif  // MXNET_RANDOM_GENERATOR_H_