random/data/CuHipRngData.hh

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