eic-opticks/CSG/CSGPrim.h

0001 #pragma once
0002
0003 #include "squad.h"
0004
0005 #if defined(__CUDACC__) || defined(__CUDABE__)
0006    #define PRIM_METHOD __device__
0007 #else
0008    #define PRIM_METHOD
0009 #endif
0010
0011 #if defined(__CUDACC__) || defined(__CUDABE__)
0012 #else
0013 #include "SCSGPrimSpec.h"
0014 struct SMeshGroup ;
0015 #endif
0016
0017 #include "CSG_API_EXPORT.hh"
0018
0019
0020 /**
0021 CSGPrim : references contiguous sequence of *numNode* CSGNode starting from *nodeOffset* : complete binary tree of 1,3,7,15,... CSGNode
0022 =========================================================================================================================================
0023
0024 * although CSGPrim are uploaded to GPU by CSGFoundry::upload, instances of CSGPrim at first glance
0025   appear not to be needed GPU side because the Binding.h HitGroupData carries the same information.
0026
0027 * But that is disceptive as the uploaded CSGPrim AABB are essential for GAS construction
0028
0029 * vim replace : shift-R
0030
0031
0032     +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
0033     | q  |      x         |      y         |     z          |      w         |  notes                                          |
0034     +====+================+================+================+================+=================================================+
0035     |    |  numNode       |  nodeOffset    | tranOffset     | planOffset     |                                                 |
0036     | q0 |                |                |                |                |                                                 |
0037     |    |                |                |                |                |                                                 |
0038     +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
0039     |    | sbtIndexOffset |  meshIdx       | repeatIdx      | primIdx        |                                                 |
0040     |    |                |  (lvid)        |                |                |                                                 |
0041     | q1 |                |  (1,1)         |                |                |                                                 |
0042     |    |                |                |                |                |                                                 |
0043     |    |                |                |                |                |                                                 |
0044     +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
0045     |    |                |                |                |                |                                                 |
0046     |    |                |                |                |                |                                                 |
0047     | q2 |  BBMin_x       |  BBMin_y       |  BBMin_z       |  BBMax_x       |                                                 |
0048     |    |                |                |                |                |                                                 |
0049     |    |                |                |                |                |                                                 |
0050     +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
0051     |    |                |                |                |                |                                                 |
0052     |    |                |                |                |                |                                                 |
0053     | q3 |  BBMax_y       |  BBMax_z       |                |  globalPrimIdx |                                                 |
0054     |    |                |                |                |                |                                                 |
0055     |    |                |                |                |                |                                                 |
0056     |    |                |                |                |                |                                                 |
0057     |    |                |                |                |                |                                                 |
0058     +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
0059
0060
0061
0062 * as instanced CSGPrim are referenced repeatedly they are of limited use for carrying identity info,
0063   to map back to source nidx for example
0064
0065 * but the global CSGPrim from the remainder CSGSolid 0 are not repeated, so in that case the
0066   CSGPrim is the natural place to carry such mapping identity info
0067
0068
0069
0070
0071 **/
0072
0073
0074
0075 struct CSG_API CSGPrim
0076 {
0077     quad q0 ;
0078     quad q1 ;
0079     quad q2 ;
0080     quad q3 ;
0081
0082     // ----  numNode and nodeOffset are fundamental to the meaning of CSGPrim
0083
0084     PRIM_METHOD int  numNode() const    { return q0.i.x ; }
0085     PRIM_METHOD int  nodeOffset() const { return q0.i.y ; }
0086     PRIM_METHOD void setNumNode(   int numNode){    q0.i.x = numNode ; }
0087     PRIM_METHOD void setNodeOffset(int nodeOffset){ q0.i.y = nodeOffset ; }
0088
0089
0090
0091     // --------- sbtIndexOffset is essential for OptiX 7 SBT PrimSpec machinery, but otherwise not relevant to the geometrical meaning
0092
0093     PRIM_METHOD unsigned  sbtIndexOffset()    const { return  q1.u.x ; }
0094     PRIM_METHOD void   setSbtIndexOffset(unsigned sbtIndexOffset){  q1.u.x = sbtIndexOffset ; }
0095     PRIM_METHOD const unsigned* sbtIndexOffsetPtr() const { return &q1.u.x ; }
0096
0097     // ---------- ARE tran/plan-offset now just metadata, due to absolute referencing ?
0098
0099     PRIM_METHOD int  tranOffset() const { return q0.i.z ; }
0100     PRIM_METHOD int  planOffset() const { return q0.i.w ; }
0101
0102     PRIM_METHOD void setTranOffset(int tranOffset){ q0.i.z = tranOffset ; }
0103     PRIM_METHOD void setPlanOffset(int planOffset){ q0.i.w = planOffset ; }
0104
0105     // -------- mesh/repeat/primIdx are metadata for debugging convenience
0106
0107
0108
0109     PRIM_METHOD unsigned  meshIdx() const {           return q1.u.y ; }  // aka lvIdx
0110     PRIM_METHOD void   setMeshIdx(unsigned midx){     q1.u.y = midx ; }
0111
0112     PRIM_METHOD unsigned repeatIdx() const {          return q1.u.z ; }  // aka solidIdx/GASIdx
0113     PRIM_METHOD void   setRepeatIdx(unsigned ridx){   q1.u.z = ridx ; }
0114
0115     PRIM_METHOD unsigned primIdx() const {            return q1.u.w ; }
0116     PRIM_METHOD void   setPrimIdx(unsigned pidx){     q1.u.w = pidx ; }
0117
0118     PRIM_METHOD unsigned globalPrimIdx() const {       return q3.u.w ; }
0119     PRIM_METHOD void setGlobalPrimIdx(unsigned gpidx){ q3.u.w = gpidx ; }
0120
0121
0122     // ----------  AABB and ce needs changing when transform are applied to the nodes
0123
0124     PRIM_METHOD void setAABB(  float e  ){                                                   q2.f.x = -e     ; q2.f.y = -e   ; q2.f.z = -e   ; q2.f.w =  e   ; q3.f.x =  e   ; q3.f.y =  e ; }
0125     PRIM_METHOD void setAABB(  float x0, float  y0, float z0, float x1, float y1, float z1){      q2.f.x = x0    ; q2.f.y = y0   ; q2.f.z = z0   ; q2.f.w = x1   ; q3.f.x = y1   ; q3.f.y = z1 ; }
0126     PRIM_METHOD void getAABB( float& x0, float& y0, float& z0, float& x1, float& y1, float& z1) const { x0 = q2.f.x ; y0 = q2.f.y   ; z0 = q2.f.z   ; x1 = q2.f.w   ; y1 = q3.f.x   ; z1 = q3.f.y ; }
0127     PRIM_METHOD void setAABB( const float* a){                                               q2.f.x = a[0] ; q2.f.y = a[1] ; q2.f.z = a[2] ; q2.f.w = a[3] ; q3.f.x = a[4] ; q3.f.y = a[5] ; }
0128     PRIM_METHOD float zmax() const { return q3.f.y ; }
0129     PRIM_METHOD const float* AABB() const {  return &q2.f.x ; }
0130     PRIM_METHOD       float* AABB_()      {  return &q2.f.x ; }
0131     PRIM_METHOD const float3 mn() const {    return make_float3(q2.f.x, q2.f.y, q2.f.z) ; }
0132     PRIM_METHOD const float3 mx() const {    return make_float3(q2.f.w, q3.f.x, q3.f.y) ; }
0133
0134     PRIM_METHOD static bool AABB_Overlap( const CSGPrim* _a, const CSGPrim* _b )
0135     {
0136         const float* a = _a->AABB();
0137         const float* b = _b->AABB();
0138
0139         return (a[0] <= b[3] && a[3] >= b[0]) && // X overlap
0140                (a[1] <= b[4] && a[4] >= b[1]) && // Y overlap
0141                (a[2] <= b[5] && a[5] >= b[2]);   // Z overlap
0142     }
0143
0144
0145     PRIM_METHOD const float4 ce() const
0146     {
0147         float x0, y0, z0, x1, y1, z1 ;
0148         getAABB(x0, y0, z0, x1, y1, z1);
0149         return make_float4( (x0 + x1)/2.f,  (y0 + y1)/2.f,  (z0 + z1)/2.f,  extent() );
0150     }
0151     PRIM_METHOD float extent() const
0152     {
0153         float3 d = make_float3( q2.f.w - q2.f.x, q3.f.x - q2.f.y, q3.f.y - q2.f.z );
0154         return fmaxf(fmaxf(d.x, d.y), d.z) /2.f ;
0155     }
0156
0157
0158 #if defined(__CUDACC__) || defined(__CUDABE__)
0159 #else
0160
0161
0162     struct zmax_functor {
0163         float operator()(const CSGPrim& pr) const { return pr.zmax() ; }
0164     };
0165
0166     struct extent_functor {
0167         float operator()(const CSGPrim& pr) const { return pr.extent() ; }
0168     };
0169
0170
0171     void scaleAABB_( float scale );
0172     static void Copy(CSGPrim& b, const CSGPrim& a);
0173     static int value_offsetof_sbtIndexOffset(){ return offsetof(CSGPrim, q1.u.x)/4 ; }   // 4
0174     static int value_offsetof_AABB(){           return offsetof(CSGPrim, q2.f.x)/4 ; }   // 8
0175
0176     static void select_prim_mesh(const std::vector<CSGPrim>& prims, std::vector<CSGPrim>& select_prims, unsigned mesh_idx_ );
0177     static void select_prim_pointers_mesh(const std::vector<CSGPrim>& prims, std::vector<const CSGPrim*>& select_prims, unsigned mesh_idx_ );
0178     static unsigned count_prim_mesh(const std::vector<CSGPrim>& prims, unsigned mesh_idx_ );
0179     static std::string Desc(     std::vector<const CSGPrim*>& prim );
0180
0181     static constexpr const char* CSGPrim__DescRange_NAMEONLY = "CSGPrim__DescRange_NAMEONLY" ;
0182     static constexpr const char* CSGPrim__DescRange_NUMPY = "CSGPrim__DescRange_NUMPY" ;
0183     static constexpr const char* CSGPrim__DescRange_EXTENT_DIFF = "CSGPrim__DescRange_EXTENT_DIFF" ;
0184     static constexpr const char* CSGPrim__DescRange_EXTENT_MINI = "CSGPrim__DescRange_EXTENT_MINI" ;
0185
0186     static constexpr const char* CSGPrim__DescRange_CE_ZMIN_ZMAX = "CSGPrim__DescRange_CE_ZMIN_ZMAX" ;
0187     static constexpr const char* CSGPrim__DescRange_CE_YMIN_YMAX = "CSGPrim__DescRange_CE_YMIN_YMAX" ;
0188     static constexpr const char* CSGPrim__DescRange_CE_XMIN_XMAX = "CSGPrim__DescRange_CE_XMIN_XMAX" ;
0189
0190     static std::string DescRange(const CSGPrim* prim, int primOffset, int numPrim, const std::vector<std::string>* soname=nullptr, const SMeshGroup* mg = nullptr );
0191
0192     std::string desc() const ;
0193     std::string descRange() const ;
0194     std::string descRangeNumPy() const ;
0195     static bool IsDiff( const CSGPrim& a , const CSGPrim& b );
0196     static SCSGPrimSpec MakeSpec( const CSGPrim* prim0, unsigned primIdx, unsigned numPrim ) ;
0197 #endif
0198
0199 };
0200
0201 #if defined(__CUDACC__) || defined(__CUDABE__)
0202 #else
0203
0204
0205 inline void CSGPrim::scaleAABB_( float scale )
0206 {
0207     float* aabb = AABB_();
0208     for(int i=0 ; i < 6 ; i++ ) *(aabb+i) = *(aabb+i) * scale ;
0209 }
0210
0211 inline void CSGPrim::Copy(CSGPrim& b, const CSGPrim& a)  // static
0212 {
0213     b.q0.f.x = a.q0.f.x ; b.q0.f.y = a.q0.f.y ; b.q0.f.z = a.q0.f.z ; b.q0.f.w = a.q0.f.w ;
0214     b.q1.f.x = a.q1.f.x ; b.q1.f.y = a.q1.f.y ; b.q1.f.z = a.q1.f.z ; b.q1.f.w = a.q1.f.w ;
0215     b.q2.f.x = a.q2.f.x ; b.q2.f.y = a.q2.f.y ; b.q2.f.z = a.q2.f.z ; b.q2.f.w = a.q2.f.w ;
0216     b.q3.f.x = a.q3.f.x ; b.q3.f.y = a.q3.f.y ; b.q3.f.z = a.q3.f.z ; b.q3.f.w = a.q3.f.w ;
0217 }
0218
0219
0220 inline void CSGPrim::select_prim_mesh(const std::vector<CSGPrim>& prims, std::vector<CSGPrim>& select_prims, unsigned mesh_idx_ )
0221 {
0222     for(unsigned i=0 ; i < prims.size() ; i++)
0223     {
0224         const CSGPrim& pr = prims[i] ;
0225         unsigned mesh_idx = pr.meshIdx();
0226         if( mesh_idx_ == mesh_idx ) select_prims.push_back(pr) ;
0227     }
0228 }
0229
0230 /**
0231 CSGPrim::select_prim_pointers_mesh
0232 -----------------------------------
0233
0234 From the *prims* vector reference find prim with mesh_idx and collect the CSGPrim pointers into *select_prims*
0235 **/
0236
0237 inline void CSGPrim::select_prim_pointers_mesh(const std::vector<CSGPrim>& prims, std::vector<const CSGPrim*>& select_prims, unsigned mesh_idx_ )
0238 {
0239     for(unsigned i=0 ; i < prims.size() ; i++)
0240     {
0241         const CSGPrim* pr = prims.data() + i ;
0242         unsigned mesh_idx = pr->meshIdx();
0243         if( mesh_idx_ == mesh_idx ) select_prims.push_back(pr) ;
0244     }
0245 }
0246
0247 /**
0248 CSGPrim::count_prim_mesh
0249 ---------------------------
0250
0251 Count the number of prims with meshIdx equal to the queried mesh_idx_
0252
0253 **/
0254
0255 inline unsigned CSGPrim::count_prim_mesh(const std::vector<CSGPrim>& prims, unsigned mesh_idx_ )  // static
0256 {
0257     unsigned count = 0u ;
0258     for(unsigned i=0 ; i < prims.size() ; i++)
0259     {
0260         const CSGPrim& pr = prims[i] ;
0261         unsigned mesh_idx = pr.meshIdx();
0262         if( mesh_idx_ == mesh_idx ) count += 1 ;
0263     }
0264     return count ;
0265 }
0266
0267 inline std::string CSGPrim::Desc(std::vector<const CSGPrim*>& prim ) // static
0268 {
0269     std::stringstream ss ;
0270     ss << "[CSGPrim::Desc num_prim " << prim.size() << "\n" ;
0271     for(size_t i=0 ; i < prim.size() ; i++ ) ss << std::setw(4) << i << " : " << prim[i]->desc() << "\n" ;
0272     ss << "]CSGPrim::Desc num_prim " << prim.size() << "\n" ;
0273     std::string str = ss.str() ;
0274     return str ;
0275 }
0276
0277
0278
0279
0280
0281
0282 #endif
0283