Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2026-04-09 07:48:55

0001 #pragma once
0002 /**
0003 csg_intersect_node.h
0004 =======================
0005 
0006 This header must be included after csg_intersect_leaf.h and before csg_intersect_tree.h
0007 
0008 Providing more types of multi-union or multi-intersection allows the user to communicate 
0009 more precisely hence a simpler less general algorithm more suited 
0010 to the specific geometry can be applied. 
0011 
0012 Motivation and description in: 
0013 
0014 * https://simoncblyth.bitbucket.io/env/presentation/opticks_20220307_fixed_global_leaf_placement_issue.html 
0015 * https://simoncblyth.github.io/env/presentation/opticks_20220307_fixed_global_leaf_placement_issue.html 
0016 * https://juno.ihep.ac.cn/~blyth/env/presentation/opticks_20220307_fixed_global_leaf_placement_issue.html 
0017 
0018 
0019 CSG_DISCONTIGUOUS (possible alt names: CSG_DISJOINT, CSG_SEPERATED
0020     "multi-union" that the user guarantees to have no overlap between nodes, 
0021     for example a menagerie of different primitives arranged such that there are all entirely disconnected
0022     (useful for modelling "holes" to be subtracted from another shape)
0023 
0024 CSG_CONTIGUOUS (hmm need better name, not CSG_MULTIUNION because the constituents must be primitives)
0025     "multi-union" which is expected to be mostly overlapped but ray paths with "gaps" eg from 
0026     corners or L-shapes are permitted by doing the disjoint disqualification in the intersect alg
0027 
0028 CSG_OVERLAP
0029     "multi-intersection" where user guarantees the compound shape has the topology of a ball with no concaves     
0030     ie the paths of all possible rays crossing the geometry always find two intersects,  
0031     other than tangential edge case
0032 
0033 
0034 distance_node_list
0035 
0036 intersect_node_contiguous
0037 intersect_node_discontiguous
0038 intersect_node_overlap
0039 intersect_node
0040     switch between the above and intersect_leaf from lower level header csg_intersect_leaf.h 
0041 
0042 distance_node
0043     switch between distance_node_list OR distance_leaf from lower level header csg_intersect_leaf.h 
0044 
0045 **/
0046 
0047 
0048 #if defined(__CUDACC__) || defined(__CUDABE__)
0049 #    define INTERSECT_FUNC __forceinline__ __device__
0050 #else
0051 #    define INTERSECT_FUNC inline
0052 #endif
0053 
0054 
0055 
0056 /**
0057 distance_node_list
0058 ----------------------------
0059 
0060 1. get the number of subs from *node* which should be the list header node
0061 2. loop over those subs
0062 
0063 **/
0064 
0065 
0066 INTERSECT_FUNC
0067 float distance_node_list( unsigned typecode, const float3& pos, const CSGNode* node, const CSGNode* root, const float4* plan, const qat4* itra )
0068 {
0069     const unsigned num_sub = node->subNum() ; 
0070     const unsigned offset_sub = node->subOffset(); 
0071 
0072     float sd = typecode == CSG_OVERLAP ? -RT_DEFAULT_MAX : RT_DEFAULT_MAX ; 
0073     for(unsigned isub=0 ; isub < num_sub ; isub++)
0074     {
0075          //const CSGNode* sub_node = node+1u+isub ;  
0076          // TOFIX: the abobe is assuming the sub_node follow the node, which they do not for lists within trees
0077          const CSGNode* sub_node = root+offset_sub+isub ;
0078 #ifdef DEBUG
0079          printf("//distance_node_list num_sub %d offset_sub %d isub %d sub_node.typecode %d sub_node.typecode.name %s\n", num_sub, offset_sub, isub, sub_node->typecode(), CSG::Name(sub_node->typecode())) ;  
0080          assert( sub_node->typecode() >= CSG_LEAF ); 
0081 #endif
0082 
0083  
0084          float sub_sd = distance_leaf( pos, sub_node, plan, itra ); 
0085 
0086          switch(typecode)
0087          {
0088             case CSG_CONTIGUOUS:    sd = fminf( sd, sub_sd );   break ; 
0089             case CSG_DISCONTIGUOUS: sd = fminf( sd, sub_sd );   break ; 
0090             case CSG_OVERLAP:       sd = fmaxf( sd, sub_sd );   break ; 
0091          } 
0092 #ifdef DEBUG
0093         printf("//distance_node_list isub %d sub_sd %10.4f sd %10.4f \n", isub, sub_sd, sd );  
0094 #endif
0095     }
0096     return sd ; 
0097 }
0098 
0099 
0100 /**
0101 intersect_node_contiguous : union of shapes which combine to make single fully connected compound shape 
0102 ------------------------------------------------------------------------------------------------------------
0103 
0104 * shapes that filfil the requirements for using CSG_CONTIGUOUS (multi-union) can avoid tree overheads, 
0105   and thus benefit from simpler and faster intersection 
0106   
0107 
0108 Example of a contiguous union shape composed of constituent boxes::
0109 
0110 
0111                                            +---------------+
0112              E : ENTER                     |               |   
0113              X : EXIT                      |               |   
0114                                            |               |   
0115                              +-------------|...............|-----------+
0116                              |             .               .           |   
0117                              |             .               .           |   
0118                              |             .               .           |   
0119                              |             .................           |   
0120                              |                                         |   
0121                              |                                         |   
0122                       +------........                            .......--------+
0123                       |      .      .                            .     .        |   
0124                       |      .      .                            .  0 -1 - - - [2]  
0125                       |      .      .                            .     .        |   
0126            0 - - - - [1]- - -2      .                     0 - - -1 - - 2 - - - [3]  
0127                       E      E      .                            E     X        |   
0128                       |      .      .                            .     .        |  
0129                       +------........                            .......--------+
0130                              |                                         |   
0131                              |                                         |   
0132                              |                                   0 - - 1   
0133                              |                                         |   
0134                              |             .................           |   
0135                              |             .               .           |   
0136                      0 - - -[1]- - - - - - 2               .           |   
0137                              E             E               .           |   
0138                              +-------------.................-----------+
0139                                            |               |   
0140                                            |               |   
0141                                            |               |   
0142                                            +---------------+
0143 
0144 
0145 
0146 * can tell are outside all constituents (and hence the compound) 
0147   when *ALL* first intersects on constituents are state ENTER 
0148   [this assumes non-complemented constituents]
0149   (this is because EXIT is shielded by the corresponding ENTER) 
0150 
0151   * -> compound intersect is the closest ENTER (that fulfils t_min)
0152 
0153   * HMM: notice that having a t_min that eats into that external intersect
0154     makes the starting point effectively t_min ahead (so outside/inside talk
0155     needs to bear that in mind) 
0156 
0157   * WHAT IS RELEVANT IS THE OUTSIDE/INSIDE OF THE +t_min EFFECTIVE START POINT
0158 
0159 * when *ANY* EXIT states are obtained from first intersects this means 
0160   are inside the CONTIGUOUS compound
0161 
0162   * this is true in general for CONTIGUOUS compounds as are considering intersects 
0163     with all constituents and constituents do not "shield" each other : so will 
0164     always manage to have a first intersect that EXITs when start inside any of them 
0165 
0166   * to find the compound intersect are forced to find all the EXITS 
0167     for all the consituents that found ENTERs for  
0168 
0169 
0170 ALG ISSUE : DOES NOT HONOUR t_min CUTTING : PRESUMABLY DUE TO MIXING DISTANCE AND INTERSECT ?
0171 
0172 * fixied by checking insideness of : ray_origin + t_min*ray_direction  ?
0173 * TODO: check get expected t_min behaviour, that is: cutaway sphere in perspective projection around origin
0174   (correct t_min behaviour is a requirement to participate in CSG trees)
0175 
0176 
0177 PROBLEMS WITH USING DISTANCE FUNCTION 
0178 
0179 Using distance_node_contiguous feels "impure" as it mixes volume-centric with surface-centric approaches
0180 plus it comes with some issues:
0181 
0182 1. kinda duplicitous : as will be running intersect_leaf on all sub_node anyhow
0183 2. using a second way of getting the same info (with another float cut) feels likely to cause edge issues 
0184 3. distance functions are not yet implemented for all leaf 
0185   
0186 ADVANTAGE OF USING DISTANCE FUNCTION 
0187 
0188 Knowing inside_or_surface ahead of time allows:
0189 
0190 1. single loop over leaves (BUT there is a hidden loop inside distance_node_contiguous)
0191 2. avoids storing ENTER/EXIT states and distances for isect  
0192  
0193    * expect this to be is a big GPU performance advantage (more storage less in flight)
0194    * especially advantageous as this shape is targetting CSG boolean abuse solids with large numbers of leaves 
0195    * from future : note that cannot avoid storing t_enter to properly get furthest exit that is not disjoint 
0196      making the case to eliminate the distance function
0197 
0198 THOUGHTS ON ROOT CAUSE
0199 
0200 Need to know all isect are ENTER before can decide that do not need to find the EXITs.
0201 Put another way : if any EXIT encountered, must promote all ENTER into EXIT, 
0202 and doing this requires the ENTER distance. 
0203 If the EXIT is obtained immediately after the ENTER that avoids having to store the ENTER distance. 
0204   
0205 How important is the only ENTER optimization is unclear : would need to measure for specific geometry.
0206 BUT : really want to avoid storing all isect if at all possible. 
0207 
0208 FORGO ALL ENTER "OPTIMIZATION"
0209 
0210 * no need to know inside_or_surface 
0211 * just one loop and promote all ENTER to EXIT would also avoid the need to store all isect 
0212 
0213 
0214 TODO: implement in several different ways and test performance for a variety of shapes
0215 
0216 
0217 ISSUE WHEN USED IN TREE
0218 
0219 Checking this within a CSG tree boolean combination reveals that it is an 
0220 over simplification to just return the furthest exit as it misses intermediate exits
0221 that prevents the CSG tree intersection from working.
0222 
0223 Enters and exits that are within the compound must not raise an intersect 
0224 but when a enter or exit transitions in or out of the compound an intersect must 
0225 be returned even when there is a subsequent one further along. 
0226 
0227 How to do that without lots of distance calls ?
0228 Perhaps a depth counter can distinguish between internal 
0229 and external borders.
0230 
0231 
0232 
0233          +----------+ 
0234          |        B |
0235          |          |
0236     +----|----------|-----+ 
0237     | A  |          |     |
0238     |    |          |     |
0239     | 0 E1         X2    X3         In this simple case [3] is the last exit, but that is not generally the case 
0240     |   +1          0    -1         for more involved shapes. 
0241     |    |          |     |
0242     |    |          |     |         depth counter, starting at 0, add 1 at an Enter, subtract 1 at an Exit 
0243     |    |          |     |         when start inside must end -ve 
0244     |    |          |     |         
0245     |    |          |     |         
0246     |    |    0    X1    X2         
0247     |    |         -1    -2         
0248     |    |          |     |
0249     |    +----------+     |
0250     |                     |
0251     +---------------------+
0252     
0253     
0254 But the alg works by looping over ALL constituents
0255 
0256 * classifying Enter/Exit
0257 * when get an Enter, also get the corresponding Exit
0258 
0259 * there can be subsequnet Enters after an Exit, but that would have to be 
0260   handled by another call to the intersect with advanced tmin : because 
0261   only the next isect with the compound at t > tmin is returned
0262 
0263 
0264 So in the above:
0265 
0266    A: Exit at t=3  (-1)
0267    B: Enter at t=1 (-1+1=0), Exit at t=2 (0-1 = -1) 
0268 
0269 
0270 
0271 Now for A B C
0272 
0273 
0274          +----------+ 
0275          |        B |
0276          |          |
0277     +----|----------|-----+                      +-------------------+
0278     | A  |          |     |                      |                C  |  
0279     |    |          |     |                      |                   | 
0280     | 0 E1         X2   [X3]                    E4                  X5
0281     |   +1          0    -1                      0                  -1
0282     |    |          |     |                      |                   |                        
0283     |    |          |     |                      |                   |
0284     |    |          |     |                      |                   |
0285     |    |    0    X1   [X2]                    E3                  X4     
0286     |    |         -1    -2                     -1                  -2
0287     |    |          |     |                      |                   | 
0288     |    |          |     |                      |                   | 
0289     |    |          | 0 [X1]                    E2                  X3 
0290     |    |          |     -1                     0                   -1
0291     |    |          |     |                      |                   | 
0292     |    +----------+     |                      |                   |
0293     |                     |                      |                   |
0294     +---------------------+                      +-------------------+
0295  
0296                                                             
0297 
0298 For rays starting in A or B need to disqualify C intersects because 
0299 the C enters are at greater t than the A and B exits.
0300 
0301 If C expanded to the left to join with A and B then that would not 
0302 be the case and its exits would be allowed.  
0303 
0304 Any of the A,B,C can be disjoint and need to detect and disqualify.
0305 
0306 HMM: should any constituent with an Enter exceeding 
0307 
0308 HMM: when are inside are in need of an exit, but if a constituent
0309 gives you an enter it is under suspicion of needing disqualification.
0310 
0311 HMM: need to distinguish first-exits from second-exits (obtained after an enter) 
0312 
0313 The farthest first-exit is highly likely to be THE ONE, but 
0314 not inevitably.  So on first loop collect the farthest first exit 
0315 
0316 Perhaps can disqualify constituents with Enter that is 
0317 further away that the farthest first-exit. 
0318 
0319 
0320 What about a chain, perhaps repeated pushing 
0321 of envelope prevents the two pass approach ?
0322 
0323 
0324 
0325                  +----------------+     +-------------------+ 
0326                  |B               |     |D                  |
0327                  |                |     |                   |
0328                  |                |     |                   |
0329             +----|----+      +----|-----|----+       +------|----------+             +-----------+
0330             |A   |    |      |C   |     |    |       |E     |          |             |           | 
0331             |    |    |      |    |     |    |       |      |          |             |           |
0332             | 0 E1   (X2)    E3  X4    E5   X6      E7     X8         X9            E10         X10
0333             |    |    |      |    |     |    |       |      |          |             |           | 
0334             |    |    |      |    |     |    |       |      |          |             |           |
0335             |    |    |      |    |     |    |       |      |          |             |           | 
0336             +----|----+      +----|-----|----+       +------|----------+             +-----------+
0337                  |                |     |                   |
0338                  |                |     |                   |
0339                  |                |     |                   |
0340                  +----------------+     +-------------------+
0341 
0342                  E           E          E            E           
0343                       X           X          X              X          X       
0344 
0345 
0346 
0347    
0348 1. *first pass* : find furthest first exit 
0349 
0350    A: X2 : this is the furthest (and only) first exit 
0351    B: E1
0352    C: E3
0353    D: E5
0354    E: E7
0355 
0356 2. *second pass* : redo the undone that are all State_Enter, looping subs with enter less than furthest exit 
0357    
0358    * only E1 qualifies, looping that takes you to X4 
0359 
0360 
0361 
0362 
0363 
0364 
0365 1. *zeroth pass* : hoping that are outside just find nearest enter and count first exits 
0366 
0367 2. when no Exits are outside the compound which makes the intersect simply the closest enter, so return it  
0368 
0369 3. *first pass* : collect enter distances, isub indices and get farthest_exit of the first exits 
0370 
0371    * first exits always qualify as potential intersects, it is only exits after an enter that may be disjoint 
0372      and require contiguity checking before can qualify as candidate intersect
0373    
0374 
0375 
0376 SORTING NETWORKS, BITONIC SORT REFS
0377 
0378 Need to sort a small number of distances
0379 
0380 * https://en.wikipedia.org/wiki/Sorting_network
0381 * https://en.wikipedia.org/wiki/Bitonic_sorter
0382 * https://www.cs.kent.edu/~batcher/sort.pdf
0383 * ~/opticks_refs/batcher_bitonic_sort.pdf
0384  
0385 
0386 
0387 **/
0388 
0389 
0390 #ifdef WITH_CONTIGUOUS
0391 
0392 INTERSECT_FUNC
0393 void intersect_node_contiguous( bool& valid_isect, float4& isect, const CSGNode* node, const CSGNode* root, 
0394        const float4* plan, const qat4* itra, const float t_min , const float3& ray_origin, const float3& ray_direction, bool dumpxyz )
0395 {
0396     const int num_sub = node->subNum() ; 
0397     const int offset_sub = node->subOffset() ; 
0398 #ifdef DEBUG
0399      printf("//intersect_node_contiguous num_sub %d offset_sub %d \n", num_sub, offset_sub ); 
0400 #endif
0401 
0402 #ifdef DEBUG_DISTANCE
0403     float sd = distance_node_list( CSG_CONTIGUOUS, ray_origin + t_min*ray_direction, node, root, plan, itra ); 
0404     bool inside_or_surface = sd <= 0.f ;
0405     printf("//intersect_node_contiguous sd %10.4f inside_or_surface %d num_sub %d offset_sub %d \n", sd, inside_or_surface, num_sub, offset_sub ); 
0406 #endif
0407 
0408     float4 nearest_enter = make_float4( 0.f, 0.f, 0.f, RT_DEFAULT_MAX ) ; 
0409     float4 farthest_exit = make_float4( 0.f, 0.f, 0.f, t_min ) ; 
0410     // TODO: with split zeroth and first passes do not need both these isect at the same time, so could combine/reuse
0411 
0412     float4 sub_isect = make_float4( 0.f, 0.f, 0.f, 0.f ) ;    
0413 
0414     int exit_count = 0 ; 
0415     const float propagate_epsilon = 0.0001f ; 
0416     IntersectionState_t sub_state = State_Miss ;  
0417 
0418     // 1. *zeroth pass* : hoping that are outside just find nearest enter and count exits 
0419 
0420     for(int i=0 ; i < num_sub ; i++)
0421     {
0422         const CSGNode* sub_node = root+offset_sub+i ; 
0423 
0424         bool sub_isect_valid = false ; 
0425         intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, t_min, ray_origin, ray_direction, dumpxyz );
0426         if(sub_isect_valid)
0427         {
0428             sub_state = CSG_CLASSIFY( sub_isect, ray_direction, t_min ); 
0429 #ifdef DEBUG
0430             printf("//intersect_node_contiguous i %d state %s t %10.4f \n", i, IntersectionState::Name(sub_state), sub_isect.w );   
0431 #endif          
0432             if( sub_state == State_Enter)
0433             {
0434                 if( sub_isect.w < nearest_enter.w ) nearest_enter = sub_isect ;  
0435             }
0436             else if( sub_state == State_Exit )
0437             {
0438                 exit_count += 1 ; 
0439             }
0440         }
0441     } 
0442 
0443 #ifdef DEBUG
0444     printf("//intersect_node_contiguous exit_count %d \n", exit_count); 
0445 #endif
0446 
0447     // 2. when no Exits are outside the compound which makes the intersect simply the closest enter, so return it
0448     //
0449     // TODO: the below looks like it breaks the t_min scan-ability requirement : need to check t_min earlier ? 
0450     // HMM: actually it might well be OK as intersect_leaf is doing the t_min checking already, so the 
0451     //      the too late t_min check here is doing nothing other than being mis-leading
0452     //
0453     if(exit_count == 0) 
0454     {
0455         valid_isect = nearest_enter.w > t_min && nearest_enter.w < RT_DEFAULT_MAX ; 
0456         if(valid_isect) isect = nearest_enter ;
0457 #ifdef DEBUG
0458         printf("//intersect_node_contiguous : outside early exit   %10.4f %10.4f %10.4f %10.4f \n", isect.x, isect.y, isect.z, isect.w );  
0459 #endif
0460         return ;   
0461     }
0462 
0463 
0464     // now to work, there are Exits so are inside the compund and need to get outta here using some resources
0465     // hmm but i guess these static resources will also be in play even with the early exit ?
0466     // TODO: compare performance between combining and splitting the zeroth and first passes 
0467 
0468 
0469     int enter_count = 0 ; 
0470     float enter[8] ; 
0471     int   aux[8] ; 
0472     int   idx[8] ;   // HMM: could squeeze 16 nibbles into the 64 bits 
0473 
0474     
0475     
0476     // 3. *first pass* : collect enter distances, isub indices and get farthest_exit of the first exits 
0477     // 
0478     // * first exits always qualify as potential intersects, it is only exits after an enter that may be disjoint 
0479     //   and require contiguity checking before can qualify as candidate intersect
0480     // 
0481 
0482     for(int isub=0 ; isub < num_sub ; isub++)
0483     {
0484         const CSGNode* sub_node = root+offset_sub+isub ; 
0485 
0486         bool sub_isect_valid = false ; 
0487         intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, t_min, ray_origin, ray_direction, dumpxyz); 
0488         if(sub_isect_valid)
0489         {
0490             sub_state = CSG_CLASSIFY( sub_isect, ray_direction, t_min ); 
0491             if( sub_state == State_Enter)
0492             {
0493                 aux[enter_count] = isub ;   // record which isub this enter corresponds to : could instead store the sub_node
0494                 idx[enter_count] = enter_count ; 
0495                 enter[enter_count] = sub_isect.w ;
0496 
0497                 // TODO: very wasteful enter_count and isub usually very small values, 
0498                 //       could pack them together perhaps ?
0499                 //
0500                 // HMM: could have fixed resources for a shape using a num_sub array 
0501                 //
0502                 // HMM: trying to do two levels of indirection at once doesnt work with the indirect sort, hence have to use the aux
0503 
0504 #ifdef DEBUG
0505                 printf("//intersect_node_contiguous isub %d enter_count %d idx[enter_count] %d enter[enter_count] %10.4f \n", isub, enter_count, idx[enter_count], enter[enter_count] ); 
0506 #endif
0507                 enter_count += 1 ; 
0508             }
0509             else if( sub_state == State_Exit )  
0510             {
0511                 exit_count += 1 ; 
0512                 if( sub_isect.w > farthest_exit.w ) farthest_exit = sub_isect ;  
0513             }
0514         }
0515     } 
0516 
0517 #ifdef DEBUG
0518     if(enter_count > 8)
0519     { 
0520         // maybe could use integer template specialization to tailor the limit to each geometry combined 
0521         // with nvrtc runtime-compilation to allow resource use customization during geometry conversion
0522         printf("//intersect_node_contiguous enter_count %d exceeds limit of 8 \n", enter_count );  
0523     }
0524     assert( enter_count <= 8 ) ; 
0525 #endif
0526 
0527 
0528     // insertionSortIndirectSentinel : 
0529     // 4. order the enter indices so that they would make enter ascend 
0530     // see sysrap/tests/sorting/insertionSortIndirect.sh to understand the sort 
0531     // ordering the idx (isub indices) to make the enter values ascend 
0532 
0533 #ifdef DEBUG
0534     printf("//intersect_node_contiguous enter_count %d  (before sort) \n", enter_count ); 
0535     for(int i=0 ; i < enter_count ; i++) printf("//intersect_node_contiguous  i %2d idx[i] %2d  aux[i] %2d enter[i] %10.4f \n", i, idx[i], aux[i], enter[i] ); 
0536 #endif
0537 
0538     for (int i = 1; i < enter_count ; i++)
0539     {   
0540         int key = idx[i] ;  // hold idx[1] out of the pack    
0541         int j = i - 1 ;
0542         
0543         // descending j below i whilst find out of order  
0544         while( j >= 0 && enter[idx[j]] > enter[key] )   
0545         {   
0546             idx[j+1] = idx[j] ;    // i=1,j=0,idx[1]=idx[0]   assuming not ascending
0547             
0548             j = j - 1;
0549             
0550             // sliding values (actually the isub indices) that are greater than the key one upwards
0551             // no need to "swap" as are holding the key out of the pack
0552             // ready to place it into the slot opened by the slide up   
0553         }
0554         
0555         idx[j+1] = key ;       // i=1,j=-1, idx[0]=key
0556 
0557         // i=1, j->0, when enter[j] <= enter[i]  (already ascending) 
0558         // the while block doesnt run 
0559         //   => pointless idx[1] = idx[1]   
0560         // so puts the key back in the pack at the same place it came from  
0561     }
0562 
0563 
0564 #ifdef DEBUG
0565     printf("//intersect_node_contiguous enter_count %d  (after sort) \n", enter_count ); 
0566     for(int i=0 ; i < enter_count ; i++) printf("//intersect_node_contiguous  i %2d idx[i] %2d  aux[i] %2d enter[i] %10.4f enter[idx[i]] %10.4f  \n", i, idx[i], aux[i], enter[i], enter[idx[i]] ); 
0567 #endif
0568 
0569  
0570     // *2nd pass* : loop over enters in t order  
0571     // only find exits for Enters that qualify as contiguous (thus excluding disjoint subs)
0572     // where the qualification is based on the farthest_exit.w, which is updated 
0573     // as each exit is found. 
0574     //
0575     // I think that because the enters are t ordered there is no need for rerunning this
0576     // despite each turn of the loop potentially pushing the envelope of permissible enters
0577     //
0578     // suspect can break on first disqualification ? think line of disjoint boxes
0579     //
0580 
0581     for(int i=0 ; i < enter_count ; i++)
0582     {
0583         // idx[i]       : enter index with ascending t_enter
0584         // aux[idx[i]]  : isub index with ascending t_enter      
0585 
0586         //float tminAdvanced = enter[i] + propagate_epsilon ;    // <= without ordered enters get internal spurious 
0587         float tminAdvanced = enter[idx[i]] + propagate_epsilon ; 
0588 
0589         int isub = aux[idx[i]];  // rather than shuffling aux, are just the mapping from enter index to isub index 
0590         const CSGNode* sub_node = root+offset_sub+isub ; 
0591 
0592 #ifdef DEBUG
0593         printf("//intersect_node_contiguous i/enter_count/isub %d/%d/%d tminAdvanced %10.4f farthest_exit.w %10.4f  \n", i, enter_count, isub, tminAdvanced, farthest_exit.w ); 
0594 #endif
0595 
0596         // note that discontiguous constituents will not enter into farthest_exit 
0597         // (should call "farthest_contiguous_exit")
0598 
0599         if(tminAdvanced < farthest_exit.w)  // enter[idx[i]]+propagate_epsilon < "farthest_contiguous_exit" so far    
0600         {  
0601             bool sub_isect_valid = false ; 
0602             intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, tminAdvanced , ray_origin, ray_direction, dumpxyz);
0603             if(sub_isect_valid)
0604             {
0605                 sub_state = CSG_CLASSIFY( sub_isect, ray_direction, tminAdvanced ); 
0606                 if( sub_state == State_Exit ) 
0607                 {
0608                     exit_count += 1 ; 
0609                     if( sub_isect.w > farthest_exit.w ) farthest_exit = sub_isect ;  
0610                 }
0611 #ifdef DEBUG
0612                
0613                 assert( sub_state == State_Exit ); 
0614 #endif          
0615             }
0616         }
0617     }
0618 
0619     valid_isect = exit_count > 0 && farthest_exit.w > t_min ; 
0620     if(valid_isect) isect = farthest_exit ;  
0621 #ifdef DEBUG
0622      printf("//intersect_node_contiguous valid_isect %d  (%10.4f %10.4f %10.4f %10.4f) \n", valid_isect, isect.x, isect.y, isect.z, isect.w ); 
0623 #endif
0624 }
0625 
0626 #endif
0627 
0628 
0629 /**
0630 intersect_node_discontiguous : union of disjoint leaves with absolutely no overlapping
0631 ----------------------------------------------------------------------------------------
0632 
0633 The guarantee that all sub-nodes do not overlap other sub-nodes
0634 makes the implementation straightforward and hence fast:
0635 
0636 * closest ENTER or EXIT 
0637 
0638 
0639 
0640      +-------+          +-------+          +-------+          +-------+         +-------+      
0641      |       |          |       |          |       |          |       |         |       |      
0642      |       |          |       |          |       |          |       |         |       |      
0643      +-------+          +-------+          +-------+          +-------+         +-------+       
0644 
0645      +-------+          +-------+          +-------+          +-------+         +-------+      
0646      |       |          |       |          |       |          |       |         |       |      
0647      |       |          |       |          |       |          |       |         |       |      
0648      +-------+          +-------+          +-------+          +-------+         +-------+       
0649 
0650 
0651 **/
0652 
0653 
0654 INTERSECT_FUNC
0655 void intersect_node_discontiguous(bool& valid_isect,  float4& isect, const CSGNode* node, const CSGNode* root, 
0656      const float4* plan, const qat4* itra, const float t_min , const float3& ray_origin, const float3& ray_direction, bool dumpxyz )
0657 {
0658     const unsigned num_sub = node->subNum() ; 
0659     const unsigned offset_sub = node->subOffset() ; 
0660 
0661     float4 closest = make_float4( 0.f, 0.f, 0.f, RT_DEFAULT_MAX ) ; 
0662     float4 sub_isect = make_float4( 0.f, 0.f, 0.f, 0.f ) ;    
0663 
0664     for(unsigned isub=0 ; isub < num_sub ; isub++)
0665     {
0666         const CSGNode* sub_node = root+offset_sub+isub ; 
0667         bool sub_isect_valid(false); 
0668         intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, t_min, ray_origin, ray_direction, dumpxyz);
0669         if(sub_isect_valid)
0670         {
0671             if( sub_isect.w < closest.w ) closest = sub_isect ;  
0672         }
0673     }
0674 
0675     valid_isect = closest.w < RT_DEFAULT_MAX ; 
0676     if(valid_isect) 
0677     {
0678         isect = closest ;  
0679     }
0680 
0681 #ifdef DEBUG
0682     printf("//intersect_node_discontiguous num_sub %d  closest.w %10.4f \n", 
0683        num_sub, closest.w ); 
0684 #endif
0685 
0686 }
0687  
0688 
0689 
0690 
0691 
0692 
0693 
0694 /**
0695 intersect_node_overlap : intersection of leaves
0696 ----------------------------------------------------
0697 
0698 * HMM: could general sphere be implemented using CSG_OVERLAP multi-intersection 
0699   of inner and outer sphere, phi planes and theta cones ?
0700 
0701 * shapes that can be described entirely with intersection can avoid tree overheads 
0702 
0703 Imagine the CSG_OVERLAP of 3 shapes A,B,C the resulting ABC shape is the portion that is inside all of them.
0704 
0705 Algorithm loops over constituent sub nodes classifying intersects as ENTER/EXIT/MISS whilst 
0706 updating : nearest_exit, farthest_enter. 
0707 
0708 * from outside the compound ABC shape the desired intersect is the deepest enter ENTER
0709   which you can find by shooting A,B,C and 
0710 
0711 
0712                                                        1X,2M -> MISS
0713 
0714                                         +----------------1X-------------+
0715                                         |               /             C |
0716                                         |              /                |
0717                                         |             0                 |
0718                                         |                               |
0719                               +---------------------------------+       |  
0720                               |         |                     B |       |  
0721                               |         |                       |       |  
0722                    +----------------------------+         0- - -1 - - - 2     2X,1M -> MISS
0723                    | A        |         | . ABC |               X       X  
0724                    |          |         | . . . |               |       |  
0725                    |          |         | . . . |               |       |  
0726           0 - - - -1- - - - - 2- - - - [3]- - - 4               |       |  
0727                    E          E         E . . . X               |       |  
0728                    |          |         | . . . |               |       |  
0729                    |          |         | . . . |               |       |  
0730                    |          |         | .0- -[1]- - - - - - - 2 - - - 3 - - 
0731                    |          |         | . . . X               X       X  
0732                    |          |         | . . . |               |       |  
0733                    |          |         +-------------------------------+  
0734                    |          |                 |               |          
0735                    |          |                 |               |          
0736              0 - - 1 - - - - -2 - - - - - - - - 3 - - - - - - - 4          
0737                    E          E   (2E,1M)-> M   X               X   (2X,1M) -> M
0738                    |          |                 |               |          
0739                    |          +---------------------------------+          
0740                    |                            |                          
0741                    |                            |                                
0742                    +----------------------------+                          
0743              
0744 
0745 
0746 
0747 The below arrangements of constituents are not permissable as there is no common overlap
0748 giving such a shape to a CSG_OVERLAP may abort and will give incorrect or no intersects.      
0749 
0750 
0751              +------+  
0752          +---|------|-----+
0753          |   |      |     |
0754          |   |      |     |
0755          |   +------+     | 
0756          |           +----|-----+
0757          +-----------|----+     |
0758                      +----------+
0759 
0760       +----+      +-----+      +-----+
0761       |    |      |     |      |     |
0762    - -E - - - - - E- - - - - - E- - - - - - 
0763       |    |      |     |      |     |
0764       +----+      +-----+      +-----+
0765 
0766 
0767 
0768 If eed to compare the farthest enter with the nearest enter 
0769 if nearest enter is closer than farthest exit then its a HIT ?
0770 
0771 
0772 
0773 Thinking of ROTH diagrams for intersection
0774 
0775                            
0776                        |--------------|
0777                        
0778                             |------------------|
0779 
0780                                   |-----------------|
0781 
0782            Enters:     |    |     | 
0783 
0784            Exits:                     |        |    |
0785 
0786 
0787                   farthest enter  |   |
0788                                   |   |    nearest exit
0789 
0790 For an overlap need::
0791 
0792                       farthest_enter < nearest_exit  
0793 
0794 
0795 For disjoint::
0796 
0797 
0798                           |-----------|     
0799 
0800                                            |------------|
0801   
0802             Enters:       |                |
0803 
0804             Exits:                    |                 |              
0805                          
0806   
0807                          nearest exit |    | farthest enter 
0808 
0809 No overlap as::
0810 
0811                          nearest_exit   < farthest_enter 
0812 
0813 
0814 
0815 
0816 **/
0817 
0818 
0819 INTERSECT_FUNC
0820 void intersect_node_overlap( bool& valid_isect, float4& isect, const CSGNode* node, const CSGNode* root, 
0821         const float4* plan, const qat4* itra, const float t_min , const float3& ray_origin, const float3& ray_direction, bool dumpxyz)
0822 {
0823     const unsigned num_sub = node->subNum() ; 
0824     const unsigned offset_sub = node->subOffset() ; 
0825 
0826     float4 farthest_enter = make_float4( 0.f, 0.f, 0.f, t_min ) ; 
0827     float4 nearest_exit  = make_float4( 0.f, 0.f, 0.f, RT_DEFAULT_MAX ) ; 
0828     float4 sub_isect = make_float4( 0.f, 0.f, 0.f, 0.f ) ;    
0829 
0830     float propagate_epsilon = 0.0001f ; 
0831     unsigned enter_count = 0 ; 
0832     unsigned exit_count = 0 ; 
0833     IntersectionState_t sub_state = State_Miss ; 
0834 
0835     for(unsigned isub=0 ; isub < num_sub ; isub++)
0836     {
0837         const CSGNode* sub_node = root+offset_sub+isub ; 
0838 
0839         bool sub_isect_valid(false);   
0840         intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, t_min, ray_origin, ray_direction, dumpxyz );
0841         if(sub_isect_valid)
0842         {
0843             sub_state = CSG_CLASSIFY( sub_isect, ray_direction, t_min ); 
0844             if(sub_state == State_Enter)
0845             {
0846                 enter_count += 1 ;  
0847                 if( sub_isect.w > farthest_enter.w ) farthest_enter = sub_isect ;  
0848 
0849                 float tminAdvanced = sub_isect.w + propagate_epsilon ; 
0850 
0851                 intersect_leaf( sub_isect_valid, sub_isect, sub_node, plan, itra, tminAdvanced , ray_origin, ray_direction, dumpxyz);
0852                 if(sub_isect_valid)
0853                 {
0854                     sub_state = CSG_CLASSIFY( sub_isect, ray_direction, tminAdvanced ); 
0855                     if( sub_state == State_Exit ) 
0856                     {
0857                         exit_count += 1 ;  
0858                         if( sub_isect.w < nearest_exit.w ) nearest_exit = sub_isect ;  
0859                     } 
0860                 }
0861             }
0862             else if(sub_state == State_Exit)
0863             {
0864                 exit_count += 1 ; 
0865                 if( sub_isect.w < nearest_exit.w ) nearest_exit = sub_isect ;  
0866             }
0867         } 
0868     }
0869 
0870 
0871     // if no Enter encountered farthest_enter.w stays t_min
0872     // if no Exit encountered nearest_exit.w  stays RT_DEFAULT_MAX 
0873 
0874     valid_isect = false ;  
0875     bool overlap_all = farthest_enter.w < nearest_exit.w && max(enter_count, exit_count) == num_sub ; 
0876     if(overlap_all)  
0877     {
0878         if( farthest_enter.w > t_min && farthest_enter.w  < RT_DEFAULT_MAX )
0879         {
0880             valid_isect = true ; 
0881             isect = farthest_enter ; 
0882         }
0883         else if( nearest_exit.w > t_min && nearest_exit.w < RT_DEFAULT_MAX )
0884         {
0885             valid_isect = true ; 
0886             isect = nearest_exit ; 
0887         }
0888     }
0889 
0890 #ifdef DEBUG
0891     printf("//intersect_node_overlap num_sub %d  enter_count %d exit_count %d overlap_all %d valid_isect %d farthest_enter.w %10.4f nearest_exit.w %10.4f \n", 
0892        num_sub, enter_count, exit_count, overlap_all, valid_isect, farthest_enter.w, nearest_exit.w ); 
0893 #endif
0894 
0895 }
0896  
0897 
0898 
0899 
0900 
0901 
0902 
0903 
0904 /**
0905 intersect_node : some node are compound eg CSG_CONTIGUOUS, but mostly leaf nodes
0906 ----------------------------------------------------------------------------------
0907 
0908 Three level layout tree-node-leaf is needed to avoid intersect_node recursion which OptiX disallows. 
0909 
0910 **/
0911 
0912 
0913 INTERSECT_FUNC
0914 void intersect_node( bool& valid_isect, float4& isect, const CSGNode* node, const CSGNode* root, 
0915        const float4* plan, const qat4* itra, const float t_min , const float3& ray_origin , const float3& ray_direction, bool dumpxyz )
0916 {
0917     const unsigned typecode = node->typecode() ;  
0918 
0919 #ifdef DEBUG
0920     printf("//intersect_node typecode %d name %s \n", typecode, CSG::Name(typecode) ); 
0921 #endif
0922 
0923     valid_isect = false ; 
0924     switch(typecode)
0925     {
0926 #ifdef WITH_CONTIGUOUS
0927        case CSG_CONTIGUOUS:     intersect_node_contiguous(    valid_isect, isect, node, root, plan, itra, t_min, ray_origin, ray_direction, dumpxyz )  ; break ; 
0928 #endif
0929        case CSG_OVERLAP:        intersect_node_overlap(       valid_isect, isect, node, root, plan, itra, t_min, ray_origin, ray_direction, dumpxyz )  ; break ; 
0930        case CSG_DISCONTIGUOUS:  intersect_node_discontiguous( valid_isect, isect, node, root, plan, itra, t_min, ray_origin, ray_direction, dumpxyz )  ; break ; 
0931                    default:     intersect_leaf(               valid_isect, isect, node,       plan, itra, t_min, ray_origin, ray_direction, dumpxyz )  ; break ; 
0932     }
0933 
0934 #ifdef DEBUG_PIDXYZ
0935     //if(dumpxyz) printf("//intersect_node valid_isect:%d \n", valid_isect ); 
0936 #endif
0937 
0938 }
0939 
0940 /**
0941 distance_node
0942 ----------------
0943 
0944 **/
0945 
0946 INTERSECT_FUNC
0947 float distance_node( const float3& global_position, const CSGNode* node, const CSGNode* root, const float4* plan, const qat4* itra )
0948 {
0949     const unsigned typecode = node->typecode() ;  
0950     float distance ; 
0951     switch(typecode)
0952     {
0953 #ifdef WITH_CONTIGUOUS
0954         case CSG_CONTIGUOUS:       distance = distance_node_list( typecode,  global_position, node, root, plan, itra )  ; break ; 
0955 #endif
0956         case CSG_OVERLAP:          distance = distance_node_list( typecode,  global_position, node, root, plan, itra )  ; break ; 
0957         case CSG_DISCONTIGUOUS:    distance = distance_node_list( typecode,  global_position, node, root, plan, itra )  ; break ; 
0958         default:                   distance = distance_leaf(                 global_position, node, plan, itra )  ; break ; 
0959     }
0960     return distance ; 
0961 }
0962