Opticks : GPU Optical Photon Simulation for Particle Physics with NVIDIA OptiX

Outline

• Context and Problem
  • p2: Jiangmen Underground Neutrino Observatory (JUNO)
  • p3: JUNO Optical Photon Simulation Problem...
  • p4: Optical Photon Simulation ≈ Ray Traced Image Rendering
• NVIDIA Tools to create Solution
  • p5: NVIDIA Ampere : 2nd Generation RTX
  • p6: NVIDIA OptiX Ray Tracing Engine
  • p7: NVIDIA OptiX 7 : Entirely new thin API
• Opticks : Introduction + New Features
  • p8,9: Geant4 + Opticks Hybrid Workflow : External Optical Photon Simulation
  • p10: New : Collection Efficiency Culling on GPU
  • p11: New : CSGFoundry Geometry Model
• Opticks : 1st JUNO Renders with NVIDIA OptiX 7
  • p12:Ray trace renders
  • p13:Geometry factorization
  • p15:Ray trace times for various geometries
• p17: Summary + Links
• p18: Acknowledgement : Opticks "Hackathon" Series for NVIDIA OptiX 6->7
• p19: LZ with Opticks

JUNO_Intro_2

Optical Photon Simulation Problem...

Optical Photon Simulation ≈ Ray Traced Image Rendering

Not a Photo, a Calculation

http://on-demand.gputechconf.com/siggraph/2013/presentation/SG3106-Building-Ray-Tracing-Applications-OptiX.pdf

simulation

photon parameters at sensors (PMTs)

rendering

pixel values at image plane

Much in common : geometry, light sources, optical physics

• both limited by ray geometry intersection, aka ray tracing

Many Applications of ray tracing :

• advertising, design, architecture, films, games,...
• -> huge efforts to improve hw+sw over 30 yrs

Ampere : 2nd Generation RTX

NVIDIA Ampere (2020):

"...triple double over Turing (2018, 10 GigaRays/s)..."

• RT Core : ray trace dedicated GPU hardware
• NVIDIA GeForce RTX 3090
  • 10,496 CUDA Cores, 28GB VRAM, USD 1499
• ray trace performance continues rapid improvement

NVIDIA® OptiX™ Ray Tracing Engine -- http://developer.nvidia.com/optix

OptiX Raytracing Pipeline

Analogous to OpenGL rasterization pipeline:

OptiX makes GPU ray tracing accessible

• accelerates ray-geometry intersections
• simple : single-ray programming model
• "...free to use within any application..."
• access RT Cores[1] with OptiX 6+ via RTX™ mode

NVIDIA expertise:

• OptiX pre-7 : ~linear scaling up to 4 GPUs
• acceleration structure creation + traversal (Blue)
• instanced sharing of geometry + acceleration structures
• compiler optimized for GPU ray tracing

https://developer.nvidia.com/rtx

User provides (Yellow):

• ray generation
• geometry bounding box, intersects

[1] Turing+ GPUs eg NVIDIA TITAN RTX

NVIDIA OptiX 7 : Entirely new thin API (Introduced Aug 2019)

NVIDIA OptiX 6->7 : drastically slimmed down

• headers only (no library, just Driver)
• low-level CUDA-centric thin API (Vulkan-ized)
• Minimal host state, All host functions are thread-safe
• GPU launches : explicit, asynchronous (CUDA streams)
• near perfect scaling to 4 GPUs, for free
• Shared CPU/GPU geometry context
• GPU memory management
• Multi-GPU support

Advantages

More control/flexibility over everything.

• Fully benefit from future GPUs
• Keep pace with state-of-the-art GPU ray tracing

Disadvantages

Demands much more developer effort than OptiX 6

• Major re-implementation of Opticks required

NEWS: Major Progress on 6->7 (p11,..)

Geant4OpticksWorkflow

Geant4OpticksWorkflow 2

Opticks + JUNO Progress : Efficiency Hit Culling on GPU

On GPU Efficiency Culling

• upload sensor angular efficiencies
• -> COLLECT OR CULL hit flags

CUDA Thrust stream compaction

• thrust::count_if
• thrust::copy_if
• Download hits with : hitmask = COLLECT

• Reduces hit CPU memory by factor of the efficiency
• reduce size of GPU->CPU downloads by factor of efficiency
• all downloaded hits -> electronics simulation

Mock angular efficiency test of GPU texture machinery :

• texturing sensor surfaces with the efficiency
• mockup striped theta efficiency, two random categories
• cosine phi efficiency variation for one category

New "Foundry" Model : Shared CPU/GPU Geometry Context

struct CSGFoundry
{
   void upload(); // to GPU 
...
   std::vector<CSGSolid>  solid ; // compounds (eg PMT)
   std::vector<CSGPrim>   prim ;
   std::vector<CSGNode>   node ; // shapes, operators

   std::vector<float4> plan ; // planes
   std::vector<qat4>   tran ; // CSG transforms
   std::vector<qat4>   itra ; // inverse CSG transforms
   std::vector<qat4>   inst ; // instance transforms

   // entire geometry in four GPU allocations
   CSGPrim*    d_prim ;
   CSGNode*    d_node ;
   float4*     d_plan ;
   qat4*       d_itra ;
 };

• replaces geometry context dropped in OptiX 6->7
• array-based -> simple, inherent serialization + persisting
• entire geometry in 4 GPU allocations

Simple intersect headers, common CPU/GPU types

• use with : pre-7, 7 + testing on CPU

https://github.com/simoncblyth/CSG "Foundry" model

csg_intersect_tree.h/csg_intersect_node.h/...

simple headers common to pre-7/7/CPU-testing

https://github.com/simoncblyth/CSG_GGeo

Convert Opticks/GGeo -> CSGFoundry

https://github.com/simoncblyth/CSGOptiX

OptiX 7 + pre-7 rendering

GAS : Geometry Acceleration Structure

IAS : Instance Acceleration Structure

CSG : Constructive Solid Geometry

[9]cxr_i0_t8,_-1 : EXCLUDE SLOWEST

Current JUNO Geometry : Auto-Factorized by "progeny digest"

• ridx:0 "remainder" Prim
  • Prim that did not pass instancing criteria, on number of repeats + complexity
  • TODO: tune criteria to instance more, reducing remainder Prim (Expect: 3084->~ 84)
• ridx:1,2,3,4
  • four types of PMT, all with 5 Prim
• ridx:5,6,7,8
  • same 590x assembly but not grouped together : as siblings (not parent-child like PMTs)
  • TODO: implement instancing of siblings, combining 4 -> 1

Increasing instancing : reduces memory for geometry -> improved performance

JUNO OptiX 7 : "Foundry" Geometry Scan

JUNO Geometry : OptiX 7 Ray Trace Times ~2M pixels : TITAN RTX

JUNO ALL PMTs : 2M ray traced pixels in 0.0097 s : NVIDIA TITAN RTX, NVIDIA OptiX 7.0.0, Opticks

Summary and Links

Opticks : state-of-the-art GPU ray traced optical simulation integrated with Geant4. Geometry progress with OptiX 7 suggests integration within months is achievable.

• Efficiency culling decision moved to GPU, reducing CPU hit memory
• "Foundry" geometry model implemented to support NVIDIA OptiX 7 API.
• First JUNO OptiX 7 renders achieved.

Acknowledgement : Opticks "Hackathon" Series for NVIDIA OptiX 6->7

• Suggested and organized by LZ, LBNL, NERSC
• ~fortnightly meetings since Feb 2021
• Valuable advice from NVIDIA engineers
• see next talk by Oisin Creaner, LZ with Opticks

LZ with Opticks (Images from Sam Eriksen, University of Bristol)