|
|
|||
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
| [ Source navigation ] | [ Diff markup ] | [ Identifier search ] | [ general search ] |
|
This page was automatically generated by the 2.3.7 LXR engine. The LXR team |
|