eic-opticks/qudarap/qrng.h

0001 #pragma once
0002 /**
0003 qrng.h
0004 =======
0005
0006 Despite differences between template specializations regarding *uploaded_states*
0007 all specializations have the same init_with_skipahead signature.
0008 The ctor is only compiled on CPU, as the instance get instanciated on CPU
0009 before being uploaded to GPU.
0010
0011 changes
0012 ~~~~~~~~
0013
0014 * 2024/10/20 : photon_idx event_idx args changed from "unsigned" to "unsigned long long"
0015
0016
0017
0018 init_with_skipahead
0019 ---------------------
0020
0021 1. With XORWOW copy the *photon_idx* element of uploaded_states array
0022    to the curandState reference argument, with others curand_init is cheap,
0023    so is used directly.  For XORWOW the curand_init is done in separate
0024    launches creating the chunked states files with QCurandState.cu
0025
0026    * hence only qrng<XORWOW> template specialization has *uploaded_states* member
0027
0028 2. skipahead the curandState by skipahead_event_offset*event_idx
0029    The offset can be configured using the OPTICKS_EVENT_SKIPAHEAD envvar.
0030
0031    Ideally the offset value should be more than maximum number of random values consumed
0032    in any photon of any event. In practice the number of random values consumed per photon will
0033    have a very long tail, so setting the event skipahead offset value to for example 10000
0034    should prevent any clumping issues from the repeated use of the same randoms in every event.
0035
0036
0037 curand_init(seed, subsequence, offset, &rng)
0038 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
0039
0040 seed
0041    use as simulation constant
0042
0043 subsequence
0044    uses photon_idx
0045    [current usage focusses on this RNG "dimension"]
0046
0047    * TODO: check implications of using different RNG "dimensions" for different purposes
0048
0049
0050 offset
0051    use as simulation constant
0052
0053
0054 **/
0055
0056 #if defined(__CUDACC__) || defined(__CUDABE__)
0057    #define QRNG_METHOD __device__
0058 #else
0059    #define QRNG_METHOD
0060 #endif
0061
0062 #include "srng.h"
0063
0064 using ULL = unsigned long long ;
0065
0066 template<typename T> struct qrng {} ;
0067
0068 template<>
0069 struct qrng<XORWOW>
0070 {
0071     ULL  seed ;
0072     ULL  offset ;
0073     ULL  skipahead_event_offset ;
0074
0075
0076 #if defined(__CUDACC__) || defined(__CUDABE__)
0077     XORWOW*   uploaded_states ;
0078 #else
0079     void*     uploaded_states ;
0080 #endif
0081
0082
0083 #if defined(__CUDACC__) || defined(__CUDABE__)
0084     QRNG_METHOD void init(XORWOW& rng, unsigned long long event_idx, unsigned long long photon_idx )
0085     {
0086         rng = uploaded_states[photon_idx] ;
0087         ULL skipahead_ = skipahead_event_offset*event_idx ;
0088         skipahead( skipahead_, &rng );
0089     }
0090 #else
0091     qrng(ULL seed_, ULL offset_, ULL skipahead_event_offset_ )
0092         :
0093         seed(seed_),
0094         offset(offset_),
0095         skipahead_event_offset(skipahead_event_offset_),
0096         uploaded_states(nullptr)
0097     {
0098     }
0099
0100     void set_uploaded_states( void*  uploaded_states_ )
0101     {
0102         uploaded_states = uploaded_states_ ;
0103     }
0104
0105 #endif
0106 };
0107
0108
0109 /**
0110 qrng<Philox>
0111 -------------
0112
0113 With Philox the curand_init does skipahead and skipahead_sequence advancing ctr.xyzw:
0114
0115 skipahead(offset,&rng)
0116    ctr.xyzw
0117
0118 skipahead_sequence(subsequence,&rng)
0119    ctr.zw
0120
0121 **/
0122
0123 template<>
0124 struct qrng<Philox>
0125 {
0126     ULL  seed ;
0127     ULL  offset ;
0128     ULL  skipahead_event_offset ;
0129
0130 #if defined(__CUDACC__) || defined(__CUDABE__)
0131     QRNG_METHOD void init(Philox& rng, unsigned long long event_idx, unsigned long long photon_idx )
0132     {
0133         ULL subsequence_ = photon_idx ;
0134         curand_init( seed, subsequence_, offset, &rng ) ;
0135         ULL skipahead_ = skipahead_event_offset*event_idx ;
0136         skipahead( skipahead_, &rng );
0137     }
0138 #else
0139     qrng(ULL seed_, ULL offset_, ULL skipahead_event_offset_ )
0140         :
0141         seed(seed_),
0142         offset(offset_),
0143         skipahead_event_offset(skipahead_event_offset_)
0144     {
0145     }
0146     void set_uploaded_states( void* ){}
0147 #endif
0148 };
0149
0150
0151 #ifdef RNG_PHILITEOX
0152 template<>
0153 struct qrng<PhiloxLite>
0154 {
0155     ULL  seed ;
0156     ULL  offset ;
0157     ULL  skipahead_event_offset ;
0158
0159 #if defined(__CUDACC__) || defined(__CUDABE__)
0160     QRNG_METHOD void init(PhiloxLite& rng, unsigned long long event_idx, unsigned long long photon_idx )
0161     {
0162         ULL subsequence_ = photon_idx ;
0163         curand_init( seed, subsequence_, offset, &rng ) ;
0164         ULL skipahead_ = skipahead_event_offset*event_idx ;
0165         skipahead( skipahead_, &rng );
0166     }
0167 #else
0168     qrng(ULL seed_, ULL offset_, ULL skipahead_event_offset_ )
0169         :
0170         seed(seed_),
0171         offset(offset_),
0172         skipahead_event_offset(skipahead_event_offset_)
0173     {
0174     }
0175     void set_uploaded_states( void* ){}
0176 #endif
0177 };
0178 #endif
0179
0180