eic-opticks/sysrap/sseq.h

0001 #pragma once
0002 /**
0003 sseq.h : photon level step-by-step history and material recording seqhis/seqmat using NSEQ 64 bit uint
0004 ========================================================================================================
0005
0006 Note that needing NSEQ=2 (to allow storing 32 flags) follows from SEventConfig::MaxBounce
0007 being 32 (rather than the long held 16).
0008
0009 Could consider changing NSEQ based on config somehow, but that is not worthy of
0010 the complexity as sseq recording is a debugging activity not a production activity.
0011 So memory usage and performance of sseq is not important, as it will not
0012 be used in production running by definition.
0013
0014 For persisting srec arrays use::
0015
0016    NP* seq = NP<unsigned long long>::Make(num_photon, 2)
0017
0018 **/
0019
0020 #include "scuda.h"
0021
0022 #if defined(__CUDACC__) || defined(__CUDABE__)
0023 #    define SSEQ_METHOD __device__ __forceinline__
0024 #else
0025 #    define SSEQ_METHOD inline
0026 #endif
0027
0028 #if defined(__CUDACC__) || defined(__CUDABE__)
0029 #define FFS(x)   (__ffs(x))
0030 #define FFSLL(x) (__ffsll(x))
0031 #else
0032 #define FFS(x)   (ffs(x))
0033 #define FFSLL(x) (ffsll(x))
0034 #endif
0035
0036 #define rFFS(x)   ( x == 0 ? 0 : 0x1 << (x - 1) )
0037 #define rFFSLL(x) ( x == 0 ? 0 : 0x1ull << (x - 1) )
0038 // use rFFSLL when values of x exceed about 31
0039
0040
0041 #if defined(__CUDACC__) || defined(__CUDABE__)
0042 #else
0043 #include <string>
0044 #include <iostream>
0045 #include <iomanip>
0046 #include <sstream>
0047 #include "smath.h"
0048 #include "OpticksPhoton.hh"
0049 #include "sstr.h"
0050 #endif
0051
0052
0053
0054 struct sseq
0055 {
0056     static constexpr const unsigned NSEQ = 2 ;
0057     static constexpr const unsigned BITS = 4 ;
0058     static constexpr const unsigned long long MASK = ( 0x1ull << BITS ) - 1ull ;
0059     static constexpr const unsigned SLOTMAX = 64/BITS ;     // 16
0060     static constexpr const unsigned SLOTS = SLOTMAX*NSEQ ;  // 32
0061
0062     typedef unsigned long long ULL ;
0063     ULL seqhis[NSEQ] ;
0064     ULL seqbnd[NSEQ] ;
0065
0066     static SSEQ_METHOD unsigned GetNibble(const ULL& seq, unsigned slot) ;
0067     static SSEQ_METHOD void     ClearNibble(    ULL& seq, unsigned slot) ;
0068     static SSEQ_METHOD void     SetNibble(      ULL& seq, unsigned slot, unsigned value ) ;
0069
0070     SSEQ_METHOD unsigned get_flag(unsigned slot) const ;
0071     SSEQ_METHOD void     set_flag(unsigned slot, unsigned flag) ;
0072
0073     SSEQ_METHOD int      seqhis_nibbles() const ;
0074     SSEQ_METHOD int      seqbnd_nibbles() const ;
0075
0076     SSEQ_METHOD void zero();
0077     SSEQ_METHOD void add_nibble( unsigned slot, unsigned flag, unsigned boundary );
0078
0079
0080 #if defined(__CUDACC__) || defined(__CUDABE__)
0081 #else
0082     SSEQ_METHOD std::string desc() const ;
0083     SSEQ_METHOD std::string desc_seqhis() const ;
0084     SSEQ_METHOD std::string brief() const ;
0085     SSEQ_METHOD std::string seqhis_() const ;
0086     SSEQ_METHOD bool operator< (const sseq& other) const ;
0087     SSEQ_METHOD bool operator==(const sseq& other) const ;
0088 #endif
0089 };
0090
0091
0092 SSEQ_METHOD unsigned sseq::GetNibble(const unsigned long long& seq, unsigned slot)
0093 {
0094     return ( seq >> 4*slot ) & 0xfull ;
0095 }
0096 SSEQ_METHOD void sseq::ClearNibble(unsigned long long& seq, unsigned slot)
0097 {
0098     seq &= ~( 0xfull << 4*slot ) ;
0099 }
0100 SSEQ_METHOD void sseq::SetNibble(unsigned long long& seq, unsigned slot, unsigned value)
0101 {
0102     seq =  ( seq & ~( 0xfull << 4*slot )) | ( (value & 0xfull) << 4*slot ) ;
0103 }
0104
0105 /*
0106 SSEQ_METHOD unsigned sseq::get_flag(unsigned slot) const
0107 {
0108     unsigned f = GetNibble(seqhis, slot) ;
0109     return  f == 0 ? 0 : 0x1ull << (f - 1) ;
0110 }
0111 SSEQ_METHOD void sseq::set_flag(unsigned slot, unsigned flag)
0112 {
0113     SetNibble(seqhis, slot, FFS(flag)) ;
0114 }
0115 */
0116
0117 SSEQ_METHOD unsigned sseq::get_flag(unsigned slot) const
0118 {
0119     unsigned iseq = slot/SLOTMAX ;
0120     unsigned f = iseq < NSEQ ? ( seqhis[iseq] >> BITS*(slot - iseq*SLOTMAX) ) & MASK : 0  ;
0121     return  f == 0 ? 0 : 0x1ull << (f - 1) ;   // reconstitute flag from bitpos, inverse of FFS
0122 }
0123 SSEQ_METHOD void sseq::set_flag(unsigned slot, unsigned flag)
0124 {
0125     unsigned iseq = slot/SLOTMAX ;  // iseq:element to write to
0126     if(iseq < NSEQ) seqhis[iseq] |=  (( FFS(flag) & MASK ) << BITS*(slot - iseq*SLOTMAX) );
0127
0128     // NB: note that this does not clear first, so it requires starting with zeroed elements
0129     //
0130     // slot ranges across all elements, so subtracting total slots for preceding elements (iseq*SLOTMAX)
0131     // gives the appropriate shift for the iseq element
0132 }
0133
0134 SSEQ_METHOD int sseq::seqhis_nibbles() const
0135 {
0136     int count = 0 ;
0137     for(unsigned i=0 ; i < NSEQ ; i++) count += smath::count_nibbles(seqhis[i]) ;
0138     return count ;
0139 }
0140 SSEQ_METHOD int sseq::seqbnd_nibbles() const
0141 {
0142     int count = 0 ;
0143     for(unsigned i=0 ; i < NSEQ ; i++) count += smath::count_nibbles(seqbnd[i]) ;
0144     return count ;
0145 }
0146
0147 SSEQ_METHOD void sseq::zero()
0148 {
0149     for(unsigned i=0 ; i < NSEQ ; i++)
0150     {
0151         seqhis[i] = 0ull ;
0152         seqbnd[i] = 0ull ;
0153     }
0154 }
0155
0156 /**
0157 sseq::add_nibble
0158 ------------------
0159
0160 Populates one nibble of the seqhis+seqbnd bitfields
0161
0162 Hmm signing the boundary for each step would eat into bits too much, perhaps
0163 just collect material, as done in old workflow ?
0164
0165 Have observed (see test_desc_seqhis) that when adding more than 16 nibbles
0166 into the 64 bit ULL (which will not fit), get unexpected "mixed" wraparound
0167 not just simply overwriting.
0168
0169 0xfull is needed to avoid all bits above 32 getting set
0170 NB: nibble restriction of each "slot" means there is absolute no need for FFSLL
0171
0172 **/
0173
0174 SSEQ_METHOD void sseq::add_nibble(unsigned slot, unsigned flag, unsigned boundary )
0175 {
0176     unsigned iseq = slot/SLOTMAX ;  // iseq:element to write to, 0 or 1 for NSEQ 2
0177     unsigned shift = BITS*(slot - iseq*SLOTMAX) ;
0178
0179     if(iseq < NSEQ)
0180     {
0181         seqhis[iseq] |=  (( FFS(flag) & MASK ) << shift );
0182         seqbnd[iseq] |=  (( boundary  & MASK ) << shift );
0183     }
0184 }
0185
0186 #if defined(__CUDACC__) || defined(__CUDABE__)
0187 #else
0188
0189 SSEQ_METHOD std::string sseq::desc() const
0190 {
0191     std::stringstream ss ;
0192     ss << " seqhis " ;
0193     for(unsigned i=0 ; i < NSEQ ; i++) ss << " " << std::setw(16) << std::hex << seqhis[NSEQ-1-i] << std::dec  ;
0194     ss << " seqbnd " ;
0195     for(unsigned i=0 ; i < NSEQ ; i++) ss << " " << std::setw(16) << std::hex << seqhis[NSEQ-1-i] << std::dec  ;
0196     std::string s = ss.str();
0197     return s ;
0198 }
0199
0200 SSEQ_METHOD std::string sseq::desc_seqhis() const
0201 {
0202     std::string fseq = OpticksPhoton::FlagSequence_(&seqhis[0], NSEQ) ;  // HMM: bring this within sseq ?
0203
0204     std::stringstream ss ;
0205     for(unsigned i=0 ; i < NSEQ ; i++) ss << " " << std::setw(16) << std::hex << seqhis[NSEQ-1-i] << std::dec  ;
0206     ss << " nib " << std::setw(2) << seqhis_nibbles()
0207        << " " << sstr::TrimTrailing(fseq.c_str())
0208        ;
0209     std::string s = ss.str();
0210     return s ;
0211 }
0212 SSEQ_METHOD std::string sseq::brief() const
0213 {
0214     std::string fseq = OpticksPhoton::FlagSequence_(&seqhis[0], NSEQ) ;  // HMM: bring this within sseq ?
0215     return sstr::TrimTrailing(fseq.c_str()) ;
0216 }
0217
0218 SSEQ_METHOD std::string sseq::seqhis_() const
0219 {
0220     return OpticksPhoton::FlagSequence_(&seqhis[0], NSEQ) ;  // HMM: bring this within sseq ?
0221 }
0222 SSEQ_METHOD bool sseq::operator< (const sseq& other) const
0223 {
0224     return seqhis[1] == other.seqhis[1] ? seqhis[0] < other.seqhis[0] : seqhis[1] < other.seqhis[1] ;
0225 }
0226
0227 /**
0228 sseq::operator==
0229 ------------------
0230
0231 NB only seqhis is used to check equality, not seqbnd
0232
0233 **/
0234
0235 SSEQ_METHOD bool sseq::operator==(const sseq& other) const
0236 {
0237     return seqhis[1] == other.seqhis[1] && seqhis[0] == other.seqhis[0] ;
0238 }
0239
0240
0241
0242
0243 /**
0244 sseq hash specialization allowing sseq to be used as a map key
0245 **/
0246 template<>
0247 struct std::hash<sseq>
0248 {
0249     std::size_t operator()(const sseq& k) const
0250     {
0251         std::size_t h1 = std::hash<uint64_t>()(k.seqhis[1]);
0252         std::size_t h2 = std::hash<uint64_t>()(k.seqhis[0]);
0253         return h1 ^ (h2 << 1);
0254     }
0255 };
0256
0257
0258
0259
0260
0261
0262 #endif
0263
0264