File indexing completed on 2026-04-10 07:49:39
0001 #pragma once
0002
0003
0004
0005
0006
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
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
0072
0073
0074
0075
0076
0077
0078
0079
0080
0081
0082
0083
0084
0085
0086
0087
0088
0089
0090
0091
0092
0093
0094
0095
0096
0097
0098
0099
0100
0101
0102
0103 inline QBND_METHOD float4 qbnd::boundary_lookup( float nm, unsigned line, unsigned k )
0104 {
0105
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) ;
0114
0115 unsigned iy = _BOUNDARY_NUM_FLOAT4*line + k ;
0116 float y = (float(iy)+0.5f)/float(ny) ;
0117
0118
0119 float4 props = tex2D<float4>( boundary_tex, x, y );
0120
0121
0122
0123
0124 return props ;
0125 }
0126
0127
0128
0129
0130
0131
0132
0133
0134
0135
0136
0137
0138
0139
0140
0141
0142
0143
0144
0145
0146
0147
0148
0149
0150
0151
0152
0153
0154
0155
0156
0157
0158
0159
0160
0161
0162
0163
0164
0165
0166
0167
0168
0169
0170
0171
0172
0173
0174
0175
0176
0177
0178
0179
0180
0181
0182
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 ;
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);
0196 s.m1group2 = boundary_lookup( wavelength, m1_line, 1);
0197 s.material2 = boundary_lookup( wavelength, m2_line, 0);
0198 s.surface = boundary_lookup( wavelength, su_line, 0);
0199 s.set_material2_group_velocity(boundary_lookup(wavelength, m2_line, 1).x);
0200
0201 s.optical = optical[su_line].u ;
0202
0203 s.index.x = optical[m1_line].u.x ;
0204 s.index.y = optical[m2_line].u.x ;
0205 s.index.z = optical[su_line].u.x ;
0206 s.index.w = 0u ;
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