Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2026-04-10 07:49:39

0001 #pragma once
0002 /**
0003 qbnd.h
0004 =========
0005 
0006 bnd and optical are closely related so they must be kept together
0007 
0008 **/
0009 
0010 enum { _BOUNDARY_NUM_MATSUR = 4,  _BOUNDARY_NUM_FLOAT4 = 2 };
0011 
0012 enum {
0013     OMAT,
0014     OSUR,
0015     ISUR,
0016     IMAT
0017 };
0018 
0019 
0020 #if defined(__CUDACC__) || defined(__CUDABE__)
0021    #define QBND_METHOD __device__
0022 #else
0023    #define QBND_METHOD
0024 #endif
0025 
0026 struct quad4 ;
0027 struct sstate ;
0028 
0029 
0030 struct qbnd
0031 {
0032     cudaTextureObject_t boundary_tex ;
0033     quad4*              boundary_meta ;
0034     unsigned            boundary_tex_MaterialLine_Water ;
0035     unsigned            boundary_tex_MaterialLine_LS ;
0036     quad*               optical ;
0037 
0038 #if defined(__CUDACC__) || defined(__CUDABE__) || defined( MOCK_TEXTURE) || defined(MOCK_CUDA)
0039     QBND_METHOD float4  boundary_lookup( unsigned ix, unsigned iy );
0040     QBND_METHOD float4  boundary_lookup( float nm, unsigned line, unsigned k );
0041     QBND_METHOD void    fill_state(sstate& s, unsigned boundary, float wavelength, float cosTheta, unsigned long long idx, unsigned long long base_pidx );
0042 #endif
0043 
0044 };
0045 
0046 
0047 
0048 #if defined( MOCK_TEXTURE) || defined(MOCK_CUDA)
0049 #include "stexture.h"
0050 #endif
0051 
0052 
0053 #if defined(__CUDACC__) || defined(__CUDABE__) || defined( MOCK_TEXTURE) || defined(MOCK_CUDA)
0054 /**
0055 qbnd::boundary_lookup ix iy : Low level integer addressing lookup
0056 --------------------------------------------------------------------
0057 
0058 **/
0059 
0060 inline QBND_METHOD float4 qbnd::boundary_lookup( unsigned ix, unsigned iy )
0061 {
0062     const unsigned& nx = boundary_meta->q0.u.x  ;
0063     const unsigned& ny = boundary_meta->q0.u.y  ;
0064     float x = (float(ix)+0.5f)/float(nx) ;
0065     float y = (float(iy)+0.5f)/float(ny) ;
0066     float4 props = tex2D<float4>( boundary_tex, x, y );
0067     return props ;
0068 }
0069 
0070 /**
0071 qbnd::boundary_lookup nm line k
0072 ----------------------------------
0073 
0074 nm:    float wavelength
0075 line:  4*boundary_index + OMAT/OSUR/ISUR/IMAT   (0/1/2/3)
0076 k   :  property group index 0/1
0077 
0078 return float4 props
0079 
0080 boundary_meta is required to configure access to the texture,
0081 it is uploaded by QTex::uploadMeta but requires calls to
0082 
0083 
0084 QTex::init automatically sets these from tex dimensions
0085 
0086    q0.u.x : nx width  (eg wavelength dimension)
0087    q0.u.y : ny hright (eg line dimension)
0088 
0089 QTex::setMetaDomainX::
0090 
0091    q1.f.x : nm0  wavelength minimum in nm
0092    q1.f.y : -
0093    q1.f.z : nms  wavelength step size in nm
0094 
0095 QTex::setMetaDomainY::
0096 
0097    q2.f.x :
0098    q2.f.y :
0099    q2.f.z :
0100 
0101 
0102 **/
0103 inline QBND_METHOD float4 qbnd::boundary_lookup( float nm, unsigned line, unsigned k )
0104 {
0105     //printf("//qbnd.boundary_lookup nm %10.4f line %d k %d boundary_meta %p  \n", nm, line, k, boundary_meta  );
0106 
0107     const unsigned& nx = boundary_meta->q0.u.x  ;
0108     const unsigned& ny = boundary_meta->q0.u.y  ;
0109     const float& nm0 = boundary_meta->q1.f.x ;
0110     const float& nms = boundary_meta->q1.f.z ;
0111 
0112     float fx = (nm - nm0)/nms ;
0113     float x = (fx+0.5f)/float(nx) ;   // ?? +0.5f ??
0114 
0115     unsigned iy = _BOUNDARY_NUM_FLOAT4*line + k ;    // 2*line+k (0/1)
0116     float y = (float(iy)+0.5f)/float(ny) ;
0117 
0118 
0119     float4 props = tex2D<float4>( boundary_tex, x, y );
0120 
0121     // printf("//qbnd.boundary_lookup nm %10.4f nm0 %10.4f nms %10.4f  x %10.4f nx %d ny %d y %10.4f props.x %10.4f %10.4f %10.4f %10.4f  \n",
0122     //     nm, nm0, nms, x, nx, ny, y, props.x, props.y, props.z, props.w );
0123 
0124     return props ;
0125 }
0126 
0127 
0128 /**
0129 qbnd::fill_state
0130 -------------------
0131 
0132 Formerly signed the 1-based boundary, now just keeping separate cosTheta to
0133 orient the use of the boundary so are using 0-based boundary.
0134 
0135 cosTheta < 0.f
0136    photon direction is against the surface normal, ie are entering the shape
0137 
0138    * formerly this corresponded to -ve boundary
0139    * line+OSUR is relevant surface
0140    * line+OMAT is relevant first material
0141 
0142 cosTheta > 0.f
0143    photon direction is with the surface normal, ie are exiting the shape
0144 
0145    * formerly this corresponded to +ve boundary
0146    * line+ISUR is relevant surface
0147    * line+IMAT is relevant first material
0148 
0149 
0150 NB the line is above the details of the payload (ie how many float4 per matsur) it is just::
0151 
0152     boundaryIndex*4  + 0/1/2/3     for OMAT/OSUR/ISUR/IMAT
0153 
0154 
0155 The optical buffer is 4 times the length of the bnd, which allows
0156 convenient access to the material and surface indices starting
0157 from a texture line.
0158 
0159 Optical buffer is populated by sstandard::make_optical
0160 former impl::
0161 
0162     GBndLib::createOpticalBuffer
0163     GBndLib::getOpticalBuf
0164 
0165 Notice that s.optical.x and s.index.z are the same thing.
0166 So half of s.index is extraneous and the m1 index and m2 index
0167 is not much used.
0168 
0169 Also only one element of m1group2 was formerly used. The y slot now carries
0170 material2 group velocity so transmitted segments can carry the correct active
0171 medium timing across boundary crossings.
0172 
0173 
0174 s.optical.x
0175     used to distinguish between : boundary, surface (and in future multifilm)
0176 
0177     * currently contains 1-based surface index with 0 meaning "boundary" and anything else "surface"
0178 
0179     * TODO: encode boundary type enum into the high bits of s.optical.x for three way split
0180       (perhaps use trigger strings like MULTIFILM in the boundary spec to configure)
0181       THIS WILL NEED TO BE DONE AT x4 translation level (and repeated in QBndOptical for
0182       dynamic boundary adding)
0183 
0184 **/
0185 
0186 inline QBND_METHOD void qbnd::fill_state(sstate& s, unsigned boundary, float wavelength, float cosTheta, unsigned long long idx, unsigned long long base_pidx  )
0187 {
0188     const int line = boundary*_BOUNDARY_NUM_MATSUR ;      // now that are not signing boundary use 0-based
0189 
0190     const int m1_line = cosTheta > 0.f ? line + IMAT : line + OMAT ;
0191     const int m2_line = cosTheta > 0.f ? line + OMAT : line + IMAT ;
0192     const int su_line = cosTheta > 0.f ? line + ISUR : line + OSUR ;
0193 
0194 
0195     s.material1 = boundary_lookup( wavelength, m1_line, 0);   // refractive_index, absorption_length, scattering_length, reemission_prob
0196     s.m1group2  = boundary_lookup( wavelength, m1_line, 1);   // x: material1 group_velocity, y: material2 group_velocity, z/w unused
0197     s.material2 = boundary_lookup( wavelength, m2_line, 0);   // refractive_index, (absorption_length, scattering_length, reemission_prob) only m2:refractive index actually used
0198     s.surface   = boundary_lookup( wavelength, su_line, 0);   // detect,         , absorb            , (reflect_specular), reflect_diffuse     [they add to 1. so one not used]
0199     s.set_material2_group_velocity(boundary_lookup(wavelength, m2_line, 1).x);
0200 
0201     s.optical = optical[su_line].u ;   // 1-based-surface-index-0-meaning-boundary/type/finish/value  (type,finish,value not used currently)
0202 
0203     s.index.x = optical[m1_line].u.x ; // m1 index (1-based, see sstandard::make_optical)
0204     s.index.y = optical[m2_line].u.x ; // m2 index (1-based, see sstandard::make_optical)
0205     s.index.z = optical[su_line].u.x ; // su index (1-based, see sstandard::make_optical)
0206     s.index.w = 0u ;                   // avoid undefined memory comparison issues
0207 
0208 #if !defined(PRODUCTION) && defined(DEBUG_PIDX)
0209     if( idx == base_pidx )
0210     {
0211         printf("//qbnd.fill_state idx %7lld boundary %d line %d wavelength %10.4f m1_line %d m2_line %d su_line %d s.optical.x %d  \n", idx, boundary, line, wavelength, m1_line, m2_line, su_line, s.optical.x );
0212         printf("//qbnd.fill_state idx %7lld boundary %d [s.index.x-1](m1_index) %d [s.index.y-1](m2_index) %d [s.index.z-1](su_index) %d \n", idx, boundary, s.index.x-1u, s.index.y-1u, s.index.z-1u );
0213     }
0214 #endif
0215 
0216 
0217 }
0218 
0219 #endif