celeritas/random/CuHipRngData.hh

0001 //----------------------------------*-C++-*----------------------------------//
0002 // Copyright 2021-2024 UT-Battelle, LLC, and other Celeritas developers.
0003 // See the top-level COPYRIGHT file for details.
0004 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
0005 //---------------------------------------------------------------------------//
0006 //! \file celeritas/random/CuHipRngData.hh
0007 //---------------------------------------------------------------------------//
0008 #pragma once
0009
0010 #include <random>
0011
0012 #include "corecel/Config.hh"
0013 #include "corecel/DeviceRuntimeApi.hh"
0014
0015 #include "corecel/Assert.hh"
0016 #include "corecel/Macros.hh"
0017 #include "corecel/Types.hh"
0018 #include "corecel/data/Collection.hh"
0019 #include "corecel/sys/Device.hh"
0020
0021 /*!
0022  * \def CELER_RNG_PREFIX
0023  *
0024  * Add a prefix "hip", "cu", or "mock" to a code token. This is used to
0025  * toggle between the different RNG options.
0026  */
0027 #if CELERITAS_USE_CUDA
0028 // Override an undocumented cuRAND API definition to enable usage in host code.
0029 #    define QUALIFIERS static __forceinline__ __host__ __device__
0030 #    include <curand_kernel.h>
0031
0032 #    define CELER_RNG_PREFIX(TOK) cu##TOK
0033 #elif CELERITAS_USE_HIP
0034 #    if (HIP_VERSION_MAJOR > 5 \
0035          || (HIP_VERSION_MAJOR == 5 && HIP_VERSION_MINOR >= 3))
0036 // Override an undocumented hipRAND API definition to enable usage in host
0037 // code.
0038 #        define QUALIFIERS __forceinline__ __host__ __device__
0039 #    else
0040 // Override an older version of that macro
0041 #        define FQUALIFIERS __forceinline__ __host__ __device__
0042 #    endif
0043 #    pragma clang diagnostic push
0044 // "Disabled inline asm, because the build target does not support it."
0045 #    pragma clang diagnostic ignored "-W#warnings"
0046 // "ignoring return value of function declared with 'nodiscard' attribute"
0047 #    pragma clang diagnostic ignored "-Wunused-result"
0048 #    include <hiprand/hiprand_kernel.h>
0049 #    pragma clang diagnostic pop
0050 #    define CELER_RNG_PREFIX(TOK) hip##TOK
0051 #else
0052 // CuHipRng is invalid
0053 #    include "detail/mockrand.hh"
0054
0055 #    define CELER_RNG_PREFIX(TOK) mock##TOK
0056 #endif
0057
0058 namespace celeritas
0059 {
0060 //---------------------------------------------------------------------------//
0061 #if !CELER_USE_DEVICE
0062 using mockrandState_t = detail::MockRandState;
0063 #endif
0064
0065 //! RNG state type: curandState_t, hiprandState_t, mockrandState_t
0066 using CuHipRngThreadState = CELER_RNG_PREFIX(randState_t);
0067
0068 //---------------------------------------------------------------------------//
0069 /*!
0070  * Properties of the global random number generator.
0071  */
0072 template<Ownership W, MemSpace M>
0073 struct CuHipRngParamsData
0074 {
0075     //// DATA ////
0076
0077     unsigned int seed = 12345;
0078
0079     //// METHODS ////
0080
0081     //! Any settings are valid
0082     explicit CELER_FUNCTION operator bool() const { return true; }
0083
0084     //! Assign from another set of data
0085     template<Ownership W2, MemSpace M2>
0086     CuHipRngParamsData& operator=(CuHipRngParamsData<W2, M2> const& other)
0087     {
0088         seed = other.seed;
0089         return *this;
0090     }
0091 };
0092
0093 //---------------------------------------------------------------------------//
0094 /*!
0095  * Initialize an RNG.
0096  */
0097 struct CuHipRngInitializer
0098 {
0099     ull_int seed{0};
0100     ull_int subsequence{0};
0101     ull_int offset{0};
0102 };
0103
0104 //---------------------------------------------------------------------------//
0105 /*!
0106  * RNG state data.
0107  */
0108 template<Ownership W, MemSpace M>
0109 struct CuHipRngStateData
0110 {
0111     //// TYPES ////
0112
0113     template<class T>
0114     using StateItems = StateCollection<T, W, M>;
0115
0116     //// DATA ////
0117
0118     StateItems<CuHipRngThreadState> rng;
0119
0120     //// METHODS ////
0121
0122     //! True if assigned
0123     explicit CELER_FUNCTION operator bool() const { return !rng.empty(); }
0124
0125     //! State size
0126     CELER_FUNCTION size_type size() const { return rng.size(); }
0127
0128     //! Assign from another set of data
0129     template<Ownership W2, MemSpace M2>
0130     CuHipRngStateData& operator=(CuHipRngStateData<W2, M2>& other)
0131     {
0132         CELER_EXPECT(other);
0133         rng = other.rng;
0134         return *this;
0135     }
0136 };
0137
0138 //---------------------------------------------------------------------------//
0139 /*!
0140  * Resize and initialize with the seed stored in params.
0141  */
0142 template<MemSpace M>
0143 void resize(CuHipRngStateData<Ownership::value, M>* state,
0144             HostCRef<CuHipRngParamsData> const& params,
0145             StreamId stream,
0146             size_type size);
0147
0148 }  // namespace celeritas