diff --git a/src/cunumeric/random/bitgenerator.cc b/src/cunumeric/random/bitgenerator.cc index 1038a5309..eadbb18a3 100644 --- a/src/cunumeric/random/bitgenerator.cc +++ b/src/cunumeric/random/bitgenerator.cc @@ -14,11 +14,17 @@ * */ +// MacOS host variant: +// +#if defined(__APPLE__) && defined(__MACH__) +#define USE_STL_RANDOM_ENGINE_ +#endif + #include "cunumeric/random/bitgenerator.h" #include "cunumeric/random/bitgenerator_template.inl" #include "cunumeric/random/bitgenerator_util.h" -#include "cunumeric/random/curand_help.h" +#include "cunumeric/random/rnd_types.h" #include "cunumeric/random/randutil/randutil.h" #include "cunumeric/random/bitgenerator_curand.inl" @@ -31,6 +37,16 @@ static Logger log_curand("cunumeric.random"); Logger& randutil_log() { return log_curand; } +#ifdef USE_STL_RANDOM_ENGINE_ +void randutil_check_status(rnd_status_t error, const char* file, int line) +{ + if (error) { + randutil_log().fatal() << "Internal random engine failure with error " << (int)error + << " in file " << file << " at line " << line; + assert(false); + } +} +#else void randutil_check_curand(curandStatus_t error, const char* file, int line) { if (error != CURAND_STATUS_SUCCESS) { @@ -39,15 +55,16 @@ void randutil_check_curand(curandStatus_t error, const char* file, int line) assert(false); } } +#endif struct CPUGenerator : public CURANDGenerator { CPUGenerator(BitGeneratorType gentype, uint64_t seed, uint64_t generatorId, uint32_t flags) : CURANDGenerator(gentype, seed, generatorId) { - CHECK_CURAND(::randutilCreateGeneratorHost(&gen_, type_, seed, generatorId, flags)); + CHECK_RND_ENGINE(::randutilCreateGeneratorHost(&gen_, type_, seed, generatorId, flags)); } - virtual ~CPUGenerator() { CHECK_CURAND(::randutilDestroyGenerator(gen_)); } + virtual ~CPUGenerator() { CHECK_RND_ENGINE(::randutilDestroyGenerator(gen_)); } }; template <> diff --git a/src/cunumeric/random/bitgenerator.cu b/src/cunumeric/random/bitgenerator.cu index f63ed3ef1..4edb7a47f 100644 --- a/src/cunumeric/random/bitgenerator.cu +++ b/src/cunumeric/random/bitgenerator.cu @@ -27,19 +27,30 @@ namespace cunumeric { using namespace legate; +// required by CHECK_CURAND_DEVICE: +// +void randutil_check_curand_device(curandStatus_t error, const char* file, int line) +{ + if (error != CURAND_STATUS_SUCCESS) { + randutil_log().fatal() << "Internal CURAND failure with error " << (int)error << " in file " + << file << " at line " << line; + assert(false); + } +} + struct GPUGenerator : public CURANDGenerator { cudaStream_t stream_; GPUGenerator(BitGeneratorType gentype, uint64_t seed, uint64_t generatorId, uint32_t flags) : CURANDGenerator(gentype, seed, generatorId) { CHECK_CUDA(::cudaStreamCreate(&stream_)); - CHECK_CURAND(::randutilCreateGenerator(&gen_, type_, seed, generatorId, flags, stream_)); + CHECK_CURAND_DEVICE(::randutilCreateGenerator(&gen_, type_, seed, generatorId, flags, stream_)); } virtual ~GPUGenerator() { CHECK_CUDA(::cudaStreamSynchronize(stream_)); - CHECK_CURAND(::randutilDestroyGenerator(gen_)); + CHECK_CURAND_DEVICE(::randutilDestroyGenerator(gen_)); } }; diff --git a/src/cunumeric/random/bitgenerator_curand.inl b/src/cunumeric/random/bitgenerator_curand.inl index f74c87ad0..fab6fb1f4 100644 --- a/src/cunumeric/random/bitgenerator_curand.inl +++ b/src/cunumeric/random/bitgenerator_curand.inl @@ -25,7 +25,7 @@ #include "cunumeric/random/bitgenerator_template.inl" #include "cunumeric/random/bitgenerator_util.h" -#include "cunumeric/random/curand_help.h" +#include "cunumeric/random/rnd_types.h" #include "cunumeric/random/randutil/randutil.h" namespace cunumeric { @@ -41,11 +41,11 @@ struct CURANDGenerator { randutilGenerator_t gen_; uint64_t seed_; uint64_t generatorId_; - curandRngType type_; + randRngType type_; protected: CURANDGenerator(BitGeneratorType gentype, uint64_t seed, uint64_t generatorId) - : type_(get_curandRngType(gentype)), seed_(seed), generatorId_(generatorId) + : type_(get_rndRngType(gentype)), seed_(seed), generatorId_(generatorId) { randutil_log().debug() << "CURANDGenerator::create"; } @@ -57,217 +57,218 @@ struct CURANDGenerator { void generate_raw(uint64_t count, uint32_t* out) { - CHECK_CURAND(::randutilGenerateRawUInt32(gen_, out, count)); + CHECK_RND_ENGINE(::randutilGenerateRawUInt32(gen_, out, count)); } void generate_integer_64(uint64_t count, int64_t* out, int64_t low, int64_t high) { - CHECK_CURAND(::randutilGenerateIntegers64(gen_, out, count, low, high)); + CHECK_RND_ENGINE(::randutilGenerateIntegers64(gen_, out, count, low, high)); } void generate_integer_16(uint64_t count, int16_t* out, int16_t low, int16_t high) { - CHECK_CURAND(::randutilGenerateIntegers16(gen_, out, count, low, high)); + CHECK_RND_ENGINE(::randutilGenerateIntegers16(gen_, out, count, low, high)); } void generate_integer_32(uint64_t count, int32_t* out, int32_t low, int32_t high) { - CHECK_CURAND(::randutilGenerateIntegers32(gen_, out, count, low, high)); + CHECK_RND_ENGINE(::randutilGenerateIntegers32(gen_, out, count, low, high)); } void generate_uniform_64(uint64_t count, double* out, double low, double high) { - CHECK_CURAND(::randutilGenerateUniformDoubleEx(gen_, out, count, low, high)); + CHECK_RND_ENGINE(::randutilGenerateUniformDoubleEx(gen_, out, count, low, high)); } void generate_uniform_32(uint64_t count, float* out, float low, float high) { - CHECK_CURAND(::randutilGenerateUniformEx(gen_, out, count, low, high)); + CHECK_RND_ENGINE(::randutilGenerateUniformEx(gen_, out, count, low, high)); } void generate_lognormal_64(uint64_t count, double* out, double mean, double stdev) { - CHECK_CURAND(::randutilGenerateLogNormalDoubleEx(gen_, out, count, mean, stdev)); + CHECK_RND_ENGINE(::randutilGenerateLogNormalDoubleEx(gen_, out, count, mean, stdev)); } void generate_lognormal_32(uint64_t count, float* out, float mean, float stdev) { - CHECK_CURAND(::randutilGenerateLogNormalEx(gen_, out, count, mean, stdev)); + CHECK_RND_ENGINE(::randutilGenerateLogNormalEx(gen_, out, count, mean, stdev)); } void generate_normal_64(uint64_t count, double* out, double mean, double stdev) { - CHECK_CURAND(::randutilGenerateNormalDoubleEx(gen_, out, count, mean, stdev)); + CHECK_RND_ENGINE(::randutilGenerateNormalDoubleEx(gen_, out, count, mean, stdev)); } void generate_normal_32(uint64_t count, float* out, float mean, float stdev) { - CHECK_CURAND(::randutilGenerateNormalEx(gen_, out, count, mean, stdev)); + CHECK_RND_ENGINE(::randutilGenerateNormalEx(gen_, out, count, mean, stdev)); } void generate_poisson(uint64_t count, uint32_t* out, double lam) { - CHECK_CURAND(::randutilGeneratePoissonEx(gen_, out, count, lam)); + CHECK_RND_ENGINE(::randutilGeneratePoissonEx(gen_, out, count, lam)); } void generate_exponential_64(uint64_t count, double* out, double scale) { - CHECK_CURAND(::randutilGenerateExponentialDoubleEx(gen_, out, count, scale)); + CHECK_RND_ENGINE(::randutilGenerateExponentialDoubleEx(gen_, out, count, scale)); } void generate_exponential_32(uint64_t count, float* out, float scale) { - CHECK_CURAND(::randutilGenerateExponentialEx(gen_, out, count, scale)); + CHECK_RND_ENGINE(::randutilGenerateExponentialEx(gen_, out, count, scale)); } void generate_gumbel_64(uint64_t count, double* out, double mu, double beta) { - CHECK_CURAND(::randutilGenerateGumbelDoubleEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateGumbelDoubleEx(gen_, out, count, mu, beta)); } void generate_gumbel_32(uint64_t count, float* out, float mu, float beta) { - CHECK_CURAND(::randutilGenerateGumbelEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateGumbelEx(gen_, out, count, mu, beta)); } void generate_laplace_64(uint64_t count, double* out, double mu, double beta) { - CHECK_CURAND(::randutilGenerateLaplaceDoubleEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateLaplaceDoubleEx(gen_, out, count, mu, beta)); } void generate_laplace_32(uint64_t count, float* out, float mu, float beta) { - CHECK_CURAND(::randutilGenerateLaplaceEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateLaplaceEx(gen_, out, count, mu, beta)); } void generate_logistic_64(uint64_t count, double* out, double mu, double beta) { - CHECK_CURAND(::randutilGenerateLogisticDoubleEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateLogisticDoubleEx(gen_, out, count, mu, beta)); } void generate_logistic_32(uint64_t count, float* out, float mu, float beta) { - CHECK_CURAND(::randutilGenerateLogisticEx(gen_, out, count, mu, beta)); + CHECK_RND_ENGINE(::randutilGenerateLogisticEx(gen_, out, count, mu, beta)); } void generate_pareto_64(uint64_t count, double* out, double alpha) { - CHECK_CURAND(::randutilGenerateParetoDoubleEx(gen_, out, count, 1.0, alpha)); + CHECK_RND_ENGINE(::randutilGenerateParetoDoubleEx(gen_, out, count, 1.0, alpha)); } void generate_pareto_32(uint64_t count, float* out, float alpha) { - CHECK_CURAND(::randutilGenerateParetoEx(gen_, out, count, 1.0f, alpha)); + CHECK_RND_ENGINE(::randutilGenerateParetoEx(gen_, out, count, 1.0f, alpha)); } void generate_power_64(uint64_t count, double* out, double alpha) { - CHECK_CURAND(::randutilGeneratePowerDoubleEx(gen_, out, count, alpha)); + CHECK_RND_ENGINE(::randutilGeneratePowerDoubleEx(gen_, out, count, alpha)); } void generate_power_32(uint64_t count, float* out, float alpha) { - CHECK_CURAND(::randutilGeneratePowerEx(gen_, out, count, alpha)); + CHECK_RND_ENGINE(::randutilGeneratePowerEx(gen_, out, count, alpha)); } void generate_rayleigh_64(uint64_t count, double* out, double sigma) { - CHECK_CURAND(::randutilGenerateRayleighDoubleEx(gen_, out, count, sigma)); + CHECK_RND_ENGINE(::randutilGenerateRayleighDoubleEx(gen_, out, count, sigma)); } void generate_rayleigh_32(uint64_t count, float* out, float sigma) { - CHECK_CURAND(::randutilGenerateRayleighEx(gen_, out, count, sigma)); + CHECK_RND_ENGINE(::randutilGenerateRayleighEx(gen_, out, count, sigma)); } void generate_cauchy_64(uint64_t count, double* out, double x0, double gamma) { - CHECK_CURAND(::randutilGenerateCauchyDoubleEx(gen_, out, count, x0, gamma)); + CHECK_RND_ENGINE(::randutilGenerateCauchyDoubleEx(gen_, out, count, x0, gamma)); } void generate_cauchy_32(uint64_t count, float* out, float x0, float gamma) { - CHECK_CURAND(::randutilGenerateCauchyEx(gen_, out, count, x0, gamma)); + CHECK_RND_ENGINE(::randutilGenerateCauchyEx(gen_, out, count, x0, gamma)); } void generate_triangular_64(uint64_t count, double* out, double a, double b, double c) { - CHECK_CURAND(::randutilGenerateTriangularDoubleEx(gen_, out, count, a, b, c)); + CHECK_RND_ENGINE(::randutilGenerateTriangularDoubleEx(gen_, out, count, a, b, c)); } void generate_triangular_32(uint64_t count, float* out, float a, float b, float c) { - CHECK_CURAND(::randutilGenerateTriangularEx(gen_, out, count, a, b, c)); + CHECK_RND_ENGINE(::randutilGenerateTriangularEx(gen_, out, count, a, b, c)); } void generate_weibull_64(uint64_t count, double* out, double lam, double k) { - CHECK_CURAND(::randutilGenerateWeibullDoubleEx(gen_, out, count, lam, k)); + CHECK_RND_ENGINE(::randutilGenerateWeibullDoubleEx(gen_, out, count, lam, k)); } void generate_weibull_32(uint64_t count, float* out, float lam, float k) { - CHECK_CURAND(::randutilGenerateWeibullEx(gen_, out, count, lam, k)); + CHECK_RND_ENGINE(::randutilGenerateWeibullEx(gen_, out, count, lam, k)); } void generate_beta_64(uint64_t count, double* out, double a, double b) { - CHECK_CURAND(::randutilGenerateBetaDoubleEx(gen_, out, count, a, b)); + CHECK_RND_ENGINE(::randutilGenerateBetaDoubleEx(gen_, out, count, a, b)); } void generate_beta_32(uint64_t count, float* out, float a, float b) { - CHECK_CURAND(::randutilGenerateBetaEx(gen_, out, count, a, b)); + CHECK_RND_ENGINE(::randutilGenerateBetaEx(gen_, out, count, a, b)); } void generate_f_64(uint64_t count, double* out, double dfnum, double dfden) { - CHECK_CURAND(::randutilGenerateFisherSnedecorDoubleEx(gen_, out, count, dfnum, dfden)); + CHECK_RND_ENGINE(::randutilGenerateFisherSnedecorDoubleEx(gen_, out, count, dfnum, dfden)); } void generate_f_32(uint64_t count, float* out, float dfnum, float dfden) { - CHECK_CURAND(::randutilGenerateFisherSnedecorEx(gen_, out, count, dfnum, dfden)); + CHECK_RND_ENGINE(::randutilGenerateFisherSnedecorEx(gen_, out, count, dfnum, dfden)); } void generate_logseries(uint64_t count, uint32_t* out, double p) { - CHECK_CURAND(::randutilGenerateLogSeriesEx(gen_, out, count, p)); + CHECK_RND_ENGINE(::randutilGenerateLogSeriesEx(gen_, out, count, p)); } void generate_noncentral_f_64( uint64_t count, double* out, double dfnum, double dfden, double nonc) { - CHECK_CURAND(::randutilGenerateFisherSnedecorDoubleEx(gen_, out, count, dfnum, dfden, nonc)); + CHECK_RND_ENGINE( + ::randutilGenerateFisherSnedecorDoubleEx(gen_, out, count, dfnum, dfden, nonc)); } void generate_noncentral_f_32(uint64_t count, float* out, float dfnum, float dfden, float nonc) { - CHECK_CURAND(::randutilGenerateFisherSnedecorEx(gen_, out, count, dfnum, dfden, nonc)); + CHECK_RND_ENGINE(::randutilGenerateFisherSnedecorEx(gen_, out, count, dfnum, dfden, nonc)); } void generate_chisquare_64(uint64_t count, double* out, double df, double nonc) { - CHECK_CURAND(::randutilGenerateChiSquareDoubleEx(gen_, out, count, df, nonc)); + CHECK_RND_ENGINE(::randutilGenerateChiSquareDoubleEx(gen_, out, count, df, nonc)); } void generate_chisquare_32(uint64_t count, float* out, float df, float nonc) { - CHECK_CURAND(::randutilGenerateChiSquareEx(gen_, out, count, df, nonc)); + CHECK_RND_ENGINE(::randutilGenerateChiSquareEx(gen_, out, count, df, nonc)); } void generate_gamma_64(uint64_t count, double* out, double k, double theta) { - CHECK_CURAND(::randutilGenerateGammaDoubleEx(gen_, out, count, k, theta)); + CHECK_RND_ENGINE(::randutilGenerateGammaDoubleEx(gen_, out, count, k, theta)); } void generate_gamma_32(uint64_t count, float* out, float k, float theta) { - CHECK_CURAND(::randutilGenerateGammaEx(gen_, out, count, k, theta)); + CHECK_RND_ENGINE(::randutilGenerateGammaEx(gen_, out, count, k, theta)); } void generate_standard_t_64(uint64_t count, double* out, double df) { - CHECK_CURAND(::randutilGenerateStandardTDoubleEx(gen_, out, count, df)); + CHECK_RND_ENGINE(::randutilGenerateStandardTDoubleEx(gen_, out, count, df)); } void generate_standard_t_32(uint64_t count, float* out, float df) { - CHECK_CURAND(::randutilGenerateStandardTEx(gen_, out, count, df)); + CHECK_RND_ENGINE(::randutilGenerateStandardTEx(gen_, out, count, df)); } void generate_hypergeometric( uint64_t count, uint32_t* out, int64_t ngood, int64_t nbad, int64_t nsample) { - CHECK_CURAND(::randutilGenerateHyperGeometricEx(gen_, out, count, ngood, nbad, nsample)); + CHECK_RND_ENGINE(::randutilGenerateHyperGeometricEx(gen_, out, count, ngood, nbad, nsample)); } void generate_vonmises_64(uint64_t count, double* out, double mu, double kappa) { - CHECK_CURAND(::randutilGenerateVonMisesDoubleEx(gen_, out, count, mu, kappa)); + CHECK_RND_ENGINE(::randutilGenerateVonMisesDoubleEx(gen_, out, count, mu, kappa)); } void generate_vonmises_32(uint64_t count, float* out, float mu, float kappa) { - CHECK_CURAND(::randutilGenerateVonMisesEx(gen_, out, count, mu, kappa)); + CHECK_RND_ENGINE(::randutilGenerateVonMisesEx(gen_, out, count, mu, kappa)); } void generate_zipf(uint64_t count, uint32_t* out, double a) { - CHECK_CURAND(::randutilGenerateZipfEx(gen_, out, count, a)); + CHECK_RND_ENGINE(::randutilGenerateZipfEx(gen_, out, count, a)); } void generate_geometric(uint64_t count, uint32_t* out, double p) { - CHECK_CURAND(::randutilGenerateGeometricEx(gen_, out, count, p)); + CHECK_RND_ENGINE(::randutilGenerateGeometricEx(gen_, out, count, p)); } void generate_wald_64(uint64_t count, double* out, double mean, double scale) { - CHECK_CURAND(::randutilGenerateWaldDoubleEx(gen_, out, count, mean, scale)); + CHECK_RND_ENGINE(::randutilGenerateWaldDoubleEx(gen_, out, count, mean, scale)); } void generate_wald_32(uint64_t count, float* out, float mean, float scale) { - CHECK_CURAND(::randutilGenerateWaldEx(gen_, out, count, mean, scale)); + CHECK_RND_ENGINE(::randutilGenerateWaldEx(gen_, out, count, mean, scale)); } void generate_binomial(uint64_t count, uint32_t* out, uint32_t ntrials, double p) { - CHECK_CURAND(::randutilGenerateBinomialEx(gen_, out, count, ntrials, p)); + CHECK_RND_ENGINE(::randutilGenerateBinomialEx(gen_, out, count, ntrials, p)); } void generate_negative_binomial(uint64_t count, uint32_t* out, uint32_t ntrials, double p) { - CHECK_CURAND(::randutilGenerateNegativeBinomialEx(gen_, out, count, ntrials, p)); + CHECK_RND_ENGINE(::randutilGenerateNegativeBinomialEx(gen_, out, count, ntrials, p)); } }; diff --git a/src/cunumeric/random/curand_help.h b/src/cunumeric/random/curand_help.h index dfa457b2c..c02f13ed6 100644 --- a/src/cunumeric/random/curand_help.h +++ b/src/cunumeric/random/curand_help.h @@ -24,10 +24,26 @@ randutil_check_curand(__result__, __FILE__, __LINE__); \ } while (false) +// necessary, b/c the STL variant (host only MacOS) uses non-curand abstractions, +// hence a different checker defined in bitgenerator.cc, while the curand checker +// gets undefined; however the device code still requires the curand checker and +// attempts to link against the definition in bitgenerator.cc, which is disabled +// in this situation; the checker below fulfills that purpose: +// +#define CHECK_CURAND_DEVICE(expr) \ + do { \ + curandStatus_t __result__ = (expr); \ + randutil_check_curand_device(__result__, __FILE__, __LINE__); \ + } while (false) + namespace cunumeric { legate::Logger& randutil_log(); void randutil_check_curand(curandStatus_t error, const char* file, int line); +// required by CHECK_CURAND_DEVICE: +// +void randutil_check_curand_device(curandStatus_t error, const char* file, int line); + static inline curandRngType get_curandRngType(cunumeric::BitGeneratorType kind) { switch (kind) { diff --git a/src/cunumeric/random/randutil/generator.h b/src/cunumeric/random/randutil/generator.h index c56936503..466c32ec0 100644 --- a/src/cunumeric/random/randutil/generator.h +++ b/src/cunumeric/random/randutil/generator.h @@ -23,6 +23,8 @@ #include "randutil_curand.h" #include "randutil_impl.h" +#include "cunumeric/random/rnd_aliases.h" + namespace randutilimpl { struct basegenerator { @@ -36,33 +38,42 @@ template struct generatorid; template <> -struct generatorid { - static constexpr int rng_type = CURAND_RNG_PSEUDO_XORWOW; +struct generatorid { + static constexpr int rng_type = RND_RNG_PSEUDO_XORWOW; }; + +#ifndef USE_STL_RANDOM_ENGINE_ +// Curand *different* specializations, not possible with only one generator +// template <> -struct generatorid { - static constexpr int rng_type = CURAND_RNG_PSEUDO_PHILOX4_32_10; +struct generatorid { + static constexpr int rng_type = RND_RNG_PSEUDO_PHILOX4_32_10; }; template <> -struct generatorid { - static constexpr int rng_type = CURAND_RNG_PSEUDO_MRG32K3A; +struct generatorid { + static constexpr int rng_type = RND_RNG_PSEUDO_MRG32K3A; }; +#endif +#if 0 +// this seems dead code +// template struct generatortype; template <> -struct generatortype { - using rng_t = curandStateXORWOW_t; +struct generatortype { + using rng_t = gen_XORWOW_t; }; template <> -struct generatortype { - using rng_t = curandStatePhilox4_32_10_t; +struct generatortype { + using rng_t = gen_Philox4_32_10_t; }; template <> -struct generatortype { - using rng_t = curandStateMRG32k3a_t; +struct generatortype { + using rng_t = gen_MRG32k3a_t; }; +#endif template struct inner_generator : basegenerator { @@ -70,10 +81,19 @@ struct inner_generator : basegenerator uint64_t generatorID; gen_t generator; - inner_generator(uint64_t seed, uint64_t generatorID, cudaStream_t ignored) - : seed(seed), generatorID(generatorID) + inner_generator(uint64_t seed, uint64_t generatorID, stream_t ignored) + : seed(seed), + generatorID(generatorID) +#ifdef USE_STL_RANDOM_ENGINE_ + , + generator(seed) +#endif { +#ifdef USE_STL_RANDOM_ENGINE_ + std::srand(seed); +#else curand_init(seed, generatorID, 0, &generator); +#endif } virtual void destroy() override {} @@ -93,36 +113,36 @@ struct inner_generator : basegenerator }; template -curandStatus_t inner_dispatch_sample(basegenerator* gen, func_t func, size_t N, out_t* out) +rnd_status_t inner_dispatch_sample(basegenerator* gen, func_t func, size_t N, out_t* out) { switch (gen->generatorTypeId()) { - case CURAND_RNG_PSEUDO_XORWOW: - return static_cast*>(gen) + case RND_RNG_PSEUDO_XORWOW: + return static_cast*>(gen) ->template draw(func, N, out); - case CURAND_RNG_PSEUDO_PHILOX4_32_10: - return static_cast*>(gen) + case RND_RNG_PSEUDO_PHILOX4_32_10: + return static_cast*>(gen) ->template draw(func, N, out); - case CURAND_RNG_PSEUDO_MRG32K3A: - return static_cast*>(gen) + case RND_RNG_PSEUDO_MRG32K3A: + return static_cast*>(gen) ->template draw(func, N, out); default: LEGATE_ABORT; } - return CURAND_STATUS_INTERNAL_ERROR; + return RND_STATUS_INTERNAL_ERROR; } // template funtion with HOST and DEVICE implementations template struct dispatcher { - static curandStatus_t run(randutilimpl::basegenerator* generator, - func_t func, - size_t N, - out_t* out); + static rnd_status_t run(randutilimpl::basegenerator* generator, + func_t func, + size_t N, + out_t* out); }; // HOST-side template instantiation of generator template struct dispatcher { - static curandStatus_t run(randutilimpl::basegenerator* gen, func_t func, size_t N, out_t* out) + static rnd_status_t run(randutilimpl::basegenerator* gen, func_t func, size_t N, out_t* out) { return inner_dispatch_sample( gen, func, N, out); @@ -130,7 +150,7 @@ struct dispatcher { }; template -curandStatus_t dispatch(randutilimpl::basegenerator* gen, func_t func, size_t N, out_t* out) +rnd_status_t dispatch(randutilimpl::basegenerator* gen, func_t func, size_t N, out_t* out) { switch (gen->location()) { case randutilimpl::execlocation::HOST: @@ -141,7 +161,7 @@ curandStatus_t dispatch(randutilimpl::basegenerator* gen, func_t func, size_t N, #endif default: LEGATE_ABORT; } - return CURAND_STATUS_INTERNAL_ERROR; + return RND_STATUS_INTERNAL_ERROR; } } // namespace randutilimpl diff --git a/src/cunumeric/random/randutil/generator_cauchy.inl b/src/cunumeric/random/randutil/generator_cauchy.inl index f0b5e2b12..ece9843d8 100644 --- a/src/cunumeric/random/randutil/generator_cauchy.inl +++ b/src/cunumeric/random/randutil/generator_cauchy.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct cauchy_t; @@ -28,7 +30,7 @@ struct cauchy_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 return x0 + gamma * ::tanf(pi * (y - 0.5f)); } }; @@ -42,7 +44,7 @@ struct cauchy_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 + double y = randutilimpl::engine_uniform(gen); // y cannot be 0 return x0 + gamma * ::tan(pi * (y - 0.5)); } }; diff --git a/src/cunumeric/random/randutil/generator_create.inl b/src/cunumeric/random/randutil/generator_create.inl index 2f9cdf29b..361b2ec32 100644 --- a/src/cunumeric/random/randutil/generator_create.inl +++ b/src/cunumeric/random/randutil/generator_create.inl @@ -17,34 +17,32 @@ #include "generator.h" template -curandStatus_t randutilGenerator(randutilGenerator_t* generator, - uint64_t seed, - uint64_t generatorID, - cudaStream_t stream = nullptr) +rnd_status_t randutilGenerator(randutilGenerator_t* generator, + uint64_t seed, + uint64_t generatorID, + stream_t stream = nullptr) { randutilimpl::inner_generator* result = new randutilimpl::inner_generator(seed, generatorID, stream); *generator = (randutilGenerator_t)result; - return CURAND_STATUS_SUCCESS; + return RND_STATUS_SUCCESS; } template -static curandStatus_t inner_randutilCreateGenerator(randutilGenerator_t* generator, - curandRngType_t rng_type, - uint64_t seed, - uint64_t generatorID, - cudaStream_t stream = nullptr) +static rnd_status_t inner_randutilCreateGenerator(randutilGenerator_t* generator, + randRngType_t rng_type, + uint64_t seed, + uint64_t generatorID, + stream_t stream = nullptr) { - switch (rng_type) { - case CURAND_RNG_PSEUDO_XORWOW: - return randutilGenerator(generator, seed, generatorID, stream); - case CURAND_RNG_PSEUDO_PHILOX4_32_10: - return randutilGenerator( - generator, seed, generatorID, stream); - case CURAND_RNG_PSEUDO_MRG32K3A: - return randutilGenerator( - generator, seed, generatorID, stream); + switch (static_cast(rng_type)) { + case RND_RNG_PSEUDO_XORWOW: + return randutilGenerator(generator, seed, generatorID, stream); + case RND_RNG_PSEUDO_PHILOX4_32_10: + return randutilGenerator(generator, seed, generatorID, stream); + case RND_RNG_PSEUDO_MRG32K3A: + return randutilGenerator(generator, seed, generatorID, stream); default: LEGATE_ABORT; } - return CURAND_STATUS_TYPE_ERROR; + return RND_STATUS_TYPE_ERROR; } diff --git a/src/cunumeric/random/randutil/generator_exponential.inl b/src/cunumeric/random/randutil/generator_exponential.inl index 5bc4847b2..4a01ff40e 100644 --- a/src/cunumeric/random/randutil/generator_exponential.inl +++ b/src/cunumeric/random/randutil/generator_exponential.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct exponential_t; @@ -26,7 +28,7 @@ struct exponential_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float uni = curand_uniform(&gen); + float uni = randutilimpl::engine_uniform(gen); return -::logf(uni) * scale; } }; @@ -38,7 +40,7 @@ struct exponential_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double uni = curand_uniform_double(&gen); + double uni = randutilimpl::engine_uniform(gen); return -::logf(uni) * scale; } }; diff --git a/src/cunumeric/random/randutil/generator_gumbel.inl b/src/cunumeric/random/randutil/generator_gumbel.inl index 4a5d7b0b1..32a3b0d05 100644 --- a/src/cunumeric/random/randutil/generator_gumbel.inl +++ b/src/cunumeric/random/randutil/generator_gumbel.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct gumbel_t; @@ -27,7 +29,8 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be zero + auto y = randutilimpl::engine_uniform(gen); // y cannot be zero + if (y == 1.0f) return mu; float lny = ::logf(y); return mu - beta * ::logf(-lny); @@ -41,7 +44,8 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be zero + auto y = randutilimpl::engine_uniform(gen); // y cannot be zero + if (y == 1.0) return mu; double lny = ::log(y); return mu - beta * ::log(-lny); diff --git a/src/cunumeric/random/randutil/generator_host.cc b/src/cunumeric/random/randutil/generator_host.cc index 50bcfeefc..93330c1ab 100644 --- a/src/cunumeric/random/randutil/generator_host.cc +++ b/src/cunumeric/random/randutil/generator_host.cc @@ -14,8 +14,15 @@ * */ +// MacOS host variant: +// +#if defined(__APPLE__) && defined(__MACH__) +#define USE_STL_RANDOM_ENGINE_ +#endif + #include "generator.h" #include "generator_create.inl" +#include "cunumeric/random/rnd_aliases.h" #if !defined(LEGATE_USE_CUDA) // the host code of cuRAND try to extern these variables out of nowhere, @@ -24,29 +31,29 @@ const dim3 blockDim{}; const uint3 threadIdx{}; #endif -extern "C" curandStatus_t randutilCreateGeneratorHost(randutilGenerator_t* generator, - curandRngType_t rng_type, - uint64_t seed, - uint64_t generatorID, - uint32_t flags) +extern "C" rnd_status_t randutilCreateGeneratorHost(randutilGenerator_t* generator, + randRngType_t rng_type, + uint64_t seed, + uint64_t generatorID, + uint32_t flags) { return inner_randutilCreateGenerator( generator, rng_type, seed, generatorID, nullptr); } -extern "C" curandStatus_t randutilDestroyGenerator(randutilGenerator_t generator) +extern "C" rnd_status_t randutilDestroyGenerator(randutilGenerator_t generator) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; gen->destroy(); delete gen; - return CURAND_STATUS_SUCCESS; + return RND_STATUS_SUCCESS; } #pragma region integers #include "generator_integers.inl" -extern "C" curandStatus_t randutilGenerateIntegers16( +extern "C" rnd_status_t randutilGenerateIntegers16( randutilGenerator_t generator, int16_t* outputPtr, size_t n, int16_t low, int16_t high) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -56,7 +63,7 @@ extern "C" curandStatus_t randutilGenerateIntegers16( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateIntegers32( +extern "C" rnd_status_t randutilGenerateIntegers32( randutilGenerator_t generator, int32_t* outputPtr, size_t n, int32_t low, int32_t high) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -66,7 +73,7 @@ extern "C" curandStatus_t randutilGenerateIntegers32( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateIntegers64( +extern "C" rnd_status_t randutilGenerateIntegers64( randutilGenerator_t generator, int64_t* outputPtr, size_t n, int64_t low, int64_t high) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -82,7 +89,7 @@ extern "C" curandStatus_t randutilGenerateIntegers64( #include "generator_lognormal.inl" -extern "C" curandStatus_t randutilGenerateLogNormalEx( +extern "C" rnd_status_t randutilGenerateLogNormalEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mean, float stddev) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -92,7 +99,7 @@ extern "C" curandStatus_t randutilGenerateLogNormalEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateLogNormalDoubleEx( +extern "C" rnd_status_t randutilGenerateLogNormalDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mean, double stddev) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -108,7 +115,7 @@ extern "C" curandStatus_t randutilGenerateLogNormalDoubleEx( #include "generator_normal.inl" -extern "C" curandStatus_t randutilGenerateNormalEx( +extern "C" rnd_status_t randutilGenerateNormalEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mean, float stddev) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -118,7 +125,7 @@ extern "C" curandStatus_t randutilGenerateNormalEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateNormalDoubleEx( +extern "C" rnd_status_t randutilGenerateNormalDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mean, double stddev) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -134,10 +141,10 @@ extern "C" curandStatus_t randutilGenerateNormalDoubleEx( #include "generator_poisson.inl" -extern "C" curandStatus_t randutilGeneratePoissonEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double lambda) +extern "C" rnd_status_t randutilGeneratePoissonEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double lambda) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; poisson func; @@ -151,9 +158,9 @@ extern "C" curandStatus_t randutilGeneratePoissonEx(randutilGenerator_t generato #include "generator_raw.inl" -extern "C" curandStatus_t randutilGenerateRawUInt32(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n) +extern "C" rnd_status_t randutilGenerateRawUInt32(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; raw func; @@ -166,7 +173,7 @@ extern "C" curandStatus_t randutilGenerateRawUInt32(randutilGenerator_t generato #include "generator_uniform.inl" -extern "C" curandStatus_t randutilGenerateUniformEx( +extern "C" rnd_status_t randutilGenerateUniformEx( randutilGenerator_t generator, float* outputPtr, size_t n, float low, float high) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -177,7 +184,7 @@ extern "C" curandStatus_t randutilGenerateUniformEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateUniformDoubleEx( +extern "C" rnd_status_t randutilGenerateUniformDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double low, double high) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; diff --git a/src/cunumeric/random/randutil/generator_host_advanced.cc b/src/cunumeric/random/randutil/generator_host_advanced.cc index 6f836403f..af3788933 100644 --- a/src/cunumeric/random/randutil/generator_host_advanced.cc +++ b/src/cunumeric/random/randutil/generator_host_advanced.cc @@ -14,13 +14,19 @@ * */ +// MacOS host variant: +// +#if defined(__APPLE__) && defined(__MACH__) +#define USE_STL_RANDOM_ENGINE_ +#endif + #include "generator.h" #pragma region beta #include "generator_beta.inl" -extern "C" curandStatus_t randutilGenerateBetaEx( +extern "C" rnd_status_t randutilGenerateBetaEx( randutilGenerator_t generator, float* outputPtr, size_t n, float a, float b) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -30,7 +36,7 @@ extern "C" curandStatus_t randutilGenerateBetaEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateBetaDoubleEx( +extern "C" rnd_status_t randutilGenerateBetaDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double a, double b) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -46,12 +52,12 @@ extern "C" curandStatus_t randutilGenerateBetaDoubleEx( #include "generator_f.inl" -extern "C" curandStatus_t randutilGenerateFisherSnedecorEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float dfnum, - float dfden, - float nonc) // 0.0f is F distribution +extern "C" rnd_status_t randutilGenerateFisherSnedecorEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float dfnum, + float dfden, + float nonc) // 0.0f is F distribution { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; if (nonc == 0.0f) { @@ -68,7 +74,7 @@ extern "C" curandStatus_t randutilGenerateFisherSnedecorEx(randutilGenerator_t g } } -extern "C" curandStatus_t randutilGenerateFisherSnedecorDoubleEx( +extern "C" rnd_status_t randutilGenerateFisherSnedecorDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, @@ -97,10 +103,10 @@ extern "C" curandStatus_t randutilGenerateFisherSnedecorDoubleEx( #include "generator_logseries.inl" -extern "C" curandStatus_t randutilGenerateLogSeriesEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double p) +extern "C" rnd_status_t randutilGenerateLogSeriesEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double p) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; logseries_t func; @@ -114,7 +120,7 @@ extern "C" curandStatus_t randutilGenerateLogSeriesEx(randutilGenerator_t genera #include "generator_chisquare.inl" -extern "C" curandStatus_t randutilGenerateChiSquareEx( +extern "C" rnd_status_t randutilGenerateChiSquareEx( randutilGenerator_t generator, float* outputPtr, size_t n, @@ -134,7 +140,7 @@ extern "C" curandStatus_t randutilGenerateChiSquareEx( } } -extern "C" curandStatus_t randutilGenerateChiSquareDoubleEx( +extern "C" rnd_status_t randutilGenerateChiSquareDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, @@ -160,11 +166,11 @@ extern "C" curandStatus_t randutilGenerateChiSquareDoubleEx( #include "generator_gamma.inl" -extern "C" curandStatus_t randutilGenerateGammaEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float shape, - float scale) // = 1.0f is standard_gamma +extern "C" rnd_status_t randutilGenerateGammaEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float shape, + float scale) // = 1.0f is standard_gamma { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; gamma_t func; @@ -173,11 +179,11 @@ extern "C" curandStatus_t randutilGenerateGammaEx(randutilGenerator_t generator, return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateGammaDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double shape, - double scale) // = 1.0 is standard_gamma +extern "C" rnd_status_t randutilGenerateGammaDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double shape, + double scale) // = 1.0 is standard_gamma { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; gamma_t func; @@ -192,10 +198,10 @@ extern "C" curandStatus_t randutilGenerateGammaDoubleEx(randutilGenerator_t gene #include "generator_standard_t.inl" -extern "C" curandStatus_t randutilGenerateStandardTEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float df) +extern "C" rnd_status_t randutilGenerateStandardTEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float df) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; standard_t_t func; @@ -203,10 +209,10 @@ extern "C" curandStatus_t randutilGenerateStandardTEx(randutilGenerator_t genera return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateStandardTDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double df) +extern "C" rnd_status_t randutilGenerateStandardTDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double df) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; standard_t_t func; @@ -220,7 +226,7 @@ extern "C" curandStatus_t randutilGenerateStandardTDoubleEx(randutilGenerator_t #include "generator_vonmises.inl" -extern "C" curandStatus_t randutilGenerateVonMisesEx( +extern "C" rnd_status_t randutilGenerateVonMisesEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float kappa) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -230,7 +236,7 @@ extern "C" curandStatus_t randutilGenerateVonMisesEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateVonMisesDoubleEx( +extern "C" rnd_status_t randutilGenerateVonMisesDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double kappa) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -246,12 +252,12 @@ extern "C" curandStatus_t randutilGenerateVonMisesDoubleEx( #include "generator_hypergeometric.inl" -extern "C" curandStatus_t randutilGenerateHyperGeometricEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - int64_t ngood, - int64_t nbad, - int64_t nsample) +extern "C" rnd_status_t randutilGenerateHyperGeometricEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + int64_t ngood, + int64_t nbad, + int64_t nsample) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; hypergeometric_t func; @@ -267,10 +273,10 @@ extern "C" curandStatus_t randutilGenerateHyperGeometricEx(randutilGenerator_t g #include "generator_zipf.inl" -extern "C" curandStatus_t randutilGenerateZipfEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double a) +extern "C" rnd_status_t randutilGenerateZipfEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double a) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; zipf_t func; @@ -284,10 +290,10 @@ extern "C" curandStatus_t randutilGenerateZipfEx(randutilGenerator_t generator, #include "generator_geometric.inl" -extern "C" curandStatus_t randutilGenerateGeometricEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double p) +extern "C" rnd_status_t randutilGenerateGeometricEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double p) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; geometric_t func; @@ -301,7 +307,7 @@ extern "C" curandStatus_t randutilGenerateGeometricEx(randutilGenerator_t genera #include "generator_wald.inl" -extern "C" curandStatus_t randutilGenerateWaldEx( +extern "C" rnd_status_t randutilGenerateWaldEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float lambda) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -311,7 +317,7 @@ extern "C" curandStatus_t randutilGenerateWaldEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateWaldDoubleEx( +extern "C" rnd_status_t randutilGenerateWaldDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double lambda) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -327,7 +333,7 @@ extern "C" curandStatus_t randutilGenerateWaldDoubleEx( #include "generator_binomial.inl" -extern "C" curandStatus_t randutilGenerateBinomialEx( +extern "C" rnd_status_t randutilGenerateBinomialEx( randutilGenerator_t generator, uint32_t* outputPtr, size_t n, uint32_t ntrials, double p) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -343,7 +349,7 @@ extern "C" curandStatus_t randutilGenerateBinomialEx( #include "generator_negative_binomial.inl" -extern "C" curandStatus_t randutilGenerateNegativeBinomialEx( +extern "C" rnd_status_t randutilGenerateNegativeBinomialEx( randutilGenerator_t generator, uint32_t* outputPtr, size_t n, uint32_t ntrials, double p) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; diff --git a/src/cunumeric/random/randutil/generator_host_straightforward.cc b/src/cunumeric/random/randutil/generator_host_straightforward.cc index 894ad5571..900c666e9 100644 --- a/src/cunumeric/random/randutil/generator_host_straightforward.cc +++ b/src/cunumeric/random/randutil/generator_host_straightforward.cc @@ -14,16 +14,22 @@ * */ +// MacOS host variant: +// +#if defined(__APPLE__) && defined(__MACH__) +#define USE_STL_RANDOM_ENGINE_ +#endif + #include "generator.h" #pragma region exponential #include "generator_exponential.inl" -extern "C" curandStatus_t randutilGenerateExponentialEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float scale) +extern "C" rnd_status_t randutilGenerateExponentialEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float scale) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; exponential_t func; @@ -31,10 +37,10 @@ extern "C" curandStatus_t randutilGenerateExponentialEx(randutilGenerator_t gene return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateExponentialDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double scale) +extern "C" rnd_status_t randutilGenerateExponentialDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double scale) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; exponential_t func; @@ -48,7 +54,7 @@ extern "C" curandStatus_t randutilGenerateExponentialDoubleEx(randutilGenerator_ #include "generator_gumbel.inl" -extern "C" curandStatus_t randutilGenerateGumbelEx( +extern "C" rnd_status_t randutilGenerateGumbelEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -58,7 +64,7 @@ extern "C" curandStatus_t randutilGenerateGumbelEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateGumbelDoubleEx( +extern "C" rnd_status_t randutilGenerateGumbelDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -74,7 +80,7 @@ extern "C" curandStatus_t randutilGenerateGumbelDoubleEx( #include "generator_laplace.inl" -extern "C" curandStatus_t randutilGenerateLaplaceEx( +extern "C" rnd_status_t randutilGenerateLaplaceEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -84,7 +90,7 @@ extern "C" curandStatus_t randutilGenerateLaplaceEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateLaplaceDoubleEx( +extern "C" rnd_status_t randutilGenerateLaplaceDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -100,7 +106,7 @@ extern "C" curandStatus_t randutilGenerateLaplaceDoubleEx( #include "generator_logistic.inl" -extern "C" curandStatus_t randutilGenerateLogisticEx( +extern "C" rnd_status_t randutilGenerateLogisticEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -110,7 +116,7 @@ extern "C" curandStatus_t randutilGenerateLogisticEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateLogisticDoubleEx( +extern "C" rnd_status_t randutilGenerateLogisticDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -126,7 +132,7 @@ extern "C" curandStatus_t randutilGenerateLogisticDoubleEx( #include "generator_pareto.inl" -extern "C" curandStatus_t randutilGenerateParetoEx( +extern "C" rnd_status_t randutilGenerateParetoEx( randutilGenerator_t generator, float* outputPtr, size_t n, float xm, float alpha) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -136,7 +142,7 @@ extern "C" curandStatus_t randutilGenerateParetoEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateParetoDoubleEx( +extern "C" rnd_status_t randutilGenerateParetoDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double xm, double alpha) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -152,10 +158,10 @@ extern "C" curandStatus_t randutilGenerateParetoDoubleEx( #include "generator_power.inl" -extern "C" curandStatus_t randutilGeneratePowerEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float alpha) +extern "C" rnd_status_t randutilGeneratePowerEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float alpha) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; power_t func; @@ -163,10 +169,10 @@ extern "C" curandStatus_t randutilGeneratePowerEx(randutilGenerator_t generator, return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGeneratePowerDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double alpha) +extern "C" rnd_status_t randutilGeneratePowerDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double alpha) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; power_t func; @@ -180,10 +186,10 @@ extern "C" curandStatus_t randutilGeneratePowerDoubleEx(randutilGenerator_t gene #include "generator_rayleigh.inl" -extern "C" curandStatus_t randutilGenerateRayleighEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float sigma) +extern "C" rnd_status_t randutilGenerateRayleighEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float sigma) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; rayleigh_t func; @@ -191,10 +197,10 @@ extern "C" curandStatus_t randutilGenerateRayleighEx(randutilGenerator_t generat return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateRayleighDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double sigma) +extern "C" rnd_status_t randutilGenerateRayleighDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double sigma) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; rayleigh_t func; @@ -208,7 +214,7 @@ extern "C" curandStatus_t randutilGenerateRayleighDoubleEx(randutilGenerator_t g #include "generator_cauchy.inl" -extern "C" curandStatus_t randutilGenerateCauchyEx( +extern "C" rnd_status_t randutilGenerateCauchyEx( randutilGenerator_t generator, float* outputPtr, size_t n, float x0, float gamma) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -218,7 +224,7 @@ extern "C" curandStatus_t randutilGenerateCauchyEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateCauchyDoubleEx( +extern "C" rnd_status_t randutilGenerateCauchyDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double x0, double gamma) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -234,7 +240,7 @@ extern "C" curandStatus_t randutilGenerateCauchyDoubleEx( #include "generator_triangular.inl" -extern "C" curandStatus_t randutilGenerateTriangularEx( +extern "C" rnd_status_t randutilGenerateTriangularEx( randutilGenerator_t generator, float* outputPtr, size_t n, float a, float b, float c) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -245,7 +251,7 @@ extern "C" curandStatus_t randutilGenerateTriangularEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateTriangularDoubleEx( +extern "C" rnd_status_t randutilGenerateTriangularDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double a, double b, double c) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -262,7 +268,7 @@ extern "C" curandStatus_t randutilGenerateTriangularDoubleEx( #include "generator_weibull.inl" -extern "C" curandStatus_t randutilGenerateWeibullEx( +extern "C" rnd_status_t randutilGenerateWeibullEx( randutilGenerator_t generator, float* outputPtr, size_t n, float lam, float k) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; @@ -272,7 +278,7 @@ extern "C" curandStatus_t randutilGenerateWeibullEx( return randutilimpl::dispatch(gen, func, n, outputPtr); } -extern "C" curandStatus_t randutilGenerateWeibullDoubleEx( +extern "C" rnd_status_t randutilGenerateWeibullDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double lam, double k) { randutilimpl::basegenerator* gen = (randutilimpl::basegenerator*)generator; diff --git a/src/cunumeric/random/randutil/generator_integers.inl b/src/cunumeric/random/randutil/generator_integers.inl index cafb5e591..7227ad63c 100644 --- a/src/cunumeric/random/randutil/generator_integers.inl +++ b/src/cunumeric/random/randutil/generator_integers.inl @@ -15,45 +15,42 @@ */ #include "generator.h" +#include -template -struct integers; - -template <> -struct integers { - int16_t from; - int16_t to; +#include "randomizer.h" - template - RANDUTIL_QUALIFIERS int32_t operator()(gen_t& gen) - { - return (int16_t)(curand(&gen) % (uint16_t)(to - from)) + from; - } -}; +template +struct integers; -template <> -struct integers { - int32_t from; - int32_t to; +template +struct integers< + field_t, + std::enable_if_t || std::is_same_v>> { + using ufield_t = std::conditional_t, uint16_t, uint32_t>; + field_t from; + field_t to; template - RANDUTIL_QUALIFIERS int32_t operator()(gen_t& gen) + RANDUTIL_QUALIFIERS field_t operator()(gen_t& gen) { - return (int32_t)(curand(&gen) % (uint32_t)(to - from)) + from; + auto y = randutilimpl::engine_rand(gen); + return (field_t)(y % (ufield_t)(to - from)) + from; } }; -template <> -struct integers { - int64_t from; - int64_t to; +template +struct integers>> { + using ufield_t = uint64_t; + field_t from; + field_t to; template - RANDUTIL_QUALIFIERS int64_t operator()(gen_t& gen) + RANDUTIL_QUALIFIERS field_t operator()(gen_t& gen) { // take two draws to get a 64 bits value - unsigned low = curand(&gen); - unsigned high = curand(&gen); - return (int64_t)((((uint64_t)high << 32) | (uint64_t)low) % (to - from)) + from; + unsigned low = randutilimpl::engine_rand(gen); + unsigned high = randutilimpl::engine_rand(gen); + + return (field_t)((((ufield_t)high << 32) | (ufield_t)low) % (to - from)) + from; } }; diff --git a/src/cunumeric/random/randutil/generator_laplace.inl b/src/cunumeric/random/randutil/generator_laplace.inl index 0d4c7b1ae..9177c948f 100644 --- a/src/cunumeric/random/randutil/generator_laplace.inl +++ b/src/cunumeric/random/randutil/generator_laplace.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct laplace_t; @@ -26,7 +28,7 @@ struct laplace_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be zero + float y = randutilimpl::engine_uniform(gen); // y cannot be zero if (y == 0.5f) return mu; if (y < 0.5f) return mu + beta * ::logf(2.0f * y); @@ -42,7 +44,7 @@ struct laplace_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be zero + double y = randutilimpl::engine_uniform(gen); // y cannot be zero if (y == 0.5) return mu; if (y < 0.5) return mu + beta * ::log(2.0 * y); diff --git a/src/cunumeric/random/randutil/generator_logistic.inl b/src/cunumeric/random/randutil/generator_logistic.inl index 7405c18a3..d37d515ca 100644 --- a/src/cunumeric/random/randutil/generator_logistic.inl +++ b/src/cunumeric/random/randutil/generator_logistic.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct logistic_t; @@ -26,7 +28,7 @@ struct logistic_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 float t = 1.0f / y - 1.0f; if (t == 0) t = 1.0f; return mu - beta * ::logf(t); @@ -40,8 +42,8 @@ struct logistic_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - float y = curand_uniform_double(&gen); // y cannot be 0 - float t = 1.0 / y - 1.0; + auto y = randutilimpl::engine_uniform(gen); // y cannot be 0 + auto t = 1.0 / y - 1.0; if (t == 0) t = 1.0; return mu - beta * ::log(t); } diff --git a/src/cunumeric/random/randutil/generator_lognormal.inl b/src/cunumeric/random/randutil/generator_lognormal.inl index 972d2ce1c..ba1138ac6 100644 --- a/src/cunumeric/random/randutil/generator_lognormal.inl +++ b/src/cunumeric/random/randutil/generator_lognormal.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct lognormal_t; @@ -27,7 +29,7 @@ struct lognormal_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - return curand_log_normal(&gen, mean, stddev); + return randutilimpl::engine_log_normal(gen, mean, stddev); } }; @@ -39,6 +41,6 @@ struct lognormal_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - return curand_log_normal_double(&gen, mean, stddev); + return randutilimpl::engine_log_normal(gen, mean, stddev); } }; diff --git a/src/cunumeric/random/randutil/generator_negative_binomial.inl b/src/cunumeric/random/randutil/generator_negative_binomial.inl index 241f21dea..5521af420 100644 --- a/src/cunumeric/random/randutil/generator_negative_binomial.inl +++ b/src/cunumeric/random/randutil/generator_negative_binomial.inl @@ -17,6 +17,8 @@ #include "generator.h" #include "random_distributions.h" +#include "randomizer.h" + template struct negative_binomial_t; @@ -29,6 +31,6 @@ struct negative_binomial_t { RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { double lambda = rk_standard_gamma(&gen, (double)n) * ((1 - p) / p); - return curand_poisson(&gen, lambda); + return static_cast(randutilimpl::engine_poisson(gen, lambda)); } }; diff --git a/src/cunumeric/random/randutil/generator_normal.inl b/src/cunumeric/random/randutil/generator_normal.inl index bbe2d8ecb..a834b9d39 100644 --- a/src/cunumeric/random/randutil/generator_normal.inl +++ b/src/cunumeric/random/randutil/generator_normal.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct normal_t; @@ -27,7 +29,7 @@ struct normal_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - return stddev * curand_normal(&gen) + mean; + return stddev * randutilimpl::engine_normal(gen) + mean; } }; @@ -39,6 +41,6 @@ struct normal_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - return stddev * curand_normal_double(&gen) + mean; + return stddev * randutilimpl::engine_normal(gen) + mean; } }; diff --git a/src/cunumeric/random/randutil/generator_pareto.inl b/src/cunumeric/random/randutil/generator_pareto.inl index 2a29f2479..073e957f6 100644 --- a/src/cunumeric/random/randutil/generator_pareto.inl +++ b/src/cunumeric/random/randutil/generator_pareto.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct pareto_t; @@ -26,8 +28,8 @@ struct pareto_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 - return xm * ::expf(-::logf(y) * invalpha) - 1.0f; // here, use -1.0f to align with numpy + auto y = randutilimpl::engine_uniform(gen); // y cannot be 0 + return xm * ::expf(-::logf(y) * invalpha) - 1.0f; // here, use -1.0f to align with numpy } }; @@ -38,7 +40,7 @@ struct pareto_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 - return xm * ::exp(-::log(y) * invalpha) - 1.0; // here, use -1.0 to align with numpy + auto y = randutilimpl::engine_uniform(gen); // y cannot be 0 + return xm * ::exp(-::log(y) * invalpha) - 1.0; // here, use -1.0 to align with numpy } }; diff --git a/src/cunumeric/random/randutil/generator_poisson.inl b/src/cunumeric/random/randutil/generator_poisson.inl index c441ebac4..f8f238f20 100644 --- a/src/cunumeric/random/randutil/generator_poisson.inl +++ b/src/cunumeric/random/randutil/generator_poisson.inl @@ -16,12 +16,14 @@ #include "generator.h" +#include "randomizer.h" + struct poisson { double lambda = 1.0; template RANDUTIL_QUALIFIERS unsigned operator()(gen_t& gen) { - return curand_poisson(&gen, lambda); + return randutilimpl::engine_poisson(gen, lambda); } }; diff --git a/src/cunumeric/random/randutil/generator_power.inl b/src/cunumeric/random/randutil/generator_power.inl index e17eafc55..7c3c278a2 100644 --- a/src/cunumeric/random/randutil/generator_power.inl +++ b/src/cunumeric/random/randutil/generator_power.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct power_t; @@ -26,7 +28,7 @@ struct power_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 return ::expf(::logf(y) * invalpha); } }; @@ -38,7 +40,7 @@ struct power_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 -- use y as 1-cdf(x) + double y = randutilimpl::engine_uniform(gen); // y cannot be 0 -- use y as 1-cdf(x) return ::exp(::log(y) * invalpha); } }; diff --git a/src/cunumeric/random/randutil/generator_raw.inl b/src/cunumeric/random/randutil/generator_raw.inl index 5414e270e..5db4a203a 100644 --- a/src/cunumeric/random/randutil/generator_raw.inl +++ b/src/cunumeric/random/randutil/generator_raw.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct raw; @@ -24,6 +26,6 @@ struct raw { template RANDUTIL_QUALIFIERS uint32_t operator()(gen_t& gen) { - return (uint32_t)curand(&gen); + return (uint32_t)randutilimpl::engine_rand(gen); } }; diff --git a/src/cunumeric/random/randutil/generator_rayleigh.inl b/src/cunumeric/random/randutil/generator_rayleigh.inl index 3d34043b4..d1a7b2cef 100644 --- a/src/cunumeric/random/randutil/generator_rayleigh.inl +++ b/src/cunumeric/random/randutil/generator_rayleigh.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct rayleigh_t; @@ -26,7 +28,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; y cannot be 0 return sigma * ::sqrtf(-2.0f * ::logf(y)); } }; @@ -38,7 +40,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; y cannot be 0 return sigma * ::sqrt(-2.0 * ::log(y)); } }; diff --git a/src/cunumeric/random/randutil/generator_triangular.inl b/src/cunumeric/random/randutil/generator_triangular.inl index 3bfe77e28..318d569eb 100644 --- a/src/cunumeric/random/randutil/generator_triangular.inl +++ b/src/cunumeric/random/randutil/generator_triangular.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct triangular_t; @@ -26,7 +28,7 @@ struct triangular_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 if (y <= ((c - a) / (b - a))) { float delta = (y * (b - a) * (c - a)); if (delta < 0.0f) delta = 0.0f; @@ -46,7 +48,7 @@ struct triangular_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 + double y = randutilimpl::engine_uniform(gen); // y cannot be 0 if (y <= ((c - a) / (b - a))) { double delta = (y * (b - a) * (c - a)); if (delta < 0.0) delta = 0.0; diff --git a/src/cunumeric/random/randutil/generator_uniform.inl b/src/cunumeric/random/randutil/generator_uniform.inl index d6eecd52f..456f82034 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct uniform_t; @@ -26,7 +28,8 @@ struct uniform_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - return offset + mult * curand_uniform(&gen); + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; + return offset + mult * y; } }; @@ -37,6 +40,7 @@ struct uniform_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - return offset + mult * curand_uniform_double(&gen); + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; + return offset + mult * y; } }; diff --git a/src/cunumeric/random/randutil/generator_wald.inl b/src/cunumeric/random/randutil/generator_wald.inl index dec310834..20706e6ce 100644 --- a/src/cunumeric/random/randutil/generator_wald.inl +++ b/src/cunumeric/random/randutil/generator_wald.inl @@ -17,6 +17,8 @@ #include "generator.h" #include "random_distributions.h" +#include "randomizer.h" + template struct wald_t; @@ -27,11 +29,11 @@ struct wald_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float v = curand_normal(&gen); + float v = randutilimpl::engine_normal(gen); float y = v * v; float x = mu + (mu * mu * y) / (2.0f * lambda) - (mu / (2.0f * lambda)) * ::sqrtf(mu * y * (4.0f * lambda + mu * y)); - float z = curand_uniform(&gen); + float z = randutilimpl::engine_uniform(gen); if (z <= (mu) / (mu + x)) return x; else @@ -46,11 +48,11 @@ struct wald_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double v = curand_normal(&gen); + double v = randutilimpl::engine_normal(gen); double y = v * v; double x = mu + (mu * mu * y) / (2.0 * lambda) - (mu / (2.0 * lambda)) * ::sqrtf(mu * y * (4.0 * lambda + mu * y)); - double z = curand_uniform(&gen); + double z = randutilimpl::engine_uniform(gen); if (z <= (mu) / (mu + x)) return x; else diff --git a/src/cunumeric/random/randutil/generator_weibull.inl b/src/cunumeric/random/randutil/generator_weibull.inl index 5565c7411..ad1ec836d 100644 --- a/src/cunumeric/random/randutil/generator_weibull.inl +++ b/src/cunumeric/random/randutil/generator_weibull.inl @@ -16,6 +16,8 @@ #include "generator.h" +#include "randomizer.h" + template struct weibull_t; @@ -26,9 +28,9 @@ struct weibull_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 // log(y) can be zero ! - float lny = ::logf(y); + auto lny = ::logf(y); if (lny == 0.0f) return 0.0f; return lambda * ::expf(::logf(-lny) * invk); } @@ -41,9 +43,9 @@ struct weibull_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 + double y = randutilimpl::engine_uniform(gen); // y cannot be 0 // log(y) can be zero ! - float lny = ::log(y); + auto lny = ::log(y); if (lny == 0.0f) return 0.0f; return lambda * ::exp(::log(-lny) * invk); } diff --git a/src/cunumeric/random/randutil/random_distributions.h b/src/cunumeric/random/randutil/random_distributions.h index db4e609e4..5307a5858 100644 --- a/src/cunumeric/random/randutil/random_distributions.h +++ b/src/cunumeric/random/randutil/random_distributions.h @@ -49,10 +49,12 @@ #pragma once +#include "randomizer.h" + template RANDUTIL_QUALIFIERS double rk_double(rk_state* gen) { - return curand_uniform_double(gen); + return randutilimpl::engine_uniform(*gen); // returns (0, 1]; } RANDUTIL_QUALIFIERS double loggam(double x) @@ -120,7 +122,7 @@ RANDUTIL_QUALIFIERS double rk_gauss(rk_state* state) return f*x2; } #endif - return curand_normal(state); + return randutilimpl::engine_normal(*state); } template @@ -260,7 +262,7 @@ RANDUTIL_QUALIFIERS long rk_poisson(rk_state* state, double lam) return rk_poisson_mult(state, lam); } #endif - return curand_poisson(state, lam); + return randutilimpl::engine_poisson(*state, lam); } template @@ -646,4 +648,4 @@ RANDUTIL_QUALIFIERS unsigned rk_binomial(rk_state* state, unsigned n, double p) } } -#pragma endregion \ No newline at end of file +#pragma endregion diff --git a/src/cunumeric/random/randutil/randomizer.h b/src/cunumeric/random/randutil/randomizer.h new file mode 100644 index 000000000..c9bcda026 --- /dev/null +++ b/src/cunumeric/random/randutil/randomizer.h @@ -0,0 +1,132 @@ +/* Copyright 2023 NVIDIA Corporation + * + * Licensed 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. + * + */ + +#pragma once + +#include +#include +#include + +#include + +namespace randutilimpl { + +// trampoline functions for branching-off CURAND vs. STL +// implementations (STL used on platforms that don't support CUDA) +// +template +RANDUTIL_QUALIFIERS decltype(auto) engine_uniform(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0, 1); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + return 1 - y; +#else + if constexpr (std::is_same_v) + return curand_uniform(&gen); // returns (0, 1]; + else + return curand_uniform_double(&gen); // returns (0, 1]; +#endif +} + +template +RANDUTIL_QUALIFIERS decltype(auto) engine_poisson(gen_t& gen, double lambda) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::poisson_distribution dis(lambda); + return dis(gen); +#else + return curand_poisson(&gen, lambda); +#endif +} + +template +RANDUTIL_QUALIFIERS decltype(auto) engine_normal(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::normal_distribution dis(0, 1); + return dis(gen); +#else + if constexpr (std::is_same_v) + return curand_normal(&gen); + else + return curand_normal_double(&gen); +#endif +} + +template +RANDUTIL_QUALIFIERS decltype(auto) engine_log_normal(gen_t& gen, element_t mean, element_t stddev) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::lognormal_distribution dis{mean, stddev}; + return dis(gen); +#else + if constexpr (std::is_same_v) + return curand_log_normal(&gen, mean, stddev); + else + return curand_log_normal_double(&gen, mean, stddev); +#endif +} + +template +RANDUTIL_QUALIFIERS decltype(auto) engine_rand(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + return std::rand(); +#else + return curand(&gen); +#endif +} + +#ifdef EXPERIMENTAL_STL_BRANCH_OFF_ +enum class random_client_t : int { CURAND = 0, STL }; + +template +struct randomizer_t; + +// Curand randomizer maintains a state: +// +template +struct randomizer_t { + randomizer_t(state_t state, std::function f) : state_(state), f_(f) {} + + ret_t run(void) { return std::invoke(f_, state_); } + + private: + state_t state_; + std::function f_; +}; + +// STL randomizer uses Distribution-specific info (DSI); e.g., (low, high) for uniform: +// +template +struct randomizer_t { + randomizer_t(dsi_data_t const& data, std::function f) + : data_(data), f_(f) + { + } + + ret_t run(void) const { return std::invoke(f_, data_); } + + private: + dsi_data_t data_; + std::function f_; +}; +#endif + +} // namespace randutilimpl diff --git a/src/cunumeric/random/randutil/randutil.h b/src/cunumeric/random/randutil/randutil.h index f0871e14f..f3d84c1f6 100644 --- a/src/cunumeric/random/randutil/randutil.h +++ b/src/cunumeric/random/randutil/randutil.h @@ -16,7 +16,9 @@ #pragma once #include -#include +// #include + +#include "cunumeric/random/rnd_aliases.h" typedef void* randutilGenerator_t; @@ -24,149 +26,149 @@ typedef void* randutilGenerator_t; // CUDA-ONLY API #ifdef LEGATE_USE_CUDA -extern "C" curandStatus_t randutilCreateGenerator(randutilGenerator_t* generator, - curandRngType_t rng_type, - uint64_t seed, - uint64_t generatorID, - uint32_t flags, - cudaStream_t stream); +extern "C" rnd_status_t randutilCreateGenerator(randutilGenerator_t* generator, + randRngType_t rng_type, + uint64_t seed, + uint64_t generatorID, + uint32_t flags, + stream_t stream); #endif -extern "C" curandStatus_t randutilCreateGeneratorHost(randutilGenerator_t* generator, - curandRngType_t rng_type, - uint64_t seed, - uint64_t generatorID, - uint32_t flags); -extern "C" curandStatus_t randutilDestroyGenerator(randutilGenerator_t generator); +extern "C" rnd_status_t randutilCreateGeneratorHost(randutilGenerator_t* generator, + randRngType_t rng_type, + uint64_t seed, + uint64_t generatorID, + uint32_t flags); +extern "C" rnd_status_t randutilDestroyGenerator(randutilGenerator_t generator); /* curand distributions */ -extern "C" curandStatus_t randutilGenerateIntegers16(randutilGenerator_t generator, - int16_t* outputPtr, - size_t num, - int16_t low /* inclusive */, - int16_t high /* exclusive */); - -extern "C" curandStatus_t randutilGenerateIntegers32(randutilGenerator_t generator, - int32_t* outputPtr, - size_t num, - int32_t low /* inclusive */, - int32_t high /* exclusive */); -extern "C" curandStatus_t randutilGenerateIntegers64(randutilGenerator_t generator, - int64_t* outputPtr, - size_t num, - int64_t low /* inclusive */, - int64_t high /* exclusive */); -extern "C" curandStatus_t randutilGenerateRawUInt32(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t num); - -extern "C" curandStatus_t randutilGenerateUniformEx(randutilGenerator_t generator, - float* outputPtr, - size_t num, - float low = 0.0f, /* inclusive */ - float high = 1.0f /* exclusive */); +extern "C" rnd_status_t randutilGenerateIntegers16(randutilGenerator_t generator, + int16_t* outputPtr, + size_t num, + int16_t low /* inclusive */, + int16_t high /* exclusive */); + +extern "C" rnd_status_t randutilGenerateIntegers32(randutilGenerator_t generator, + int32_t* outputPtr, + size_t num, + int32_t low /* inclusive */, + int32_t high /* exclusive */); +extern "C" rnd_status_t randutilGenerateIntegers64(randutilGenerator_t generator, + int64_t* outputPtr, + size_t num, + int64_t low /* inclusive */, + int64_t high /* exclusive */); +extern "C" rnd_status_t randutilGenerateRawUInt32(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t num); + +extern "C" rnd_status_t randutilGenerateUniformEx(randutilGenerator_t generator, + float* outputPtr, + size_t num, + float low = 0.0f, /* inclusive */ + float high = 1.0f /* exclusive */); -extern "C" curandStatus_t randutilGenerateUniformDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t num, - double low = 0.0, /* inclusive */ - double high = 1.0 /* exclusive */); +extern "C" rnd_status_t randutilGenerateUniformDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t num, + double low = 0.0, /* inclusive */ + double high = 1.0 /* exclusive */); -extern "C" curandStatus_t randutilGenerateLogNormalEx( +extern "C" rnd_status_t randutilGenerateLogNormalEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mean, float stddev); -extern "C" curandStatus_t randutilGenerateLogNormalDoubleEx( +extern "C" rnd_status_t randutilGenerateLogNormalDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mean, double stddev); -extern "C" curandStatus_t randutilGenerateNormalEx( +extern "C" rnd_status_t randutilGenerateNormalEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mean, float stddev); -extern "C" curandStatus_t randutilGenerateNormalDoubleEx( +extern "C" rnd_status_t randutilGenerateNormalDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mean, double stddev); -extern "C" curandStatus_t randutilGeneratePoissonEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double lambda); +extern "C" rnd_status_t randutilGeneratePoissonEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double lambda); /* Straightforward Distributions */ -extern "C" curandStatus_t randutilGenerateExponentialEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float scale); -extern "C" curandStatus_t randutilGenerateExponentialDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double scale); +extern "C" rnd_status_t randutilGenerateExponentialEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float scale); +extern "C" rnd_status_t randutilGenerateExponentialDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double scale); -extern "C" curandStatus_t randutilGenerateGumbelEx( +extern "C" rnd_status_t randutilGenerateGumbelEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta); -extern "C" curandStatus_t randutilGenerateGumbelDoubleEx( +extern "C" rnd_status_t randutilGenerateGumbelDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta); -extern "C" curandStatus_t randutilGenerateLaplaceEx( +extern "C" rnd_status_t randutilGenerateLaplaceEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta); -extern "C" curandStatus_t randutilGenerateLaplaceDoubleEx( +extern "C" rnd_status_t randutilGenerateLaplaceDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta); -extern "C" curandStatus_t randutilGenerateLogisticEx( +extern "C" rnd_status_t randutilGenerateLogisticEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float beta); -extern "C" curandStatus_t randutilGenerateLogisticDoubleEx( +extern "C" rnd_status_t randutilGenerateLogisticDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double beta); -extern "C" curandStatus_t randutilGenerateParetoEx( +extern "C" rnd_status_t randutilGenerateParetoEx( randutilGenerator_t generator, float* outputPtr, size_t n, float xm, float alpha); -extern "C" curandStatus_t randutilGenerateParetoDoubleEx( +extern "C" rnd_status_t randutilGenerateParetoDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double xm, double alpha); -extern "C" curandStatus_t randutilGeneratePowerEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float alpha); -extern "C" curandStatus_t randutilGeneratePowerDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double alpha); - -extern "C" curandStatus_t randutilGenerateRayleighEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float sigma); -extern "C" curandStatus_t randutilGenerateRayleighDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double sigma); - -extern "C" curandStatus_t randutilGenerateCauchyEx( +extern "C" rnd_status_t randutilGeneratePowerEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float alpha); +extern "C" rnd_status_t randutilGeneratePowerDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double alpha); + +extern "C" rnd_status_t randutilGenerateRayleighEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float sigma); +extern "C" rnd_status_t randutilGenerateRayleighDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double sigma); + +extern "C" rnd_status_t randutilGenerateCauchyEx( randutilGenerator_t generator, float* outputPtr, size_t n, float x0, float gamma); -extern "C" curandStatus_t randutilGenerateCauchyDoubleEx( +extern "C" rnd_status_t randutilGenerateCauchyDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double x0, double gamma); -extern "C" curandStatus_t randutilGenerateTriangularEx( +extern "C" rnd_status_t randutilGenerateTriangularEx( randutilGenerator_t generator, float* outputPtr, size_t n, float a, float b, float c); -extern "C" curandStatus_t randutilGenerateTriangularDoubleEx( +extern "C" rnd_status_t randutilGenerateTriangularDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double a, double b, double c); -extern "C" curandStatus_t randutilGenerateWeibullEx( +extern "C" rnd_status_t randutilGenerateWeibullEx( randutilGenerator_t generator, float* outputPtr, size_t n, float lam, float k); -extern "C" curandStatus_t randutilGenerateWeibullDoubleEx( +extern "C" rnd_status_t randutilGenerateWeibullDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double lam, double k); /* more advanced distributions */ -extern "C" curandStatus_t randutilGenerateBetaEx( +extern "C" rnd_status_t randutilGenerateBetaEx( randutilGenerator_t generator, float* outputPtr, size_t n, float a, float b); -extern "C" curandStatus_t randutilGenerateBetaDoubleEx( +extern "C" rnd_status_t randutilGenerateBetaDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double a, double b); -extern "C" curandStatus_t randutilGenerateFisherSnedecorEx( +extern "C" rnd_status_t randutilGenerateFisherSnedecorEx( randutilGenerator_t generator, float* outputPtr, size_t n, float dfnum, float dfden, float nonc = 0.0f); // 0.0f is F distribution -extern "C" curandStatus_t randutilGenerateFisherSnedecorDoubleEx( +extern "C" rnd_status_t randutilGenerateFisherSnedecorDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, @@ -174,57 +176,57 @@ extern "C" curandStatus_t randutilGenerateFisherSnedecorDoubleEx( double dfden, double nonc = 0.0); // 0.0 is F distribution -extern "C" curandStatus_t randutilGenerateLogSeriesEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double p); +extern "C" rnd_status_t randutilGenerateLogSeriesEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double p); -extern "C" curandStatus_t randutilGenerateChiSquareEx( +extern "C" rnd_status_t randutilGenerateChiSquareEx( randutilGenerator_t generator, float* outputPtr, size_t n, float df, float nonc = 0.0); -extern "C" curandStatus_t randutilGenerateChiSquareDoubleEx( +extern "C" rnd_status_t randutilGenerateChiSquareDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double df, double nonc = 0.0); -extern "C" curandStatus_t randutilGenerateGammaEx( +extern "C" rnd_status_t randutilGenerateGammaEx( randutilGenerator_t generator, float* outputPtr, size_t n, float shape, float scale = 1.0f); // scale = 1.0 is standard_gamma -extern "C" curandStatus_t randutilGenerateGammaDoubleEx( +extern "C" rnd_status_t randutilGenerateGammaDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double shape, double scale = 1.0); -extern "C" curandStatus_t randutilGenerateStandardTDoubleEx(randutilGenerator_t generator, - double* outputPtr, - size_t n, - double df); -extern "C" curandStatus_t randutilGenerateStandardTEx(randutilGenerator_t generator, - float* outputPtr, - size_t n, - float df); -extern "C" curandStatus_t randutilGenerateHyperGeometricEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - int64_t ngood, - int64_t nbad, - int64_t nsample); -extern "C" curandStatus_t randutilGenerateVonMisesDoubleEx( +extern "C" rnd_status_t randutilGenerateStandardTDoubleEx(randutilGenerator_t generator, + double* outputPtr, + size_t n, + double df); +extern "C" rnd_status_t randutilGenerateStandardTEx(randutilGenerator_t generator, + float* outputPtr, + size_t n, + float df); +extern "C" rnd_status_t randutilGenerateHyperGeometricEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + int64_t ngood, + int64_t nbad, + int64_t nsample); +extern "C" rnd_status_t randutilGenerateVonMisesDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double kappa); -extern "C" curandStatus_t randutilGenerateVonMisesEx( +extern "C" rnd_status_t randutilGenerateVonMisesEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float kappa); -extern "C" curandStatus_t randutilGenerateZipfEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double a); -extern "C" curandStatus_t randutilGenerateGeometricEx(randutilGenerator_t generator, - uint32_t* outputPtr, - size_t n, - double p); -extern "C" curandStatus_t randutilGenerateWaldDoubleEx( +extern "C" rnd_status_t randutilGenerateZipfEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double a); +extern "C" rnd_status_t randutilGenerateGeometricEx(randutilGenerator_t generator, + uint32_t* outputPtr, + size_t n, + double p); +extern "C" rnd_status_t randutilGenerateWaldDoubleEx( randutilGenerator_t generator, double* outputPtr, size_t n, double mu, double lambda); -extern "C" curandStatus_t randutilGenerateWaldEx( +extern "C" rnd_status_t randutilGenerateWaldEx( randutilGenerator_t generator, float* outputPtr, size_t n, float mu, float lambda); -extern "C" curandStatus_t randutilGenerateBinomialEx( +extern "C" rnd_status_t randutilGenerateBinomialEx( randutilGenerator_t generator, uint32_t* outputPtr, size_t n, uint32_t ntrials, double p); -extern "C" curandStatus_t randutilGenerateNegativeBinomialEx( +extern "C" rnd_status_t randutilGenerateNegativeBinomialEx( randutilGenerator_t generator, uint32_t* outputPtr, size_t n, uint32_t ntrials, double p); diff --git a/src/cunumeric/random/randutil/randutil_curand.h b/src/cunumeric/random/randutil/randutil_curand.h index 42e5c19de..bfca2eff2 100644 --- a/src/cunumeric/random/randutil/randutil_curand.h +++ b/src/cunumeric/random/randutil/randutil_curand.h @@ -30,6 +30,10 @@ // host generators are not compiled with nvcc #define QUALIFIERS static #define RANDUTIL_QUALIFIERS +#ifdef USE_STL_RANDOM_ENGINE_ +#include +#else #include +#endif -#endif \ No newline at end of file +#endif diff --git a/src/cunumeric/random/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h new file mode 100644 index 000000000..f98d41a53 --- /dev/null +++ b/src/cunumeric/random/rnd_aliases.h @@ -0,0 +1,75 @@ +/* Copyright 2023 NVIDIA Corporation + * + * Licensed 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. + * + */ + +#pragma once + +#ifdef USE_STL_RANDOM_ENGINE_ + +// #pragma message("************ STL path *************") + +#include +#include + +using rnd_status_t = int; +enum class randRngType : int { RND_RNG_TEST = 0, STL_MT_19937 = 1 }; +using randRngType_t = randRngType; +constexpr int RND_STATUS_SUCCESS = 0; + +// same as for curand: +// +constexpr rnd_status_t RND_STATUS_INTERNAL_ERROR = 999; +constexpr rnd_status_t RND_STATUS_TYPE_ERROR = 103; + +// namespace randutilimpl { +constexpr int RND_RNG_PSEUDO_XORWOW = static_cast(randRngType::STL_MT_19937); +constexpr int RND_RNG_PSEUDO_PHILOX4_32_10 = static_cast(randRngType::STL_MT_19937) + 1; +constexpr int RND_RNG_PSEUDO_MRG32K3A = static_cast(randRngType::STL_MT_19937) + 2; + +// cannot be same, b/c they are used for +// specializing class generatorid: +// +using gen_XORWOW_t = std::mt19937; +using gen_Philox4_32_10_t = std::mt19937; +using gen_MRG32k3a_t = std::mt19937; +//} // namespace randutilimpl + +using stream_t = void*; +#else +#include +#include + +// #pragma message("************ CURAND path ************") + +using rnd_status_t = curandStatus_t; +using randRngType = curandRngType; +using randRngType_t = curandRngType_t; +constexpr rnd_status_t RND_STATUS_SUCCESS = CURAND_STATUS_SUCCESS; +constexpr rnd_status_t RND_STATUS_INTERNAL_ERROR = CURAND_STATUS_INTERNAL_ERROR; +constexpr rnd_status_t RND_STATUS_TYPE_ERROR = CURAND_STATUS_TYPE_ERROR; + +constexpr randRngType_t RND_RNG_TEST = CURAND_RNG_TEST; + +constexpr int RND_RNG_PSEUDO_XORWOW = CURAND_RNG_PSEUDO_XORWOW; +constexpr int RND_RNG_PSEUDO_PHILOX4_32_10 = CURAND_RNG_PSEUDO_PHILOX4_32_10; +constexpr int RND_RNG_PSEUDO_MRG32K3A = CURAND_RNG_PSEUDO_MRG32K3A; + +using gen_XORWOW_t = curandStateXORWOW_t; +using gen_Philox4_32_10_t = curandStatePhilox4_32_10_t; +using gen_MRG32k3a_t = curandStateMRG32k3a_t; + +using stream_t = cudaStream_t; + +#endif diff --git a/src/cunumeric/random/rnd_types.h b/src/cunumeric/random/rnd_types.h new file mode 100644 index 000000000..50ee67355 --- /dev/null +++ b/src/cunumeric/random/rnd_types.h @@ -0,0 +1,63 @@ +/* Copyright 2023 NVIDIA Corporation + * + * Licensed 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. + * + */ + +#pragma once + +#include "cunumeric/random/rnd_aliases.h" + +#ifdef USE_STL_RANDOM_ENGINE_ + +#define CHECK_RND_ENGINE(expr) \ + do { \ + rnd_status_t __result__ = (expr); \ + randutil_check_status(__result__, __FILE__, __LINE__); \ + } while (false) + +// #define randutil_check_curand randutil_check_status + +namespace cunumeric { +legate::Logger& randutil_log(); + +void randutil_check_status(rnd_status_t error, const char* file, int line); + +static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) +{ + // for now, all generator types rerouted to STL + // would use the MT19937 generator; perhaps, + // this might become more flexible in the future; + // + switch (kind) { + case cunumeric::BitGeneratorType::DEFAULT: return randRngType::STL_MT_19937; + case cunumeric::BitGeneratorType::XORWOW: return randRngType::STL_MT_19937; + case cunumeric::BitGeneratorType::MRG32K3A: return randRngType::STL_MT_19937; + case cunumeric::BitGeneratorType::MTGP32: return randRngType::STL_MT_19937; + case cunumeric::BitGeneratorType::MT19937: return randRngType::STL_MT_19937; + case cunumeric::BitGeneratorType::PHILOX4_32_10: return randRngType::STL_MT_19937; + default: LEGATE_ABORT; + } + return randRngType::RND_RNG_TEST; +} + +} // namespace cunumeric + +#else +#include "cunumeric/random/curand_help.h" + +#define CHECK_RND_ENGINE(expr) CHECK_CURAND((expr)) +#define randutil_check_status randutil_check_curand +#define get_rndRngType get_curandRngType + +#endif