55#if defined(__CUDACC__) || defined(__CUDABE__)
66 #define PRIM_METHOD __device__
77#else
8- #define PRIM_METHOD
9- #endif
8+ #define PRIM_METHOD
9+ #endif
1010
1111#if defined(__CUDACC__) || defined(__CUDABE__)
1212#else
2020CSGPrim : references contiguous sequence of *numNode* CSGNode starting from *nodeOffset* : complete binary tree of 1,3,7,15,... CSGNode
2121=========================================================================================================================================
2222
23- * although CSGPrim is uploaded to GPU by CSGFoundry::upload, instances of CSGPrim at first glance
24- appear not to be needed GPU side because the Binding.h HitGroupData carries the same information.
23+ * although CSGPrim is uploaded to GPU by CSGFoundry::upload, instances of CSGPrim at first glance
24+ appear not to be needed GPU side because the Binding.h HitGroupData carries the same information.
2525
26- * But that is disceptive as the uploaded CSGPrim AABB are essential for GAS construction
26+ * But that is disceptive as the uploaded CSGPrim AABB are essential for GAS construction
2727
2828* vim replace : shift-R
2929
@@ -49,7 +49,7 @@ CSGPrim : references contiguous sequence of *numNode* CSGNode starting from *nod
4949 +----+----------------+----------------+----------------+----------------+-------------------------------------------------+
5050 | | | | | | |
5151 | | | | | | |
52- | q3 | BBMax_y | BBMax_z | | | |
52+ | q3 | BBMax_y | BBMax_z | | globalPrimIdx | |
5353 | | | | | | |
5454 | | | | | | |
5555 | | | | | | |
@@ -58,11 +58,11 @@ CSGPrim : references contiguous sequence of *numNode* CSGNode starting from *nod
5858
5959
6060
61- * as instanced CSGPrim are referenced repeatedly they are of limited use for carrying identity info,
61+ * as instanced CSGPrim are referenced repeatedly they are of limited use for carrying identity info,
6262 to map back to source nidx for example
6363
6464* but the global CSGPrim from the remainder CSGSolid 0 are not repeated, so in that case the
65- CSGPrim is the natural place to carry such mapping identity info
65+ CSGPrim is the natural place to carry such mapping identity info
6666
6767
6868
@@ -71,35 +71,37 @@ CSGPrim : references contiguous sequence of *numNode* CSGNode starting from *nod
7171
7272
7373
74- struct CSG_API CSGPrim
74+ struct CSG_API CSGPrim
7575{
76- quad q0 ;
77- quad q1 ;
78- quad q2 ;
79- quad q3 ;
76+ quad q0 ;
77+ quad q1 ;
78+ quad q2 ;
79+ quad q3 ;
8080
8181 // ---- numNode and nodeOffset are fundamental to the meaning of CSGPrim
8282
83- PRIM_METHOD int numNode () const { return q0.i .x ; }
84- PRIM_METHOD int nodeOffset () const { return q0.i .y ; }
83+ PRIM_METHOD int numNode () const { return q0.i .x ; }
84+ PRIM_METHOD int nodeOffset () const { return q0.i .y ; }
8585 PRIM_METHOD void setNumNode ( int numNode){ q0.i .x = numNode ; }
8686 PRIM_METHOD void setNodeOffset (int nodeOffset){ q0.i .y = nodeOffset ; }
8787
88- // --------- sbtIndexOffset is essential for OptiX 7 SBT PrimSpec machinery, but otherwise not relevant to the geometrical meaning
88+
89+
90+ // --------- sbtIndexOffset is essential for OptiX 7 SBT PrimSpec machinery, but otherwise not relevant to the geometrical meaning
8991
9092 PRIM_METHOD unsigned sbtIndexOffset () const { return q1.u .x ; }
9193 PRIM_METHOD void setSbtIndexOffset (unsigned sbtIndexOffset){ q1.u .x = sbtIndexOffset ; }
9294 PRIM_METHOD const unsigned * sbtIndexOffsetPtr () const { return &q1.u .x ; }
9395
94- // ---------- ARE tran/plan-offset now just metadata, due to absolute referencing ?
96+ // ---------- ARE tran/plan-offset now just metadata, due to absolute referencing ?
9597
96- PRIM_METHOD int tranOffset () const { return q0.i .z ; }
98+ PRIM_METHOD int tranOffset () const { return q0.i .z ; }
9799 PRIM_METHOD int planOffset () const { return q0.i .w ; }
98100
99101 PRIM_METHOD void setTranOffset (int tranOffset){ q0.i .z = tranOffset ; }
100102 PRIM_METHOD void setPlanOffset (int planOffset){ q0.i .w = planOffset ; }
101103
102- // -------- mesh/repeat/primIdx are metadata for debugging convenience
104+ // -------- mesh/repeat/primIdx are metadata for debugging convenience
103105
104106
105107
@@ -109,30 +111,34 @@ struct CSG_API CSGPrim
109111 PRIM_METHOD unsigned repeatIdx () const { return q1.u .z ; } // aka solidIdx/GASIdx
110112 PRIM_METHOD void setRepeatIdx (unsigned ridx){ q1.u .z = ridx ; }
111113
112- PRIM_METHOD unsigned primIdx () const { return q1.u .w ; }
114+ PRIM_METHOD unsigned primIdx () const { return q1.u .w ; }
113115 PRIM_METHOD void setPrimIdx (unsigned pidx){ q1.u .w = pidx ; }
114116
117+ PRIM_METHOD unsigned globalPrimIdx () const { return q3.u .w ; }
118+ PRIM_METHOD void setGlobalPrimIdx (unsigned gpidx){ q3.u .w = gpidx ; }
119+
120+
115121 // ---------- AABB and ce needs changing when transform are applied to the nodes
116122
117- 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 ; }
118- 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 ; }
123+ 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 ; }
124+ 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 ; }
119125 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 ; }
120- 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 ] ; }
126+ 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 ] ; }
121127 PRIM_METHOD const float * AABB () const { return &q2.f .x ; }
122128 PRIM_METHOD float * AABB_ () { return &q2.f .x ; }
123129 PRIM_METHOD const float3 mn () const { return make_float3 (q2.f .x , q2.f .y , q2.f .z ) ; }
124130 PRIM_METHOD const float3 mx () const { return make_float3 (q2.f .w , q3.f .x , q3.f .y ) ; }
125131
126- PRIM_METHOD const float4 ce () const
127- {
128- float x0, y0, z0, x1, y1, z1 ;
129- getAABB (x0, y0, z0, x1, y1, z1);
130- return make_float4 ( (x0 + x1)/2 .f , (y0 + y1)/2 .f , (z0 + z1)/2 .f , extent () );
132+ PRIM_METHOD const float4 ce () const
133+ {
134+ float x0, y0, z0, x1, y1, z1 ;
135+ getAABB (x0, y0, z0, x1, y1, z1);
136+ return make_float4 ( (x0 + x1)/2 .f , (y0 + y1)/2 .f , (z0 + z1)/2 .f , extent () );
131137 }
132- PRIM_METHOD float extent () const
138+ PRIM_METHOD float extent () const
133139 {
134- float3 d = make_float3 ( q2.f .w - q2.f .x , q3.f .x - q2.f .y , q3.f .y - q2.f .z );
135- return fmaxf (fmaxf (d.x , d.y ), d.z ) /2 .f ;
140+ float3 d = make_float3 ( q2.f .w - q2.f .x , q3.f .x - q2.f .y , q3.f .y - q2.f .z );
141+ return fmaxf (fmaxf (d.x , d.y ), d.z ) /2 .f ;
136142 }
137143
138144
@@ -144,43 +150,43 @@ struct CSG_API CSGPrim
144150
145151 void scaleAABB_ ( float scale )
146152 {
147- float * aabb = AABB_ ();
148- for (int i=0 ; i < 6 ; i++ ) *(aabb+i) = *(aabb+i) * scale ;
153+ float * aabb = AABB_ ();
154+ for (int i=0 ; i < 6 ; i++ ) *(aabb+i) = *(aabb+i) * scale ;
149155 }
150156
151157 static void Copy (CSGPrim& b, const CSGPrim& a)
152158 {
153- 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 ;
154- 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 ;
155- 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 ;
156- 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 ;
159+ 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 ;
160+ 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 ;
161+ 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 ;
162+ 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 ;
157163 }
158164
159165 static int value_offsetof_sbtIndexOffset (){ return offsetof (CSGPrim, q1.u .x )/4 ; } // 4
160- static int value_offsetof_AABB (){ return offsetof (CSGPrim, q2.f .x )/4 ; } // 8
166+ static int value_offsetof_AABB (){ return offsetof (CSGPrim, q2.f .x )/4 ; } // 8
161167
162168 static PRIM_METHOD void select_prim_mesh (const std::vector<CSGPrim>& prims, std::vector<CSGPrim>& select_prims, unsigned mesh_idx_ )
163169 {
164170 for (unsigned i=0 ; i < prims.size () ; i++)
165171 {
166- const CSGPrim& pr = prims[i] ;
167- unsigned mesh_idx = pr.meshIdx ();
172+ const CSGPrim& pr = prims[i] ;
173+ unsigned mesh_idx = pr.meshIdx ();
168174 if ( mesh_idx_ == mesh_idx ) select_prims.push_back (pr) ;
169175 }
170176 }
171-
177+
172178 /* *
173179 CSGPrim::select_prim_pointers_mesh
174180 -----------------------------------
175181
176- From the *prims* vector reference find prim with mesh_idx and collect the CSGPrim pointers into *select_prims*
177- **/
182+ From the *prims* vector reference find prim with mesh_idx and collect the CSGPrim pointers into *select_prims*
183+ **/
178184 static PRIM_METHOD void select_prim_pointers_mesh (const std::vector<CSGPrim>& prims, std::vector<const CSGPrim*>& select_prims, unsigned mesh_idx_ )
179185 {
180186 for (unsigned i=0 ; i < prims.size () ; i++)
181187 {
182- const CSGPrim* pr = prims.data () + i ;
183- unsigned mesh_idx = pr->meshIdx ();
188+ const CSGPrim* pr = prims.data () + i ;
189+ unsigned mesh_idx = pr->meshIdx ();
184190 if ( mesh_idx_ == mesh_idx ) select_prims.push_back (pr) ;
185191 }
186192 }
@@ -189,25 +195,25 @@ struct CSG_API CSGPrim
189195 CSGPrim::count_prim_mesh
190196 --------------------------
191197
192- Count the number of prims with meshIdx equal to the queried mesh_idx_
198+ Count the number of prims with meshIdx equal to the queried mesh_idx_
193199
194200 **/
195201 static PRIM_METHOD unsigned count_prim_mesh (const std::vector<CSGPrim>& prims, unsigned mesh_idx_ )
196202 {
197- unsigned count = 0u ;
203+ unsigned count = 0u ;
198204 for (unsigned i=0 ; i < prims.size () ; i++)
199205 {
200- const CSGPrim& pr = prims[i] ;
201- unsigned mesh_idx = pr.meshIdx ();
206+ const CSGPrim& pr = prims[i] ;
207+ unsigned mesh_idx = pr.meshIdx ();
202208 if ( mesh_idx_ == mesh_idx ) count += 1 ;
203209 }
204- return count ;
210+ return count ;
205211 }
206212
207213
208- std::string desc () const ;
214+ std::string desc () const ;
209215 static bool IsDiff ( const CSGPrim& a , const CSGPrim& b );
210- static SCSGPrimSpec MakeSpec ( const CSGPrim* prim0, unsigned primIdx, unsigned numPrim ) ;
216+ static SCSGPrimSpec MakeSpec ( const CSGPrim* prim0, unsigned primIdx, unsigned numPrim ) ;
211217#endif
212218
213219};
0 commit comments