LZ-Opticks-NVIDIA OptiX 6->7 : Notes

• https://github.com/simoncblyth/CSG/commits https://github.com/simoncblyth/CSGOptiX/commits
• https://github.com/simoncblyth/CSGOptiXGGeo/commits

PROGRESS : OptiXTest split in two, added Opticks/GGeo conversion

• CSG : CSGFoundry model, simple intersect headers : CPU testable
• CSGOptiX : OptiX 7 + pre-7 renders of CSGFoundry geometry
• CSGOptiXGGeo : loads Opticks/GGeo, converts to CSG and renders using CSGOptiX

NEXT STEPS:

• debug GGeo -> CSGFoundry conversion : relative transforms, instances
• Add CSG + CSGOptiX packages to Opticks
• performance test full geometry approaches (split/join GAS, 1 or more IAS)
• optixrap/cu/generate.cu : photon generation + propagation : expts
  • pull out simple headers common to: pre-7, 7, CPU testing
  • minimize code difference between : pre-7, 7, CPU-testing
  • trickery/mocking needed for CURAND on CPU? (templating?)

LONGTERM POSSIBILITY : Populate CSGFoundry model direct from Geant4 geometry ? [Disruptive]

CSG : CSGFoundry/CSGSolid/CSGPrim/CSGNode/csg_intersect_tree/..

struct CSGFoundry
{
...
      void upload();
...
      std::vector<CSGSolid>  solid ;
      std::vector<CSGPrim>   prim ;
      std::vector<CSGNode>   node ;
      std::vector<float4> plan ;
      std::vector<qat4>   tran ;
      std::vector<qat4>   itra ;
      std::vector<qat4>   inst ;// instance transforms

      CSGPrim*    d_prim ;
      CSGNode*    d_node ;
      float4*     d_plan ;
      qat4*       d_itra ;   // inverse CSG transforms
 };

• https://github.com/simoncblyth/CSG
• CSG : soon to become an Opticks package

All solids+constituents created via Foundry (ref by index)

• CSGSolid : 1 or more CSGPrim : ( CSGSolid -> GAS )
• CSGPrim : 1 or more CSGNode : (CSGPrim ~ G4VSolid)
  • (nodeOffset, numNode) -> SBT HitGroup record
• CSGNode : CSG constituents
  • basis shapes : sphere, box3, cone, cylinder, ...
  • boolean operators : union, intersection, difference
• qat4 : scale-rotate-translate transform
• float4 : planes (normal + distance to origin)
• csg_intersect_tree.h/csg_intersect_node.h/...
  • simple headers common to pre-7/7/CPU-testing
  • CSGScan : CPU intersect testing

CSGFoundry::upload

all solids in geometry -> only four GPU allocations : d_prim, d_node, d_plan, d_itra

https://github.com/simoncblyth/CSG/blob/main/CSGFoundry.h https://github.com/simoncblyth/CSG/blob/main/qat4.h

OptiX 7 + pre-7 intersecting same CSGFoundry geometry

24 rtBuffer prim_buffer; // geometry level context 
25
28 rtBuffer node_buffer;  // global context
29 rtBuffer itra_buffer;
30 rtBuffer plan_buffer; 
..
40 RT_PROGRAM void intersect(int primIdx)
41 {
42 const CSGPrim* prim = &prim_buffer[primIdx] ;
43 int nodeOffset = prim->nodeOffset() ;
44 int numNode = prim->numNode() ;
45 const CSGNode* node = &node_buffer[nodeOffset] ; 
46 const float4* plan = &plan_buffer[0] ;
47 const qat4*   itra = &itra_buffer[0] ;
48
49 float4 isect ;
50 if(intersect_prim(isect, numNode, node, plan, itra,
        ray.tmin , ray.origin, ray.direction ))
51 {
52     if(rtPotentialIntersection(isect.w))
53     {
55     shading_normal = make_float3( isect );
57     rtReportIntersection(0);
58     }
59 }
60 }

• https://github.com/simoncblyth/CSGOptiX

OptiX7Test.cu

150 extern "C" __global__ void __intersection__is()
151 {
152 HitGroupData* hg  = (HitGroupData*)optixGetSbtDataPointer();
153 int numNode = hg->numNode ;
154 int nodeOffset = hg->nodeOffset ;
155
156 const CSGNode* node = params.node + nodeOffset ;
157 const float4* plan = params.plan ;
158 const qat4*   itra = params.itra ;
159
160 const float  t_min = optixGetRayTmin() ;
161 const float3 ray_origin = optixGetObjectRayOrigin();
162 const float3 ray_direction = optixGetObjectRayDirection();
163
164 float4 isect ;
165 if(intersect_prim(isect, numNode, node, plan, itra,
                     t_min , ray_origin, ray_direction ))
166 {
...
175     optixReportIntersection( isect.w, hitKind, a0, a1, a2, a3 );
176 }
177 }

Minimize code split : 7, pre-7, CPU testing : same intersect_prim

CSGOptiXGGeo : loads Opticks/GGeo, converts to CSG and renders

 11 struct Converter
 12 {
 13     CSGFoundry* foundry ;
 14     const GGeo* ggeo ;
 16     float splay ;
 17
 18     Converter(CSGFoundry* foundry, const GGeo* ggeo ) ;
 19
 20     void convert(int repeatIdx,
                     int primIdx,
                     int partIdxRel );
 21     void convert_();
 22

 23     CSGSolid* convert_(unsigned repeatIdx );
 24     void addInstances(unsigned repeatIdx );
 25
 26     CSGPrim*  convert_(const GParts* comp,
                            unsigned primIdx );

 27     CSGNode*  convert_(const GParts* comp,
                           unsigned primIdx,
                           unsigned partIdxRel );
 28 };

• https://github.com/simoncblyth/CSGOptiXGGeo

 06 #include "CSGFoundry.h"
 07 #include "CSGOptiX.h"
 09 #include "Converter.h"
 10
 11 int main(int argc, char** argv)
 12 {
 13     int repeatIdx = argc > 1 ? atoi(argv[1]) : 0 ;
 ..
 19     OPTICKS_LOG(argc, argv);
 20     Opticks ok(argc, argv);
 21     ok.configure();
 22
 24     GGeo* ggeo = GGeo::Load(&ok); 
 26
 27     CSGFoundry foundry ;
 28     Converter conv(&foundry, ggeo, dump) ;
 29     conv.convert(repeatIdx, primIdx, partIdxRel);
 30
 31     CSGOptiX cx(&foundry);
 34     foundry.upload();   // uploads nodes, planes, transforms
 ..
 52     cx.setCE(ce, tmin, tmax);
 53     cx.render( tspec );
 55     return 0 ;
 56 }

• GGeo::Load geocache identified by OPTICKS_KEY envvar
• Perhaps: go direct Geant4 -> CSG ?
  • disruptive but significant simplification of NPY/NNode

CSGOptiXGGeo_0

• JUNO GGeo -> CSGFoundry
• CSGOptiX renders :
  • OptiX 5 (LHS)
  • OptiX 7 (RHS)
• Unexplained difference
  • remainder solid (~3000 volumes combined)

CSGOptiXGGeo_1

• JUNO GGeo -> CSGFoundry
• CSGOptiX renders :
  • OptiX 5 (LHS, white background)
  • OptiX 7 (RHS, grey background)

CSGOptiXGGeo_2

Converter debugging

• misses relative transforms perhaps ?
• instances not working yet

CSGOptiXGGeo_3

CSGOptiXGGeo_4

CSGOptiXGGeo_5

CSGOptiXGGeo_6

CSGOptiXGGeo_7

• cutout from tmin sphere

CSGOptiXGGeo_8

One JUNO solid (fastener) -> blank

• renders very slowly with old machinery

CSGOptiXGGeo_9

Proof that relative trans not applied

• should be many "planks" here

"Extra" Background Slides Follow

Two-Level Hierarchy : Instance transforms (TLAS) over Geometry (BLAS)

OptiX supports multiple instance levels : IAS->IAS->GAS BUT: Simple two-level is faster : works in hardware RT Cores

https://developer.nvidia.com/blog/introduction-nvidia-rtx-directx-ray-tracing/

AS

Acceleration Structure

TLAS (IAS)

4x4 transforms, refs to BLAS

BLAS (GAS)

triangles : vertices, indices

custom primitives : AABB

AABB

axis-aligned bounding box

SBT : Shader Binding Table

Flexibly binds together:

1. geometry objects
2. shader programs
3. data for shader programs

Hidden in OptiX 1-6 APIs

Optimizing Geometry : Split BLAS to avoid overlapping bbox

Optimization : deciding where to draw lines between:

1. structure and solid (IAS and GAS)
2. solids within GAS (bbox choice to minimize traversal intersection tests)

Where those lines are drawn defines the AS

https://developer.nvidia.com/blog/best-practices-using-nvidia-rtx-ray-tracing/

Optimizing Geometry : Merge BLAS when lots of overlaps

• lots of overlapping forces lots of intersections to find closest
• but too few bbox means the AS cannot help to avoid intersect tests
• balance required : needs experimentation and measurement to optimize

https://developer.nvidia.com/blog/best-practices-using-nvidia-rtx-ray-tracing/

Ray Intersection with Transformed Object -> Geometry Instancing

Fig 13.5 "Realistic Ray Tracing", Peter Shirley

Advantages apply equally to acceleration structures

Equivalent Intersects -> same t

1. ray with ellipsoid : M*p
2. M^-1 ray with sphere : p

Local Frame Advantages

1. simpler intersect (sphere vs ellipsoid)
2. closer to origin -> better precision

Geometry Instancing Advantages

• many objects share local geometry
  • orient+position with 4x4 M
• huge VRAM saving, less to copy

Requirements

• must not normalize ray direction
• normals transform differently
  • N' = N * M^-1T
  • (due to non-uniform scaling)

G4Boolean -> CUDA/OptiX Intersection Program Implementing CSG

Outside/Inside Unions

dot(normal,rayDir) -> Enter/Exit

• A + B boundary not inside other
• A * B boundary inside other

Complete Binary Tree, pick between pairs of nearest intersects:

• Nearest hit intersect algorithm [1] avoids state
  • sometimes Loop : advance t_min , re-intersect both
  • classification shows if inside/outside
• Evaluative [2] implementation emulates recursion:
  • recursion not allowed in OptiX intersect programs
  • bit twiddle traversal of complete binary tree
  • stacks of postorder slices and intersects
• Identical geometry to Geant4
  • solving the same polynomials
  • near perfect intersection match

[1] Ray Tracing CSG Objects Using Single Hit Intersections, Andrew Kensler (2006)

with corrections by author of XRT Raytracer http://xrt.wikidot.com/doc:csg

[2] https://bitbucket.org/simoncblyth/opticks/src/tip/optixrap/cu/csg_intersect_boolean.h

Similar to binary expression tree evaluation using postorder traverse.

Constructive Solid Geometry (CSG) : Shapes defined "by construction"

CSG Binary Tree

Primitives combined via binary operators

Simple by construction definition, implicit geometry.

• A, B implicit primitive solids
• A + B : union (OR)
• A * B : intersection (AND)
• A - B : difference (AND NOT)
• !B : complement (NOT) (inside <-> outside)

CSG expressions

• non-unique: A - B == A * !B
• represented by binary tree, primitives at leaves

3D Parametric Ray : ray(t) = r0 + t rDir

Ray Geometry Intersection

• primitive : find t roots of implicit eqn
• composite : pick primitive intersect, depending on CSG tree

How to pick exactly ?

CSG : Which primitive intersect to pick ?

In/On/Out transitions

Classical Roth diagram approach

• find all ray/primitive intersects
• recursively combine inside intervals using CSG operator
• works from leaves upwards

Computational requirements:

• find all intersects, store them, order them
• recursive traverse

BUT : High performance on GPU requires:

• massive parallelism -> more the merrier
• low register usage -> keep it simple
• small stack size -> avoid recursion

Classical approach not appropriate on GPU

CSG Complete Binary Tree Serialization -> simplifies GPU side

Geant4 solid -> CSG binary tree (leaf primitives, non-leaf operators, 4x4 transforms on any node)

Serialize to complete binary tree buffer:

• no need to deserialize, no child/parent pointers
• bit twiddling navigation avoids recursion
• simple approach profits from small size of binary trees
• BUT: very inefficient when unbalanced

Height 3 complete binary tree with level order indices:

                                                   depth     elevation

                     1                               0           3

          10                   11                    1           2

     100       101        110        111             2           1

 1000 1001  1010 1011  1100 1101  1110  1111         3           0

postorder_next(i,elevation) = i & 1 ? i >> 1 : (i << elevation) + (1 << elevation) ; // from pattern of bits

Postorder tree traverse visits all nodes, starting from leftmost, such that children are visited prior to their parents.

Evaluative CSG intersection Pseudocode : recursion emulated

fullTree = PACK( 1 << height, 1 >> 1 )  // leftmost, parent_of_root(=0)
tranche.push(fullTree, ray.tmin)

while (!tranche.empty)         // stack of begin/end indices 
{
    begin, end, tmin <- tranche.pop  ; node <- begin ;
    while( node != end )                   // over tranche of postorder traversal 
    {
        elevation = height - TREE_DEPTH(node) ;
        if(is_primitive(node)){ isect <- intersect_primitive(node, tmin) ;  csg.push(isect) }
        else{
            i_left, i_right = csg.pop, csg.pop            // csg stack of intersect normals, t 
            l_state = CLASSIFY(i_left, ray.direction, tmin)
            r_state = CLASSIFY(i_right, ray.direction, tmin)
            action = LUT(operator(node), leftIsCloser)(l_state, r_state)

            if(      action is ReturnLeft/Right)     csg.push(i_left or i_right)
            else if( action is LoopLeft/Right)
            {
                left = 2*node ; right = 2*node + 1 ;
                endTranche = PACK( node,  end );
                leftTranche = PACK(  left << (elevation-1), right << (elevation-1) )
                rightTranche = PACK(  right << (elevation-1),  node  )
                loopTranche = action ? leftTranche : rightTranche

                tranche.push(endTranche, tmin)
                tranche.push(loopTranche, tminAdvanced )  // subtree re-traversal with changed tmin 
                break ; // to next tranche
            }
        }
        node <- postorder_next(node, elevation)         // bit twiddling postorder 
    }
}
isect = csg.pop();         // winning intersect  

https://bitbucket.org/simoncblyth/opticks/src/tip/optixrap/cu/csg_intersect_boolean.h

CSG Deep Tree : Positivize tree using De Morgan's laws

1st step to allow balancing : Positivize : remove CSG difference di operators

                                                     ...    ...

                                               un          cy

                                       un          cy

                               un          cy

                       un          cy

               un          cy

       di          cy

   cy      cy

                                                     ...    ...

                                               un          cy

                                       un          cy

                               un          cy

                       un          cy

               un          cy

       in          cy

   cy      !cy

CSG Examples