eic-opticks/CSG/CSGParams.h

0001 #pragma once
0002
0003 /**
0004 CSGParams.h : HMM: needs better name : purpose is for scan testing
0005 ======================================================================
0006
0007 See CSGScan.cc for usage
0008
0009 Aiming to do CUDA only CSG intersect tests that follow
0010 the pattern of OptiX ray trace tests.
0011
0012 CSGOptiX7.cu revolves around the Params constant::
0013
0014     extern "C" { __constant__ Params params ;  }
0015
0016 Try to do something similar here for CUDA CSG scanning.
0017
0018 **/
0019
0020
0021
0022 #if defined(__CUDACC__) || defined(__CUDABE__)
0023    #define PARAMS_METHOD __device__
0024 #else
0025    #define PARAMS_METHOD
0026 #endif
0027
0028
0029
0030 #include "scuda.h"
0031 #include "sqat4.h"
0032 #include "squad.h"
0033
0034 #include "csg_intersect_leaf.h"
0035 #include "csg_intersect_node.h"
0036 #include "csg_intersect_tree.h"
0037
0038
0039 struct CSGParams
0040 {
0041     const CSGNode* node ;
0042     const float4*  plan ;
0043     const qat4*    itra ;
0044
0045
0046     quad4*           qq ;   // "query" rays
0047     quad4*           tt ;   // intersects
0048     int             num ;
0049     bool           devp ;   // device pointers
0050
0051
0052     PARAMS_METHOD void intersect( int idx );
0053
0054 #if defined(__CUDACC__) || defined(__CUDABE__)
0055 #else
0056     PARAMS_METHOD int num_valid_isect();
0057 #endif
0058
0059 };
0060
0061 /**
0062 CSGParams::intersect
0063 ---------------------
0064
0065 Keeping the geometry pointers together with rays and
0066 intersects all together makes this very simple...
0067
0068 Though perhaps better to split off the geometry ?
0069
0070 **/
0071
0072
0073 inline PARAMS_METHOD void CSGParams::intersect( int idx )
0074 {
0075     const quad4* q = qq + idx ;
0076     const float t_min = q->q1.f.w ;
0077     const float3* ori = q->v0();
0078     const float3* dir = q->v1();
0079
0080     float4 isect = make_float4( 0.f, 0.f, 0.f, 0.f ) ;
0081     bool dump = false ;
0082     bool valid_isect = intersect_prim(isect, node, plan, itra, t_min, *ori, *dir, dump );
0083
0084     quad4* t = tt + idx ;
0085
0086     *t = *q ;
0087
0088     t->q0.i.w = int(valid_isect) ;
0089
0090     if( valid_isect )
0091     {
0092         t->q2.f.x  = ori->x + isect.w * dir->x ;
0093         t->q2.f.y  = ori->y + isect.w * dir->y ;
0094         t->q2.f.z  = ori->z + isect.w * dir->z ;
0095
0096         t->q3.f    = isect ;
0097     }
0098 }
0099
0100
0101 #if defined(__CUDACC__) || defined(__CUDABE__)
0102 #else
0103 inline PARAMS_METHOD int CSGParams::num_valid_isect()
0104 {
0105     int n_hit = 0 ;
0106     for(int i=0 ; i < num ; i++)
0107     {
0108         const quad4& t = tt[i] ;
0109         bool hit = t.q0.i.w == 1 ;
0110         if(hit)  n_hit += 1 ;
0111     }
0112     return n_hit ;
0113 }
0114 #endif
0115
0116