From dd057489664d03476a7875282d29d3de80d3a92a Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Tue, 14 Nov 2023 17:58:10 -0600 Subject: [PATCH 01/17] Branch-off distributions on STL. Group 1. --- src/cunumeric/random/randutil/generator.h | 22 +++++++- .../random/randutil/generator_gumbel.inl | 18 ++++++ .../random/randutil/generator_integers.inl | 56 ++++++++++--------- .../randutil/generator_negative_binomial.inl | 5 ++ .../random/randutil/generator_pareto.inl | 20 ++++++- .../random/randutil/generator_poisson.inl | 5 ++ .../random/randutil/generator_rayleigh.inl | 20 ++++++- .../random/randutil/generator_uniform.inl | 34 ++++++++++- .../random/randutil/random_distributions.h | 24 +++++++- 9 files changed, 170 insertions(+), 34 deletions(-) diff --git a/src/cunumeric/random/randutil/generator.h b/src/cunumeric/random/randutil/generator.h index c56936503..1df425e52 100644 --- a/src/cunumeric/random/randutil/generator.h +++ b/src/cunumeric/random/randutil/generator.h @@ -19,6 +19,10 @@ #include #include +#ifdef USE_STL_RANDOM_ENGINE_ +#include +#endif + #include "legate.h" #include "randutil_curand.h" #include "randutil_impl.h" @@ -68,12 +72,25 @@ template struct inner_generator : basegenerator { uint64_t seed; uint64_t generatorID; + +#ifdef USE_STL_RANDOM_ENGINE_ + std::mt19937 generator; +#else gen_t generator; +#endif inner_generator(uint64_t seed, uint64_t generatorID, cudaStream_t ignored) - : seed(seed), generatorID(generatorID) + : seed(seed), + generatorID(generatorID) +#ifdef USE_STL_RANDOM_ENGINE_ + , + generator(seed), + std::srand(seed) +#endif { +#if !defined USE_STL_RANDOM_ENGINE_ curand_init(seed, generatorID, 0, &generator); +#endif } virtual void destroy() override {} @@ -105,6 +122,9 @@ curandStatus_t inner_dispatch_sample(basegenerator* gen, func_t func, size_t N, case CURAND_RNG_PSEUDO_MRG32K3A: return static_cast*>(gen) ->template draw(func, N, out); +#ifdef USE_STL_RANDOM_ENGINE_ + case STL_MT_19937: return; // TODO: add constant...somewhere +#endif default: LEGATE_ABORT; } return CURAND_STATUS_INTERNAL_ERROR; diff --git a/src/cunumeric/random/randutil/generator_gumbel.inl b/src/cunumeric/random/randutil/generator_gumbel.inl index 4a5d7b0b1..92ec172b5 100644 --- a/src/cunumeric/random/randutil/generator_gumbel.inl +++ b/src/cunumeric/random/randutil/generator_gumbel.inl @@ -27,7 +27,16 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else float y = curand_uniform(&gen); // y cannot be zero +#endif + if (y == 1.0f) return mu; float lny = ::logf(y); return mu - beta * ::logf(-lny); @@ -41,7 +50,16 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0, 1.0); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else double y = curand_uniform_double(&gen); // y cannot be zero +#endif + if (y == 1.0) return mu; double lny = ::log(y); return mu - beta * ::log(-lny); diff --git a/src/cunumeric/random/randutil/generator_integers.inl b/src/cunumeric/random/randutil/generator_integers.inl index cafb5e591..fca2414bf 100644 --- a/src/cunumeric/random/randutil/generator_integers.inl +++ b/src/cunumeric/random/randutil/generator_integers.inl @@ -15,45 +15,51 @@ */ #include "generator.h" +#include -template -struct integers; - -template <> -struct integers { - int16_t from; - int16_t to; +#include - 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; +#ifdef USE_STL_RANDOM_ENGINE_ + auto y = std::rand(); +#else + auto y = curand(&gen); +#endif + 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 +#ifdef USE_STL_RANDOM_ENGINE_ + unsigned low = std::rand(); + unsigned high = std::rand(); +#else unsigned low = curand(&gen); unsigned high = curand(&gen); - return (int64_t)((((uint64_t)high << 32) | (uint64_t)low) % (to - from)) + from; +#endif + + return (field_t)((((ufield_t)high << 32) | (ufield_t)low) % (to - from)) + from; } }; diff --git a/src/cunumeric/random/randutil/generator_negative_binomial.inl b/src/cunumeric/random/randutil/generator_negative_binomial.inl index 241f21dea..7f853e405 100644 --- a/src/cunumeric/random/randutil/generator_negative_binomial.inl +++ b/src/cunumeric/random/randutil/generator_negative_binomial.inl @@ -29,6 +29,11 @@ struct negative_binomial_t { RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { double lambda = rk_standard_gamma(&gen, (double)n) * ((1 - p) / p); +#ifdef USE_STL_RANDOM_ENGINE_ + std::poisson_distribution dis(lambda); + return dis(gen); +#else return curand_poisson(&gen, lambda); +#endif } }; diff --git a/src/cunumeric/random/randutil/generator_pareto.inl b/src/cunumeric/random/randutil/generator_pareto.inl index 2a29f2479..253716c50 100644 --- a/src/cunumeric/random/randutil/generator_pareto.inl +++ b/src/cunumeric/random/randutil/generator_pareto.inl @@ -26,7 +26,15 @@ struct pareto_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + float y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + float y = curand_uniform(&gen); // y cannot be 0 +#endif return xm * ::expf(-::logf(y) * invalpha) - 1.0f; // here, use -1.0f to align with numpy } }; @@ -38,7 +46,15 @@ struct pareto_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + double y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + double y = curand_uniform_double(&gen); // y cannot be 0 +#endif 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..4c9bf029b 100644 --- a/src/cunumeric/random/randutil/generator_poisson.inl +++ b/src/cunumeric/random/randutil/generator_poisson.inl @@ -22,6 +22,11 @@ struct poisson { template RANDUTIL_QUALIFIERS unsigned operator()(gen_t& gen) { +#ifdef USE_STL_RANDOM_ENGINE_ + std::poisson_distribution dis(lambda); + return dis(gen); +#else return curand_poisson(&gen, lambda); +#endif } }; diff --git a/src/cunumeric/random/randutil/generator_rayleigh.inl b/src/cunumeric/random/randutil/generator_rayleigh.inl index 3d34043b4..7fe58f106 100644 --- a/src/cunumeric/random/randutil/generator_rayleigh.inl +++ b/src/cunumeric/random/randutil/generator_rayleigh.inl @@ -26,7 +26,15 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = curand_uniform(&gen); // y cannot be 0 +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); // [0, 1) + auto y = dis(gen); + + // bring to (0, 1]; y cannot be 0: + y = 1 - y; +#else + auto y = curand_uniform(&gen); // returns (0, 1]; y cannot be 0 +#endif return sigma * ::sqrtf(-2.0f * ::logf(y)); } }; @@ -38,7 +46,15 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = curand_uniform_double(&gen); // y cannot be 0 +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0, 1.0); // [0, 1) + auto y = dis(gen); + + // bring to (0, 1]; y cannot be 0: + y = 1 - y; +#else + auto y = curand_uniform_double(&gen); // returns (0, 1]; y cannot be 0 +#endif return sigma * ::sqrt(-2.0 * ::log(y)); } }; diff --git a/src/cunumeric/random/randutil/generator_uniform.inl b/src/cunumeric/random/randutil/generator_uniform.inl index d6eecd52f..bd64cddd2 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -14,6 +14,18 @@ * */ +#ifndef LEGATE_USE_CUDA + +// assume on MACOS for testing: +// +#define IS_MAC_OS_ 1 + +#if IS_MAC_OS_ == 1 +#define USE_STL_RANDOM_ENGINE_ +#endif + +#endif + #include "generator.h" template @@ -26,7 +38,16 @@ struct uniform_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - return offset + mult * curand_uniform(&gen); +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + auto y = curand_uniform(&gen); // returns (0, 1]; +#endif + return offset + mult * y; } }; @@ -37,6 +58,15 @@ struct uniform_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - return offset + mult * curand_uniform_double(&gen); +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0, 1.0); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + auto y = curand_uniform_double(&gen); // returns (0, 1]; +#endif + return offset + mult * y; } }; diff --git a/src/cunumeric/random/randutil/random_distributions.h b/src/cunumeric/random/randutil/random_distributions.h index db4e609e4..2752cfbaa 100644 --- a/src/cunumeric/random/randutil/random_distributions.h +++ b/src/cunumeric/random/randutil/random_distributions.h @@ -52,7 +52,17 @@ template RANDUTIL_QUALIFIERS double rk_double(rk_state* gen) { - return curand_uniform_double(gen); +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0, 1.0); + auto y = dis(*gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + auto y = curand_uniform_double(gen); // returns (0, 1]; +#endif + + return y; } RANDUTIL_QUALIFIERS double loggam(double x) @@ -120,7 +130,12 @@ RANDUTIL_QUALIFIERS double rk_gauss(rk_state* state) return f*x2; } #endif +#ifdef USE_STL_RANDOM_ENGINE_ + std::normal_distribution dis{0.0, 1.0}; + return dis(*state); +#else return curand_normal(state); +#endif } template @@ -260,7 +275,12 @@ RANDUTIL_QUALIFIERS long rk_poisson(rk_state* state, double lam) return rk_poisson_mult(state, lam); } #endif +#ifdef USE_STL_RANDOM_ENGINE_ + std::poisson_distribution dis(lam); // curand::lambda == std::mu ? + return dis(*state); +#else return curand_poisson(state, lam); +#endif } template @@ -646,4 +666,4 @@ RANDUTIL_QUALIFIERS unsigned rk_binomial(rk_state* state, unsigned n, double p) } } -#pragma endregion \ No newline at end of file +#pragma endregion From 5c344e79ada189d5c8680377c007df2849e132d7 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Tue, 14 Nov 2023 19:56:15 -0600 Subject: [PATCH 02/17] Added an additional code insulation layer to simplify portability to STL. --- .../random/randutil/generator_gumbel.inl | 22 +-- .../random/randutil/generator_integers.inl | 17 +-- .../randutil/generator_negative_binomial.inl | 9 +- .../random/randutil/generator_pareto.inl | 26 +--- .../random/randutil/generator_poisson.inl | 9 +- .../random/randutil/generator_rayleigh.inl | 22 +-- .../random/randutil/generator_uniform.inl | 22 +-- .../random/randutil/random_distributions.h | 28 +--- src/cunumeric/random/randutil/randomizer.h | 127 ++++++++++++++++++ 9 files changed, 160 insertions(+), 122 deletions(-) create mode 100644 src/cunumeric/random/randutil/randomizer.h diff --git a/src/cunumeric/random/randutil/generator_gumbel.inl b/src/cunumeric/random/randutil/generator_gumbel.inl index 92ec172b5..ea2d1be75 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,15 +29,7 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); - auto y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - float y = curand_uniform(&gen); // y cannot be zero -#endif + auto y = randutilimpl::engine_uniform_single(gen); // y cannot be zero if (y == 1.0f) return mu; float lny = ::logf(y); @@ -50,15 +44,7 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0, 1.0); - auto y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - double y = curand_uniform_double(&gen); // y cannot be zero -#endif + auto y = randutilimpl::engine_uniform_double(gen); // y cannot be zero if (y == 1.0) return mu; double lny = ::log(y); diff --git a/src/cunumeric/random/randutil/generator_integers.inl b/src/cunumeric/random/randutil/generator_integers.inl index fca2414bf..74b4f8764 100644 --- a/src/cunumeric/random/randutil/generator_integers.inl +++ b/src/cunumeric/random/randutil/generator_integers.inl @@ -17,7 +17,7 @@ #include "generator.h" #include -#include +#include "randomizer.h" template struct integers; @@ -33,11 +33,7 @@ struct integers< template RANDUTIL_QUALIFIERS field_t operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - auto y = std::rand(); -#else - auto y = curand(&gen); -#endif + auto y = randutilimpl::engine_rand(gen); return (field_t)(y % (ufield_t)(to - from)) + from; } }; @@ -52,13 +48,8 @@ struct integers>> { RANDUTIL_QUALIFIERS field_t operator()(gen_t& gen) { // take two draws to get a 64 bits value -#ifdef USE_STL_RANDOM_ENGINE_ - unsigned low = std::rand(); - unsigned high = std::rand(); -#else - unsigned low = curand(&gen); - unsigned high = curand(&gen); -#endif + 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_negative_binomial.inl b/src/cunumeric/random/randutil/generator_negative_binomial.inl index 7f853e405..eb33a5975 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,11 +31,6 @@ struct negative_binomial_t { RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { double lambda = rk_standard_gamma(&gen, (double)n) * ((1 - p) / p); -#ifdef USE_STL_RANDOM_ENGINE_ - std::poisson_distribution dis(lambda); - return dis(gen); -#else - return curand_poisson(&gen, lambda); -#endif + return randutilimpl::engine_poisson(gen); } }; diff --git a/src/cunumeric/random/randutil/generator_pareto.inl b/src/cunumeric/random/randutil/generator_pareto.inl index 253716c50..381169037 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,16 +28,8 @@ struct pareto_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); - float y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - float y = curand_uniform(&gen); // y cannot be 0 -#endif - return xm * ::expf(-::logf(y) * invalpha) - 1.0f; // here, use -1.0f to align with numpy + auto y = randutilimpl::engine_uniform_single(gen); // y cannot be 0 + return xm * ::expf(-::logf(y) * invalpha) - 1.0f; // here, use -1.0f to align with numpy } }; @@ -46,15 +40,7 @@ struct pareto_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); - double y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - double y = curand_uniform_double(&gen); // y cannot be 0 -#endif - return xm * ::exp(-::log(y) * invalpha) - 1.0; // here, use -1.0 to align with numpy + auto y = randutilimpl::engine_uniform_double(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 4c9bf029b..7ec9e066f 100644 --- a/src/cunumeric/random/randutil/generator_poisson.inl +++ b/src/cunumeric/random/randutil/generator_poisson.inl @@ -16,17 +16,14 @@ #include "generator.h" +#include "randomizer.h" + struct poisson { double lambda = 1.0; template RANDUTIL_QUALIFIERS unsigned operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::poisson_distribution dis(lambda); - return dis(gen); -#else - return curand_poisson(&gen, lambda); -#endif + return randutilimpl::engine_poisson(gen, lambda); } }; diff --git a/src/cunumeric/random/randutil/generator_rayleigh.inl b/src/cunumeric/random/randutil/generator_rayleigh.inl index 7fe58f106..f5bc5a5a8 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,15 +28,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); // [0, 1) - auto y = dis(gen); - - // bring to (0, 1]; y cannot be 0: - y = 1 - y; -#else - auto y = curand_uniform(&gen); // returns (0, 1]; y cannot be 0 -#endif + auto y = randutilimpl::engine_uniform_single(gen); // returns (0, 1]; y cannot be 0 return sigma * ::sqrtf(-2.0f * ::logf(y)); } }; @@ -46,15 +40,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0, 1.0); // [0, 1) - auto y = dis(gen); - - // bring to (0, 1]; y cannot be 0: - y = 1 - y; -#else - auto y = curand_uniform_double(&gen); // returns (0, 1]; y cannot be 0 -#endif + auto y = randutilimpl::engine_uniform_double(gen); // returns (0, 1]; y cannot be 0 return sigma * ::sqrt(-2.0 * ::log(y)); } }; diff --git a/src/cunumeric/random/randutil/generator_uniform.inl b/src/cunumeric/random/randutil/generator_uniform.inl index bd64cddd2..9a37f279e 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -28,6 +28,8 @@ #include "generator.h" +#include "randomizer.h" + template struct uniform_t; @@ -38,15 +40,7 @@ struct uniform_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); - auto y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - auto y = curand_uniform(&gen); // returns (0, 1]; -#endif + auto y = randutilimpl::engine_uniform_single(gen); // returns (0, 1]; return offset + mult * y; } }; @@ -58,15 +52,7 @@ struct uniform_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0, 1.0); - auto y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - auto y = curand_uniform_double(&gen); // returns (0, 1]; -#endif + auto y = randutilimpl::engine_uniform_double(gen); // returns (0, 1]; return offset + mult * y; } }; diff --git a/src/cunumeric/random/randutil/random_distributions.h b/src/cunumeric/random/randutil/random_distributions.h index 2752cfbaa..4b3dbfac7 100644 --- a/src/cunumeric/random/randutil/random_distributions.h +++ b/src/cunumeric/random/randutil/random_distributions.h @@ -49,20 +49,12 @@ #pragma once +#include "randomizer.h" + template RANDUTIL_QUALIFIERS double rk_double(rk_state* gen) { -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0, 1.0); - auto y = dis(*gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - auto y = curand_uniform_double(gen); // returns (0, 1]; -#endif - - return y; + return randutilimpl::engine_uniform_double(*gen); // returns (0, 1]; } RANDUTIL_QUALIFIERS double loggam(double x) @@ -130,12 +122,7 @@ RANDUTIL_QUALIFIERS double rk_gauss(rk_state* state) return f*x2; } #endif -#ifdef USE_STL_RANDOM_ENGINE_ - std::normal_distribution dis{0.0, 1.0}; - return dis(*state); -#else - return curand_normal(state); -#endif + return randutilimpl::engine_normal(*state); } template @@ -275,12 +262,7 @@ RANDUTIL_QUALIFIERS long rk_poisson(rk_state* state, double lam) return rk_poisson_mult(state, lam); } #endif -#ifdef USE_STL_RANDOM_ENGINE_ - std::poisson_distribution dis(lam); // curand::lambda == std::mu ? - return dis(*state); -#else - return curand_poisson(state, lam); -#endif + return randutilimpl::engine_poisson(*state, lam); } template diff --git a/src/cunumeric/random/randutil/randomizer.h b/src/cunumeric/random/randutil/randomizer.h new file mode 100644 index 000000000..a638b8eb5 --- /dev/null +++ b/src/cunumeric/random/randutil/randomizer.h @@ -0,0 +1,127 @@ +/* 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 + +namespace randutilimpl { + +// trampoline functions for branching-off CURAND vs. STL +// implementations (STL used on platforms that don't support CUDA) +// +template +decltype(auto) engine_uniform_single(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + auto y = curand_uniform(&gen); // returns (0, 1]; +#endif + return y; +} + +template +decltype(auto) engine_uniform_double(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::uniform_real_distribution dis(0.0f, 1.0f); + auto y = dis(gen); // returns [0, 1); + + // bring to (0, 1]: + y = 1 - y; +#else + auto y = curand_uniform_double(&gen); // returns (0, 1]; +#endif + return y; +} + +template +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 +decltype(auto) engine_normal(gen_t& gen) +{ +#ifdef USE_STL_RANDOM_ENGINE_ + std::normal_distribution dis{0.0, 1.0}; + return dis(gen); +#else + return curand_normal(&gen); +#endif +} + +template +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 From 50b2f31d45d4bd1a97c1734f56926675ae911d9f Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 15 Nov 2023 11:54:13 -0600 Subject: [PATCH 03/17] Fixes to STL insulation layer. --- src/cunumeric/random/randutil/generator_integers.inl | 4 ++-- .../random/randutil/generator_negative_binomial.inl | 2 +- src/cunumeric/random/randutil/generator_poisson.inl | 2 +- src/cunumeric/random/randutil/randomizer.h | 10 +++++----- 4 files changed, 9 insertions(+), 9 deletions(-) diff --git a/src/cunumeric/random/randutil/generator_integers.inl b/src/cunumeric/random/randutil/generator_integers.inl index 74b4f8764..7227ad63c 100644 --- a/src/cunumeric/random/randutil/generator_integers.inl +++ b/src/cunumeric/random/randutil/generator_integers.inl @@ -22,7 +22,7 @@ template struct integers; -template +template struct integers< field_t, std::enable_if_t || std::is_same_v>> { @@ -38,7 +38,7 @@ struct integers< } }; -template +template struct integers>> { using ufield_t = uint64_t; field_t from; diff --git a/src/cunumeric/random/randutil/generator_negative_binomial.inl b/src/cunumeric/random/randutil/generator_negative_binomial.inl index eb33a5975..563565983 100644 --- a/src/cunumeric/random/randutil/generator_negative_binomial.inl +++ b/src/cunumeric/random/randutil/generator_negative_binomial.inl @@ -31,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 randutilimpl::engine_poisson(gen); + return randutilimpl::engine_poisson(gen, lambda); } }; diff --git a/src/cunumeric/random/randutil/generator_poisson.inl b/src/cunumeric/random/randutil/generator_poisson.inl index 7ec9e066f..f8f238f20 100644 --- a/src/cunumeric/random/randutil/generator_poisson.inl +++ b/src/cunumeric/random/randutil/generator_poisson.inl @@ -24,6 +24,6 @@ struct poisson { template RANDUTIL_QUALIFIERS unsigned operator()(gen_t& gen) { - return randutilimpl::engine_poisson(gen, lambda); + return randutilimpl::engine_poisson(gen, lambda); } }; diff --git a/src/cunumeric/random/randutil/randomizer.h b/src/cunumeric/random/randutil/randomizer.h index a638b8eb5..b0af68c21 100644 --- a/src/cunumeric/random/randutil/randomizer.h +++ b/src/cunumeric/random/randutil/randomizer.h @@ -27,7 +27,7 @@ namespace randutilimpl { // implementations (STL used on platforms that don't support CUDA) // template -decltype(auto) engine_uniform_single(gen_t& gen) +RANDUTIL_QUALIFIERS decltype(auto) engine_uniform_single(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ std::uniform_real_distribution dis(0.0f, 1.0f); @@ -42,7 +42,7 @@ decltype(auto) engine_uniform_single(gen_t& gen) } template -decltype(auto) engine_uniform_double(gen_t& gen) +RANDUTIL_QUALIFIERS decltype(auto) engine_uniform_double(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ std::uniform_real_distribution dis(0.0f, 1.0f); @@ -57,7 +57,7 @@ decltype(auto) engine_uniform_double(gen_t& gen) } template -decltype(auto) engine_poisson(gen_t& gen, double lambda) +RANDUTIL_QUALIFIERS decltype(auto) engine_poisson(gen_t& gen, double lambda) { #ifdef USE_STL_RANDOM_ENGINE_ std::poisson_distribution dis(lambda); @@ -68,7 +68,7 @@ decltype(auto) engine_poisson(gen_t& gen, double lambda) } template -decltype(auto) engine_normal(gen_t& gen) +RANDUTIL_QUALIFIERS decltype(auto) engine_normal(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ std::normal_distribution dis{0.0, 1.0}; @@ -79,7 +79,7 @@ decltype(auto) engine_normal(gen_t& gen) } template -decltype(auto) engine_rand(gen_t& gen) +RANDUTIL_QUALIFIERS decltype(auto) engine_rand(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ return std::rand(); From a2492bea5823a2b0b4b9726ee2ae546a5298fc3b Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 15 Nov 2023 17:39:31 -0600 Subject: [PATCH 04/17] Distributions STL split. Group 2. --- .../random/randutil/generator_cauchy.inl | 6 ++++-- .../random/randutil/generator_exponential.inl | 6 ++++-- .../random/randutil/generator_laplace.inl | 6 ++++-- .../random/randutil/generator_logistic.inl | 8 +++++--- .../random/randutil/generator_lognormal.inl | 6 ++++-- .../randutil/generator_negative_binomial.inl | 2 +- src/cunumeric/random/randutil/generator_raw.inl | 4 +++- .../random/randutil/generator_uniform.inl | 8 ++++---- .../random/randutil/generator_weibull.inl | 10 ++++++---- src/cunumeric/random/randutil/randomizer.h | 15 +++++++++++++++ 10 files changed, 50 insertions(+), 21 deletions(-) diff --git a/src/cunumeric/random/randutil/generator_cauchy.inl b/src/cunumeric/random/randutil/generator_cauchy.inl index f0b5e2b12..a76eff39b 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_single(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_double(gen); // y cannot be 0 return x0 + gamma * ::tan(pi * (y - 0.5)); } }; diff --git a/src/cunumeric/random/randutil/generator_exponential.inl b/src/cunumeric/random/randutil/generator_exponential.inl index 5bc4847b2..453020a02 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_single(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_double(gen); return -::logf(uni) * scale; } }; diff --git a/src/cunumeric/random/randutil/generator_laplace.inl b/src/cunumeric/random/randutil/generator_laplace.inl index 0d4c7b1ae..aac92dd9a 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_single(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_double(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..cb980a52a 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_single(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_double(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 563565983..5521af420 100644 --- a/src/cunumeric/random/randutil/generator_negative_binomial.inl +++ b/src/cunumeric/random/randutil/generator_negative_binomial.inl @@ -31,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 randutilimpl::engine_poisson(gen, lambda); + return static_cast(randutilimpl::engine_poisson(gen, lambda)); } }; 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_uniform.inl b/src/cunumeric/random/randutil/generator_uniform.inl index 9a37f279e..d6c4689e7 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -15,15 +15,15 @@ */ #ifndef LEGATE_USE_CUDA - +//{- // assume on MACOS for testing: // -#define IS_MAC_OS_ 1 - +#define IS_MAC_OS_ 0 +// #if IS_MAC_OS_ == 1 #define USE_STL_RANDOM_ENGINE_ #endif - +//-} #endif #include "generator.h" diff --git a/src/cunumeric/random/randutil/generator_weibull.inl b/src/cunumeric/random/randutil/generator_weibull.inl index 5565c7411..78455a60a 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_single(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_double(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/randomizer.h b/src/cunumeric/random/randutil/randomizer.h index b0af68c21..138b39f1a 100644 --- a/src/cunumeric/random/randutil/randomizer.h +++ b/src/cunumeric/random/randutil/randomizer.h @@ -18,6 +18,7 @@ #include #include +#include #include @@ -78,6 +79,20 @@ RANDUTIL_QUALIFIERS decltype(auto) engine_normal(gen_t& 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) { From f2198ce725610ca74186f70ea23eafd9178181a2 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 15 Nov 2023 18:25:17 -0600 Subject: [PATCH 05/17] Normal distribution refactored. --- src/cunumeric/random/randutil/generator_normal.inl | 6 ++++-- src/cunumeric/random/randutil/random_distributions.h | 2 +- src/cunumeric/random/randutil/randomizer.h | 9 ++++++--- 3 files changed, 11 insertions(+), 6 deletions(-) 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/random_distributions.h b/src/cunumeric/random/randutil/random_distributions.h index 4b3dbfac7..07fbc7462 100644 --- a/src/cunumeric/random/randutil/random_distributions.h +++ b/src/cunumeric/random/randutil/random_distributions.h @@ -122,7 +122,7 @@ RANDUTIL_QUALIFIERS double rk_gauss(rk_state* state) return f*x2; } #endif - return randutilimpl::engine_normal(*state); + return randutilimpl::engine_normal(*state); } template diff --git a/src/cunumeric/random/randutil/randomizer.h b/src/cunumeric/random/randutil/randomizer.h index 138b39f1a..c2897079f 100644 --- a/src/cunumeric/random/randutil/randomizer.h +++ b/src/cunumeric/random/randutil/randomizer.h @@ -68,14 +68,17 @@ RANDUTIL_QUALIFIERS decltype(auto) engine_poisson(gen_t& gen, double lambda) #endif } -template +template RANDUTIL_QUALIFIERS decltype(auto) engine_normal(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ - std::normal_distribution dis{0.0, 1.0}; + std::normal_distribution dis(0, 1); return dis(gen); #else - return curand_normal(&gen); + if constexpr (std::is_same_v) + return curand_normal(&gen); + else + return curand_normal_double(&gen); #endif } From 3df844a5d469bd39e69f6a20131d07a652ac8161 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 15 Nov 2023 19:09:32 -0600 Subject: [PATCH 06/17] Distributions STL split. Group 3. --- src/cunumeric/random/randutil/generator_power.inl | 6 ++++-- src/cunumeric/random/randutil/generator_triangular.inl | 6 ++++-- src/cunumeric/random/randutil/generator_wald.inl | 10 ++++++---- 3 files changed, 14 insertions(+), 8 deletions(-) diff --git a/src/cunumeric/random/randutil/generator_power.inl b/src/cunumeric/random/randutil/generator_power.inl index e17eafc55..0199dd85d 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_single(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_double(gen); // y cannot be 0 -- use y as 1-cdf(x) return ::exp(::log(y) * invalpha); } }; diff --git a/src/cunumeric/random/randutil/generator_triangular.inl b/src/cunumeric/random/randutil/generator_triangular.inl index 3bfe77e28..a53534bf6 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_single(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_double(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_wald.inl b/src/cunumeric/random/randutil/generator_wald.inl index dec310834..5ae7f8dfe 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_single(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_double(gen); if (z <= (mu) / (mu + x)) return x; else From e55921a162ae03e7c04182005434bab1be4b3a20 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Thu, 16 Nov 2023 16:39:00 -0600 Subject: [PATCH 07/17] Better engine_uniform() function template. --- .../random/randutil/generator_cauchy.inl | 4 +-- .../random/randutil/generator_exponential.inl | 4 +-- .../random/randutil/generator_gumbel.inl | 4 +-- .../random/randutil/generator_laplace.inl | 4 +-- .../random/randutil/generator_logistic.inl | 4 +-- .../random/randutil/generator_pareto.inl | 4 +-- .../random/randutil/generator_power.inl | 4 +-- .../random/randutil/generator_rayleigh.inl | 4 +-- .../random/randutil/generator_triangular.inl | 4 +-- .../random/randutil/generator_uniform.inl | 4 +-- .../random/randutil/generator_wald.inl | 4 +-- .../random/randutil/generator_weibull.inl | 4 +-- .../random/randutil/random_distributions.h | 2 +- src/cunumeric/random/randutil/randomizer.h | 29 +++++-------------- 14 files changed, 33 insertions(+), 46 deletions(-) diff --git a/src/cunumeric/random/randutil/generator_cauchy.inl b/src/cunumeric/random/randutil/generator_cauchy.inl index a76eff39b..ece9843d8 100644 --- a/src/cunumeric/random/randutil/generator_cauchy.inl +++ b/src/cunumeric/random/randutil/generator_cauchy.inl @@ -30,7 +30,7 @@ struct cauchy_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 return x0 + gamma * ::tanf(pi * (y - 0.5f)); } }; @@ -44,7 +44,7 @@ struct cauchy_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = randutilimpl::engine_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_exponential.inl b/src/cunumeric/random/randutil/generator_exponential.inl index 453020a02..4a01ff40e 100644 --- a/src/cunumeric/random/randutil/generator_exponential.inl +++ b/src/cunumeric/random/randutil/generator_exponential.inl @@ -28,7 +28,7 @@ struct exponential_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float uni = randutilimpl::engine_uniform_single(gen); + float uni = randutilimpl::engine_uniform(gen); return -::logf(uni) * scale; } }; @@ -40,7 +40,7 @@ struct exponential_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double uni = randutilimpl::engine_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 ea2d1be75..32a3b0d05 100644 --- a/src/cunumeric/random/randutil/generator_gumbel.inl +++ b/src/cunumeric/random/randutil/generator_gumbel.inl @@ -29,7 +29,7 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_single(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); @@ -44,7 +44,7 @@ struct gumbel_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - auto y = randutilimpl::engine_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); diff --git a/src/cunumeric/random/randutil/generator_laplace.inl b/src/cunumeric/random/randutil/generator_laplace.inl index aac92dd9a..9177c948f 100644 --- a/src/cunumeric/random/randutil/generator_laplace.inl +++ b/src/cunumeric/random/randutil/generator_laplace.inl @@ -28,7 +28,7 @@ struct laplace_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(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); @@ -44,7 +44,7 @@ struct laplace_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = randutilimpl::engine_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 cb980a52a..d37d515ca 100644 --- a/src/cunumeric/random/randutil/generator_logistic.inl +++ b/src/cunumeric/random/randutil/generator_logistic.inl @@ -28,7 +28,7 @@ struct logistic_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(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); @@ -42,7 +42,7 @@ struct logistic_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_double(gen); // y cannot be 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_pareto.inl b/src/cunumeric/random/randutil/generator_pareto.inl index 381169037..f99c74252 100644 --- a/src/cunumeric/random/randutil/generator_pareto.inl +++ b/src/cunumeric/random/randutil/generator_pareto.inl @@ -28,7 +28,7 @@ struct pareto_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_single(gen); // y cannot be 0 + 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 } }; @@ -40,7 +40,7 @@ struct pareto_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_double(gen); // y cannot be 0 + 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_power.inl b/src/cunumeric/random/randutil/generator_power.inl index 0199dd85d..7c3c278a2 100644 --- a/src/cunumeric/random/randutil/generator_power.inl +++ b/src/cunumeric/random/randutil/generator_power.inl @@ -28,7 +28,7 @@ struct power_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 return ::expf(::logf(y) * invalpha); } }; @@ -40,7 +40,7 @@ struct power_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = randutilimpl::engine_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_rayleigh.inl b/src/cunumeric/random/randutil/generator_rayleigh.inl index f5bc5a5a8..d1a7b2cef 100644 --- a/src/cunumeric/random/randutil/generator_rayleigh.inl +++ b/src/cunumeric/random/randutil/generator_rayleigh.inl @@ -28,7 +28,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_single(gen); // returns (0, 1]; y cannot be 0 + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; y cannot be 0 return sigma * ::sqrtf(-2.0f * ::logf(y)); } }; @@ -40,7 +40,7 @@ struct rayleigh_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_double(gen); // returns (0, 1]; 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 a53534bf6..318d569eb 100644 --- a/src/cunumeric/random/randutil/generator_triangular.inl +++ b/src/cunumeric/random/randutil/generator_triangular.inl @@ -28,7 +28,7 @@ struct triangular_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(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; @@ -48,7 +48,7 @@ struct triangular_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = randutilimpl::engine_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 d6c4689e7..99ce04c83 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -40,7 +40,7 @@ struct uniform_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_single(gen); // returns (0, 1]; + auto y = randutilimpl::engine_uniform(gen); // returns (0, 1]; return offset + mult * y; } }; @@ -52,7 +52,7 @@ struct uniform_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - auto y = randutilimpl::engine_uniform_double(gen); // returns (0, 1]; + 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 5ae7f8dfe..20706e6ce 100644 --- a/src/cunumeric/random/randutil/generator_wald.inl +++ b/src/cunumeric/random/randutil/generator_wald.inl @@ -33,7 +33,7 @@ struct wald_t { 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 = randutilimpl::engine_uniform_single(gen); + float z = randutilimpl::engine_uniform(gen); if (z <= (mu) / (mu + x)) return x; else @@ -52,7 +52,7 @@ struct wald_t { 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 = randutilimpl::engine_uniform_double(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 78455a60a..ad1ec836d 100644 --- a/src/cunumeric/random/randutil/generator_weibull.inl +++ b/src/cunumeric/random/randutil/generator_weibull.inl @@ -28,7 +28,7 @@ struct weibull_t { template RANDUTIL_QUALIFIERS float operator()(gen_t& gen) { - float y = randutilimpl::engine_uniform_single(gen); // y cannot be 0 + float y = randutilimpl::engine_uniform(gen); // y cannot be 0 // log(y) can be zero ! auto lny = ::logf(y); if (lny == 0.0f) return 0.0f; @@ -43,7 +43,7 @@ struct weibull_t { template RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { - double y = randutilimpl::engine_uniform_double(gen); // y cannot be 0 + double y = randutilimpl::engine_uniform(gen); // y cannot be 0 // log(y) can be zero ! auto lny = ::log(y); if (lny == 0.0f) return 0.0f; diff --git a/src/cunumeric/random/randutil/random_distributions.h b/src/cunumeric/random/randutil/random_distributions.h index 07fbc7462..5307a5858 100644 --- a/src/cunumeric/random/randutil/random_distributions.h +++ b/src/cunumeric/random/randutil/random_distributions.h @@ -54,7 +54,7 @@ template RANDUTIL_QUALIFIERS double rk_double(rk_state* gen) { - return randutilimpl::engine_uniform_double(*gen); // returns (0, 1]; + return randutilimpl::engine_uniform(*gen); // returns (0, 1]; } RANDUTIL_QUALIFIERS double loggam(double x) diff --git a/src/cunumeric/random/randutil/randomizer.h b/src/cunumeric/random/randutil/randomizer.h index c2897079f..c9bcda026 100644 --- a/src/cunumeric/random/randutil/randomizer.h +++ b/src/cunumeric/random/randutil/randomizer.h @@ -27,34 +27,21 @@ 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_single(gen_t& gen) -{ -#ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); - auto y = dis(gen); // returns [0, 1); - - // bring to (0, 1]: - y = 1 - y; -#else - auto y = curand_uniform(&gen); // returns (0, 1]; -#endif - return y; -} - -template -RANDUTIL_QUALIFIERS decltype(auto) engine_uniform_double(gen_t& gen) +template +RANDUTIL_QUALIFIERS decltype(auto) engine_uniform(gen_t& gen) { #ifdef USE_STL_RANDOM_ENGINE_ - std::uniform_real_distribution dis(0.0f, 1.0f); + std::uniform_real_distribution dis(0, 1); auto y = dis(gen); // returns [0, 1); // bring to (0, 1]: - y = 1 - y; + return 1 - y; #else - auto y = curand_uniform_double(&gen); // returns (0, 1]; + if constexpr (std::is_same_v) + return curand_uniform(&gen); // returns (0, 1]; + else + return curand_uniform_double(&gen); // returns (0, 1]; #endif - return y; } template From d34ecb5fe88e2ac2e8796e3ccce42f14de6d2136 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Mon, 20 Nov 2023 18:24:18 -0600 Subject: [PATCH 08/17] Refactoring for curand abstraction replacements. Step 1. --- src/cunumeric/random/bitgenerator.cc | 17 ++- src/cunumeric/random/bitgenerator_curand.inl | 111 ++++++++++--------- src/cunumeric/random/rnd_types.h | 63 +++++++++++ 3 files changed, 133 insertions(+), 58 deletions(-) create mode 100644 src/cunumeric/random/rnd_types.h diff --git a/src/cunumeric/random/bitgenerator.cc b/src/cunumeric/random/bitgenerator.cc index 1038a5309..c93abf4a7 100644 --- a/src/cunumeric/random/bitgenerator.cc +++ b/src/cunumeric/random/bitgenerator.cc @@ -18,7 +18,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" #include "cunumeric/random/bitgenerator_curand.inl" @@ -31,6 +31,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 +49,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_curand.inl b/src/cunumeric/random/bitgenerator_curand.inl index f74c87ad0..6f4ffa527 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 { @@ -45,7 +45,7 @@ struct CURANDGenerator { 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/rnd_types.h b/src/cunumeric/random/rnd_types.h new file mode 100644 index 000000000..48b9d9b55 --- /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 + +#ifdef USE_STL_RANDOM_ENGINE_ +using rnd_status_t = int; +enum class randRngType : int { STL_MT19937 = 1 }; + +#define CHECK_RND_ENGINE(expr) \ + do { \ + rnd_status_t __result__ = (expr); \ + randutil_check_status(__result__, __FILE__, __LINE__); \ + } while (false) + +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_MT19937; + case cunumeric::BitGeneratorType::XORWOW: return randRngType::STL_MT19937; + case cunumeric::BitGeneratorType::MRG32K3A: return randRngType::STL_MT19937; + case cunumeric::BitGeneratorType::MTGP32: return randRngType::STL_MT19937; + case cunumeric::BitGeneratorType::MT19937: return randRngType::STL_MT19937; + case cunumeric::BitGeneratorType::PHILOX4_32_10: return randRngType::STL_MT19937; + default: LEGATE_ABORT; + } + return randRngType::CURAND_RNG_TEST; +} + +} // namespace cunumeric + +#else +#include "cunumeric/random/curand_help.h" +using rnd_status_t = curandStatus_t; +using randRngType = curandRngType; + +#define CHECK_RND_ENGINE(expr) CHECK_CURAND((expr)) +#define randutil_check_status randutil_check_curand +#define get_rndRngType get_curandRngType + +#endif From 3b998488121e1e872b9a08db68aafebf44e88420 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Tue, 21 Nov 2023 11:48:28 -0600 Subject: [PATCH 09/17] Refactoring for curand abstraction replacements. Step 2. --- src/cunumeric/random/randutil/generator.h | 4 ---- src/cunumeric/random/randutil/randutil_curand.h | 6 +++++- src/cunumeric/random/rnd_types.h | 8 ++++++-- 3 files changed, 11 insertions(+), 7 deletions(-) diff --git a/src/cunumeric/random/randutil/generator.h b/src/cunumeric/random/randutil/generator.h index 1df425e52..435244560 100644 --- a/src/cunumeric/random/randutil/generator.h +++ b/src/cunumeric/random/randutil/generator.h @@ -19,10 +19,6 @@ #include #include -#ifdef USE_STL_RANDOM_ENGINE_ -#include -#endif - #include "legate.h" #include "randutil_curand.h" #include "randutil_impl.h" 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_types.h b/src/cunumeric/random/rnd_types.h index 48b9d9b55..96a12fcb2 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -19,6 +19,8 @@ #ifdef USE_STL_RANDOM_ENGINE_ using rnd_status_t = int; enum class randRngType : int { STL_MT19937 = 1 }; +using randRngType_t = randRngType; +constexpr int RND_STATUS_SUCCESS = 0; #define CHECK_RND_ENGINE(expr) \ do { \ @@ -53,8 +55,10 @@ static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) #else #include "cunumeric/random/curand_help.h" -using rnd_status_t = curandStatus_t; -using randRngType = curandRngType; +using rnd_status_t = curandStatus_t; +using randRngType = curandRngType; +using randRngType_t = curandRngType_t; +constexpr rnd_status_t RND_STATUS_SUCCESS = CURAND_STATUS_SUCCESS; #define CHECK_RND_ENGINE(expr) CHECK_CURAND((expr)) #define randutil_check_status randutil_check_curand From 3331e337d16626cc5a770e5da09be697271f4c54 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Tue, 21 Nov 2023 15:51:12 -0600 Subject: [PATCH 10/17] Refactoring for curand abstraction replacements. Step 3. --- .../random/randutil/generator_host.cc | 47 ++++++++++--------- src/cunumeric/random/rnd_aliases.h | 29 ++++++++++++ src/cunumeric/random/rnd_types.h | 13 ++--- 3 files changed, 58 insertions(+), 31 deletions(-) create mode 100644 src/cunumeric/random/rnd_aliases.h diff --git a/src/cunumeric/random/randutil/generator_host.cc b/src/cunumeric/random/randutil/generator_host.cc index 50bcfeefc..7d3f456e1 100644 --- a/src/cunumeric/random/randutil/generator_host.cc +++ b/src/cunumeric/random/randutil/generator_host.cc @@ -16,6 +16,7 @@ #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 +25,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 +57,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 +67,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 +83,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 +93,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 +109,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 +119,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 +135,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 +152,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 +167,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 +178,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/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h new file mode 100644 index 000000000..693d793fd --- /dev/null +++ b/src/cunumeric/random/rnd_aliases.h @@ -0,0 +1,29 @@ +/* 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_ +using rnd_status_t = int; +enum class randRngType : int { STL_MT19937 = 1 }; +using randRngType_t = randRngType; +constexpr int RND_STATUS_SUCCESS = 0; +#else +using rnd_status_t = curandStatus_t; +using randRngType = curandRngType; +using randRngType_t = curandRngType_t; +constexpr rnd_status_t RND_STATUS_SUCCESS = CURAND_STATUS_SUCCESS; +#endif diff --git a/src/cunumeric/random/rnd_types.h b/src/cunumeric/random/rnd_types.h index 96a12fcb2..1af40c82a 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -17,10 +17,7 @@ #pragma once #ifdef USE_STL_RANDOM_ENGINE_ -using rnd_status_t = int; -enum class randRngType : int { STL_MT19937 = 1 }; -using randRngType_t = randRngType; -constexpr int RND_STATUS_SUCCESS = 0; +#include "cunumeric/random/rnd_aliases.h" #define CHECK_RND_ENGINE(expr) \ do { \ @@ -55,10 +52,10 @@ static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) #else #include "cunumeric/random/curand_help.h" -using rnd_status_t = curandStatus_t; -using randRngType = curandRngType; -using randRngType_t = curandRngType_t; -constexpr rnd_status_t RND_STATUS_SUCCESS = CURAND_STATUS_SUCCESS; + +// in this case include aliases _after_ curand_help +// +#include "cunumeric/random/rnd_aliases.h" #define CHECK_RND_ENGINE(expr) CHECK_CURAND((expr)) #define randutil_check_status randutil_check_curand From b343896903c6d86a1e3a373011e7d8f3de2a025b Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 22 Nov 2023 14:05:57 -0600 Subject: [PATCH 11/17] Refactoring for curand abstraction replacements. Step 4. --- src/cunumeric/random/randutil/generator.h | 68 +++++++++++------------ src/cunumeric/random/rnd_aliases.h | 39 ++++++++++--- src/cunumeric/random/rnd_types.h | 12 ++-- 3 files changed, 71 insertions(+), 48 deletions(-) diff --git a/src/cunumeric/random/randutil/generator.h b/src/cunumeric/random/randutil/generator.h index 435244560..93f42706c 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,44 +38,43 @@ 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; }; 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; }; +#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 { uint64_t seed; uint64_t generatorID; - -#ifdef USE_STL_RANDOM_ENGINE_ - std::mt19937 generator; -#else gen_t generator; -#endif inner_generator(uint64_t seed, uint64_t generatorID, cudaStream_t ignored) : seed(seed), @@ -106,39 +107,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); -#ifdef USE_STL_RANDOM_ENGINE_ - case STL_MT_19937: return; // TODO: add constant...somewhere -#endif 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); @@ -146,7 +144,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: @@ -157,7 +155,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/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h index 693d793fd..81b06bd26 100644 --- a/src/cunumeric/random/rnd_aliases.h +++ b/src/cunumeric/random/rnd_aliases.h @@ -17,13 +17,38 @@ #pragma once #ifdef USE_STL_RANDOM_ENGINE_ +#include + using rnd_status_t = int; -enum class randRngType : int { STL_MT19937 = 1 }; -using randRngType_t = randRngType; -constexpr int RND_STATUS_SUCCESS = 0; +enum class randRngType : int { STL_MT_19937 = 1 }; +using randRngType_t = randRngType; +constexpr int RND_STATUS_SUCCESS = 0; +constexpr rnd_status_t RND_STATUS_INTERNAL_ERROR = 1; //??? + +namespace randutilimpl { +constexpr int RND_RNG_PSEUDO_XORWOW = randRngType::STL_MT_19937; +constexpr int RND_RNG_PSEUDO_PHILOX4_32_10 = randRngType::STL_MT_19937; +constexpr int RND_RNG_PSEUDO_MRG32K3A = randRngType::STL_MT_19937; + +using gen_XORWOW_t = std::mt19937; +using gen_Philox4_32_10_t = std::mt19937; +using gen_MRG32k3a_t = std::mt19937; +} // namespace randutilimpl #else -using rnd_status_t = curandStatus_t; -using randRngType = curandRngType; -using randRngType_t = curandRngType_t; -constexpr rnd_status_t RND_STATUS_SUCCESS = CURAND_STATUS_SUCCESS; +#include + +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 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; + #endif diff --git a/src/cunumeric/random/rnd_types.h b/src/cunumeric/random/rnd_types.h index 1af40c82a..29403a726 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -37,12 +37,12 @@ static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) // this might become more flexible in the future; // switch (kind) { - case cunumeric::BitGeneratorType::DEFAULT: return randRngType::STL_MT19937; - case cunumeric::BitGeneratorType::XORWOW: return randRngType::STL_MT19937; - case cunumeric::BitGeneratorType::MRG32K3A: return randRngType::STL_MT19937; - case cunumeric::BitGeneratorType::MTGP32: return randRngType::STL_MT19937; - case cunumeric::BitGeneratorType::MT19937: return randRngType::STL_MT19937; - case cunumeric::BitGeneratorType::PHILOX4_32_10: return randRngType::STL_MT19937; + 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::CURAND_RNG_TEST; From 231ea9277be00f6b48026274cdf6df06217f338e Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 22 Nov 2023 17:37:06 -0600 Subject: [PATCH 12/17] Refactoring for curand abstraction replacements. Step 5. --- .../random/randutil/generator_create.inl | 34 +++++++++---------- src/cunumeric/random/rnd_aliases.h | 15 ++++++-- 2 files changed, 28 insertions(+), 21 deletions(-) diff --git a/src/cunumeric/random/randutil/generator_create.inl b/src/cunumeric/random/randutil/generator_create.inl index 2f9cdf29b..d3f6f29e6 100644 --- a/src/cunumeric/random/randutil/generator_create.inl +++ b/src/cunumeric/random/randutil/generator_create.inl @@ -17,10 +17,10 @@ #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); @@ -29,22 +29,20 @@ curandStatus_t randutilGenerator(randutilGenerator_t* generator, } 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); + 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/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h index 81b06bd26..186196bfb 100644 --- a/src/cunumeric/random/rnd_aliases.h +++ b/src/cunumeric/random/rnd_aliases.h @@ -21,9 +21,13 @@ using rnd_status_t = int; enum class randRngType : int { STL_MT_19937 = 1 }; -using randRngType_t = randRngType; -constexpr int RND_STATUS_SUCCESS = 0; -constexpr rnd_status_t RND_STATUS_INTERNAL_ERROR = 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 = randRngType::STL_MT_19937; @@ -34,6 +38,8 @@ 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 @@ -42,6 +48,7 @@ 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 int RND_RNG_PSEUDO_XORWOW = CURAND_RNG_PSEUDO_XORWOW; constexpr int RND_RNG_PSEUDO_PHILOX4_32_10 = CURAND_RNG_PSEUDO_PHILOX4_32_10; @@ -51,4 +58,6 @@ 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 From 9077427305cc9f70fca2fe5144433ccbf0c9b8f9 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 22 Nov 2023 18:52:30 -0600 Subject: [PATCH 13/17] Clean-up of some headers. --- src/cunumeric/random/rnd_aliases.h | 11 +++++++++++ src/cunumeric/random/rnd_types.h | 7 ++----- 2 files changed, 13 insertions(+), 5 deletions(-) diff --git a/src/cunumeric/random/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h index 186196bfb..f6adf5210 100644 --- a/src/cunumeric/random/rnd_aliases.h +++ b/src/cunumeric/random/rnd_aliases.h @@ -16,7 +16,16 @@ #pragma once +// attempt to masquerade as MacOS on host: +// +// #ifndef LEGATE_USE_CUDA +// #define USE_STL_RANDOM_ENGINE_ +// #endif + #ifdef USE_STL_RANDOM_ENGINE_ + +// #pragma message("************ STL path *************") + #include using rnd_status_t = int; @@ -43,6 +52,8 @@ using stream_t = void*; #else #include +// #pragma message("************ CURAND path ************") + using rnd_status_t = curandStatus_t; using randRngType = curandRngType; using randRngType_t = curandRngType_t; diff --git a/src/cunumeric/random/rnd_types.h b/src/cunumeric/random/rnd_types.h index 29403a726..6e1e807e0 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -16,9 +16,10 @@ #pragma once -#ifdef USE_STL_RANDOM_ENGINE_ #include "cunumeric/random/rnd_aliases.h" +#ifdef USE_STL_RANDOM_ENGINE_ + #define CHECK_RND_ENGINE(expr) \ do { \ rnd_status_t __result__ = (expr); \ @@ -53,10 +54,6 @@ static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) #else #include "cunumeric/random/curand_help.h" -// in this case include aliases _after_ curand_help -// -#include "cunumeric/random/rnd_aliases.h" - #define CHECK_RND_ENGINE(expr) CHECK_CURAND((expr)) #define randutil_check_status randutil_check_curand #define get_rndRngType get_curandRngType From c2f75e06017b374152be67d6608cd2305a8d1cbb Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Tue, 28 Nov 2023 13:49:00 -0600 Subject: [PATCH 14/17] Branch-off STL code. Builds, but tests fail, including on device. --- src/cunumeric/random/bitgenerator.cc | 4 + src/cunumeric/random/bitgenerator_curand.inl | 2 +- src/cunumeric/random/randutil/generator.h | 14 +- .../random/randutil/generator_create.inl | 4 +- .../random/randutil/generator_host.cc | 4 + .../randutil/generator_host_advanced.cc | 110 ++++---- .../generator_host_straightforward.cc | 80 +++--- .../random/randutil/generator_uniform.inl | 12 - src/cunumeric/random/randutil/randutil.h | 264 +++++++++--------- src/cunumeric/random/rnd_aliases.h | 25 +- src/cunumeric/random/rnd_types.h | 2 +- 11 files changed, 267 insertions(+), 254 deletions(-) diff --git a/src/cunumeric/random/bitgenerator.cc b/src/cunumeric/random/bitgenerator.cc index c93abf4a7..87cce9d0f 100644 --- a/src/cunumeric/random/bitgenerator.cc +++ b/src/cunumeric/random/bitgenerator.cc @@ -14,6 +14,10 @@ * */ +// attempt to masquerade as MacOS on host: +// +#define USE_STL_RANDOM_ENGINE_ + #include "cunumeric/random/bitgenerator.h" #include "cunumeric/random/bitgenerator_template.inl" #include "cunumeric/random/bitgenerator_util.h" diff --git a/src/cunumeric/random/bitgenerator_curand.inl b/src/cunumeric/random/bitgenerator_curand.inl index 6f4ffa527..fab6fb1f4 100644 --- a/src/cunumeric/random/bitgenerator_curand.inl +++ b/src/cunumeric/random/bitgenerator_curand.inl @@ -41,7 +41,7 @@ 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) diff --git a/src/cunumeric/random/randutil/generator.h b/src/cunumeric/random/randutil/generator.h index 93f42706c..466c32ec0 100644 --- a/src/cunumeric/random/randutil/generator.h +++ b/src/cunumeric/random/randutil/generator.h @@ -41,6 +41,10 @@ template <> 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 = RND_RNG_PSEUDO_PHILOX4_32_10; @@ -49,6 +53,7 @@ template <> struct generatorid { static constexpr int rng_type = RND_RNG_PSEUDO_MRG32K3A; }; +#endif #if 0 // this seems dead code @@ -76,16 +81,17 @@ struct inner_generator : basegenerator uint64_t generatorID; gen_t generator; - inner_generator(uint64_t seed, uint64_t generatorID, cudaStream_t ignored) + inner_generator(uint64_t seed, uint64_t generatorID, stream_t ignored) : seed(seed), generatorID(generatorID) #ifdef USE_STL_RANDOM_ENGINE_ , - generator(seed), - std::srand(seed) + generator(seed) #endif { -#if !defined USE_STL_RANDOM_ENGINE_ +#ifdef USE_STL_RANDOM_ENGINE_ + std::srand(seed); +#else curand_init(seed, generatorID, 0, &generator); #endif } diff --git a/src/cunumeric/random/randutil/generator_create.inl b/src/cunumeric/random/randutil/generator_create.inl index d3f6f29e6..361b2ec32 100644 --- a/src/cunumeric/random/randutil/generator_create.inl +++ b/src/cunumeric/random/randutil/generator_create.inl @@ -25,7 +25,7 @@ rnd_status_t randutilGenerator(randutilGenerator_t* generator, randutilimpl::inner_generator* result = new randutilimpl::inner_generator(seed, generatorID, stream); *generator = (randutilGenerator_t)result; - return CURAND_STATUS_SUCCESS; + return RND_STATUS_SUCCESS; } template @@ -35,7 +35,7 @@ static rnd_status_t inner_randutilCreateGenerator(randutilGenerator_t* generator uint64_t generatorID, stream_t stream = nullptr) { - switch (rng_type) { + switch (static_cast(rng_type)) { case RND_RNG_PSEUDO_XORWOW: return randutilGenerator(generator, seed, generatorID, stream); case RND_RNG_PSEUDO_PHILOX4_32_10: diff --git a/src/cunumeric/random/randutil/generator_host.cc b/src/cunumeric/random/randutil/generator_host.cc index 7d3f456e1..384901e5c 100644 --- a/src/cunumeric/random/randutil/generator_host.cc +++ b/src/cunumeric/random/randutil/generator_host.cc @@ -14,6 +14,10 @@ * */ +// attempt to masquerade as MacOS on host: +// +#define USE_STL_RANDOM_ENGINE_ + #include "generator.h" #include "generator_create.inl" #include "cunumeric/random/rnd_aliases.h" diff --git a/src/cunumeric/random/randutil/generator_host_advanced.cc b/src/cunumeric/random/randutil/generator_host_advanced.cc index 6f836403f..1648d171e 100644 --- a/src/cunumeric/random/randutil/generator_host_advanced.cc +++ b/src/cunumeric/random/randutil/generator_host_advanced.cc @@ -14,13 +14,17 @@ * */ +// attempt to masquerade as MacOS on host: +// +#define USE_STL_RANDOM_ENGINE_ + #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 +34,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 +50,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 +72,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 +101,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 +118,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 +138,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 +164,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 +177,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 +196,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 +207,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 +224,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 +234,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 +250,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 +271,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 +288,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 +305,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 +315,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 +331,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 +347,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..703350ba9 100644 --- a/src/cunumeric/random/randutil/generator_host_straightforward.cc +++ b/src/cunumeric/random/randutil/generator_host_straightforward.cc @@ -14,16 +14,20 @@ * */ +// attempt to masquerade as MacOS on host: +// +#define USE_STL_RANDOM_ENGINE_ + #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 +35,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 +52,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 +62,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 +78,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 +88,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 +104,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 +114,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 +130,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 +140,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 +156,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 +167,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 +184,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 +195,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 +212,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 +222,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 +238,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 +249,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 +266,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 +276,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_uniform.inl b/src/cunumeric/random/randutil/generator_uniform.inl index 99ce04c83..456f82034 100644 --- a/src/cunumeric/random/randutil/generator_uniform.inl +++ b/src/cunumeric/random/randutil/generator_uniform.inl @@ -14,18 +14,6 @@ * */ -#ifndef LEGATE_USE_CUDA -//{- -// assume on MACOS for testing: -// -#define IS_MAC_OS_ 0 -// -#if IS_MAC_OS_ == 1 -#define USE_STL_RANDOM_ENGINE_ -#endif -//-} -#endif - #include "generator.h" #include "randomizer.h" 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/rnd_aliases.h b/src/cunumeric/random/rnd_aliases.h index f6adf5210..f98d41a53 100644 --- a/src/cunumeric/random/rnd_aliases.h +++ b/src/cunumeric/random/rnd_aliases.h @@ -16,20 +16,15 @@ #pragma once -// attempt to masquerade as MacOS on host: -// -// #ifndef LEGATE_USE_CUDA -// #define USE_STL_RANDOM_ENGINE_ -// #endif - #ifdef USE_STL_RANDOM_ENGINE_ // #pragma message("************ STL path *************") +#include #include using rnd_status_t = int; -enum class randRngType : int { STL_MT_19937 = 1 }; +enum class randRngType : int { RND_RNG_TEST = 0, STL_MT_19937 = 1 }; using randRngType_t = randRngType; constexpr int RND_STATUS_SUCCESS = 0; @@ -38,19 +33,23 @@ constexpr int RND_STATUS_SUCCESS = 0; 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 = randRngType::STL_MT_19937; -constexpr int RND_RNG_PSEUDO_PHILOX4_32_10 = randRngType::STL_MT_19937; -constexpr int RND_RNG_PSEUDO_MRG32K3A = randRngType::STL_MT_19937; +// 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 +//} // namespace randutilimpl using stream_t = void*; #else #include +#include // #pragma message("************ CURAND path ************") @@ -61,6 +60,8 @@ 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; diff --git a/src/cunumeric/random/rnd_types.h b/src/cunumeric/random/rnd_types.h index 6e1e807e0..a3063519a 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -46,7 +46,7 @@ static inline randRngType get_rndRngType(cunumeric::BitGeneratorType kind) case cunumeric::BitGeneratorType::PHILOX4_32_10: return randRngType::STL_MT_19937; default: LEGATE_ABORT; } - return randRngType::CURAND_RNG_TEST; + return randRngType::RND_RNG_TEST; } } // namespace cunumeric From 72a5df0061e97bc76cea9f71938fe574a4553120 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 29 Nov 2023 10:26:59 -0600 Subject: [PATCH 15/17] Branch-off STL code. Builds, preliminary tests pass. Device fail. --- src/cunumeric/random/bitgenerator.cu | 15 +++++++++++++-- src/cunumeric/random/curand_help.h | 16 ++++++++++++++++ src/cunumeric/random/rnd_types.h | 2 ++ 3 files changed, 31 insertions(+), 2 deletions(-) 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/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/rnd_types.h b/src/cunumeric/random/rnd_types.h index a3063519a..50ee67355 100644 --- a/src/cunumeric/random/rnd_types.h +++ b/src/cunumeric/random/rnd_types.h @@ -26,6 +26,8 @@ randutil_check_status(__result__, __FILE__, __LINE__); \ } while (false) +// #define randutil_check_curand randutil_check_status + namespace cunumeric { legate::Logger& randutil_log(); From 4a4ca088b773e3fd9d54d5d389729ad5f9f3c867 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Tue, 5 Dec 2023 14:49:12 +0000 Subject: [PATCH 16/17] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- src/cunumeric/random/randutil/generator_pareto.inl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/cunumeric/random/randutil/generator_pareto.inl b/src/cunumeric/random/randutil/generator_pareto.inl index f99c74252..073e957f6 100644 --- a/src/cunumeric/random/randutil/generator_pareto.inl +++ b/src/cunumeric/random/randutil/generator_pareto.inl @@ -41,6 +41,6 @@ struct pareto_t { RANDUTIL_QUALIFIERS double operator()(gen_t& gen) { 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 + return xm * ::exp(-::log(y) * invalpha) - 1.0; // here, use -1.0 to align with numpy } }; From 770de2ff2eeceb237d86975c86e4f0643b7f4201 Mon Sep 17 00:00:00 2001 From: Andrei Schaffer Date: Wed, 20 Dec 2023 10:40:05 -0600 Subject: [PATCH 17/17] MACOS macro guards. --- src/cunumeric/random/bitgenerator.cc | 4 +++- src/cunumeric/random/randutil/generator_host.cc | 4 +++- src/cunumeric/random/randutil/generator_host_advanced.cc | 4 +++- .../random/randutil/generator_host_straightforward.cc | 4 +++- 4 files changed, 12 insertions(+), 4 deletions(-) diff --git a/src/cunumeric/random/bitgenerator.cc b/src/cunumeric/random/bitgenerator.cc index 87cce9d0f..eadbb18a3 100644 --- a/src/cunumeric/random/bitgenerator.cc +++ b/src/cunumeric/random/bitgenerator.cc @@ -14,9 +14,11 @@ * */ -// attempt to masquerade as MacOS on host: +// 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" diff --git a/src/cunumeric/random/randutil/generator_host.cc b/src/cunumeric/random/randutil/generator_host.cc index 384901e5c..93330c1ab 100644 --- a/src/cunumeric/random/randutil/generator_host.cc +++ b/src/cunumeric/random/randutil/generator_host.cc @@ -14,9 +14,11 @@ * */ -// attempt to masquerade as MacOS on host: +// MacOS host variant: // +#if defined(__APPLE__) && defined(__MACH__) #define USE_STL_RANDOM_ENGINE_ +#endif #include "generator.h" #include "generator_create.inl" diff --git a/src/cunumeric/random/randutil/generator_host_advanced.cc b/src/cunumeric/random/randutil/generator_host_advanced.cc index 1648d171e..af3788933 100644 --- a/src/cunumeric/random/randutil/generator_host_advanced.cc +++ b/src/cunumeric/random/randutil/generator_host_advanced.cc @@ -14,9 +14,11 @@ * */ -// attempt to masquerade as MacOS on host: +// MacOS host variant: // +#if defined(__APPLE__) && defined(__MACH__) #define USE_STL_RANDOM_ENGINE_ +#endif #include "generator.h" diff --git a/src/cunumeric/random/randutil/generator_host_straightforward.cc b/src/cunumeric/random/randutil/generator_host_straightforward.cc index 703350ba9..900c666e9 100644 --- a/src/cunumeric/random/randutil/generator_host_straightforward.cc +++ b/src/cunumeric/random/randutil/generator_host_straightforward.cc @@ -14,9 +14,11 @@ * */ -// attempt to masquerade as MacOS on host: +// MacOS host variant: // +#if defined(__APPLE__) && defined(__MACH__) #define USE_STL_RANDOM_ENGINE_ +#endif #include "generator.h"