eic-opticks/CSGOptiX/CSGOptiX.cc

0001 /**
0002 CSGOptiX.cc
0003 ============
0004
0005 * NOTE : < 7 BRANCH NO LONGER VIABLE BUT ITS EXPEDIENT TO KEEP
0006   IT FOR LAPTOP COMPILATION
0007
0008 This code contains two branches for old (OptiX < 7) and new (OptiX 7+) API
0009
0010 Branched aspects:
0011
0012 1. CSGOptiX7.cu vs CSGOptiX6.cu
0013 2. Ctx/SBT/PIP  vs Six
0014 3. new workflow uses uploaded params extensively from CUDA device code,
0015    old workflow with *Six* uses hostside params to populate the optix context
0016    which then gets used on device
0017
0018    CSGOptiX::prepareParam (called just before launch)
0019
0020    * new workflow: uploads param
0021    * old workflow: Six::updateContext passes hostside param into optix context variables
0022
0023    Six MATCHING NEEDS DUPLICATION FROM Params INTO CONTEXT VARIABLES AND BUFFERS
0024
0025 HMM: looking like getting qudarap/qsim.h to work with OptiX < 7 is more effort than it is worth
0026
0027 * would have to shadow it into context variables
0028 * CUDA textures would not work without optix textureSamplers
0029
0030 **/
0031
0032 #include <iostream>
0033 #include <cstdlib>
0034 #include <chrono>
0035
0036 #include <optix.h>
0037 #if OPTIX_VERSION < 70000
0038 #else
0039 #include <optix_stubs.h>
0040 #endif
0041
0042 #include <cuda_runtime.h>
0043 #include <glm/glm.hpp>
0044
0045 // sysrap
0046 #include "sproc.h"
0047 #include "ssys.h"
0048 #include "spath.h"
0049 #include "smeta.h"
0050 #include "SProf.hh"
0051
0052 #include "SGLM.h"
0053 #include "NP.hh"
0054 #include "SRG.h"
0055 #include "SCAM.h"
0056 #include "SEventConfig.hh"
0057 #include "SGeoConfig.hh"
0058 #include "SSim.hh"
0059 #include "SStr.hh"
0060 #include "SEvt.hh"
0061 #include "SMeta.hh"
0062 #include "SPath.hh"
0063 #include "SVec.hh"
0064 #include "SLOG.hh"
0065 #include "scuda.h"
0066 #include "squad.h"
0067 #include "sframe.h"
0068
0069 // csg
0070 #include "CSGPrim.h"
0071 #include "CSGFoundry.h"
0072 #include "CSGView.h"
0073
0074 // qudarap
0075 #include "qrng.h"
0076 #include "QU.hh"
0077 #include "QSim.hh"
0078 #include "qsim.h"
0079 #include "QEvt.hh"
0080
0081 // CSGOptiX
0082 #include "Frame.h"
0083 #include "Params.h"
0084 #include "config.h"
0085
0086 #if OPTIX_VERSION < 70000
0087 #include "Six.h"
0088 #else
0089 #include "Ctx.h"
0090 #include "CUDA_CHECK.h"
0091 #include "OPTIX_CHECK.h"
0092 #include "PIP.h"
0093 #include "SBT.h"
0094 #endif
0095
0096 #include "CSGOptiX.h"
0097
0098 const plog::Severity CSGOptiX::LEVEL = SLOG::EnvLevel("CSGOptiX", "DEBUG" );
0099 CSGOptiX* CSGOptiX::INSTANCE = nullptr ;
0100 CSGOptiX* CSGOptiX::Get()
0101 {
0102     return INSTANCE ;
0103 }
0104
0105
0106 int CSGOptiX::Version()
0107 {
0108     int vers = 0 ;
0109 #if OPTIX_VERSION < 70000
0110     vers = 6 ;
0111 #else
0112     vers = 7 ;
0113 #endif
0114     return vers ;
0115 }
0116
0117
0118 /**
0119 CSGOptiX::RenderMain CSGOptiX::SimtraceMain CSGOptiX::SimulateMain
0120 ---------------------------------------------------------------------
0121
0122 These three mains are use by the minimal main tests:
0123
0124 +-------------+---------------------------+---------------------+
0125 |  script     | mains                     | notes               |
0126 +=============+===========================+=====================+
0127 | cxr_min.sh  | tests/CSGOptiXRMTest.cc   | minimal render      |
0128 +-------------+---------------------------+---------------------+
0129 | cxt_min.sh  | tests/CSGOptiXTMTest.cc   | minimal simtrace    |
0130 +-------------+---------------------------+---------------------+
0131 | cxs_min.sh  | tests/CSGOptiXSMTest.cc   | minimal simulate    |
0132 +-------------+---------------------------+---------------------+
0133
0134 Note that SEvt setup and frame hookup formerly
0135 done in the main is now moved into CSGFoundry::AfterLoadOrCreate
0136 and invokation of SEvt::beginOfEvent SEvt::endOfEvent is done from
0137 QSim
0138
0139 Note that currently rendering persisting does not use SEvt in the
0140 same way as simtrace and simulate, but it could do in future.
0141
0142 **/
0143
0144 int CSGOptiX::RenderMain() // static
0145 {
0146     SEventConfig::SetRGModeRender();
0147     CSGFoundry* fd = CSGFoundry::Load();
0148     CSGOptiX* cx = CSGOptiX::Create(fd) ;
0149     cx->render();
0150     delete cx ;
0151     return 0 ;
0152 }
0153
0154 /**
0155 CSGOptiX::SimtraceMain
0156 -------------------------
0157
0158 Q: Where is the frame needed when simtracing ?
0159 A: Its needed for genstep preparation with SEvt
0160    so the setup by SSim::afterLoadOrCreate giving it to SEvt
0161    should be OK if that happens before gensteps are prepared
0162
0163
0164
0165 **/
0166
0167
0168 int CSGOptiX::SimtraceMain()
0169 {
0170     SEventConfig::SetRGModeSimtrace();
0171     CSGFoundry* fd = CSGFoundry::Load();
0172     CSGOptiX* cx = CSGOptiX::Create(fd) ;
0173     cx->simtrace(0);
0174     delete cx ;
0175     return 0 ;
0176 }
0177
0178
0179 /**
0180 CSGOptiX::SimulateMain
0181 -----------------------
0182
0183 **/
0184
0185 int CSGOptiX::SimulateMain() // static
0186 {
0187     SProf::Add("CSGOptiX__SimulateMain_HEAD");
0188     SEventConfig::SetRGModeSimulate();
0189     CSGFoundry* fd = CSGFoundry::Load();
0190     CSGOptiX* cx = CSGOptiX::Create(fd) ;
0191     bool reset = true ;
0192     for(int i=0 ; i < SEventConfig::NumEvent() ; i++) cx->simulate(i, reset);
0193     SProf::UnsetTag();
0194     SProf::Add("CSGOptiX__SimulateMain_TAIL");
0195     SProf::Write();
0196     cx->write_Ctx_log();
0197     delete cx ;
0198     return 0 ;
0199 }
0200
0201
0202
0203
0204
0205 const char* CSGOptiX::Desc()
0206 {
0207     std::stringstream ss ;
0208     ss << "CSGOptiX::Desc"
0209        << " Version " << Version()
0210 #ifdef WITH_CUSTOM4
0211        << " WITH_CUSTOM4 "
0212 #else
0213        << " NOT:WITH_CUSTOM4 "
0214 #endif
0215        ;
0216     std::string str = ss.str();
0217     return strdup(str.c_str());
0218 }
0219
0220
0221 const char* CSGOptiX::desc() const
0222 {
0223     std::stringstream ss ;
0224     ss << Desc() ;
0225 #if OPTIX_VERSION < 70000
0226 #else
0227     ss << pip->desc() ;
0228 #endif
0229     std::string s = ss.str();
0230     return strdup(s.c_str());
0231 }
0232
0233
0234
0235
0236 /**
0237 CSGOptiX::InitEvt  TODO : THIS DOES NOT USE GPU : SO SHOULD BE ELSEWHERE
0238 --------------------------------------------------------------------------
0239
0240 Invoked from CSGOptiX::Create
0241
0242
0243 Q: Why the SEvt geometry connection ?
0244 A: Needed for global to local transform conversion
0245
0246 Q: What uses SEVt::setGeo (SGeo) ?
0247 A: Essential set_matline of Cerenkov Genstep
0248
0249 **/
0250
0251 void CSGOptiX::InitEvt( CSGFoundry* fd  )
0252 {
0253     SEvt* sev = SEvt::CreateOrReuse(SEvt::EGPU) ;
0254
0255 #ifdef WITH_OLD_FRAME
0256     sev->setGeo((SGeo*)fd);
0257 #else
0258     const SSim* ssim = fd->getSim();
0259     sev->setSim(ssim);
0260 #endif
0261
0262     std::string* rms = SEvt::RunMetaString() ;
0263     assert(rms);
0264
0265     bool stamp = false ;
0266     smeta::Collect(*rms, "CSGOptiX__InitEvt", stamp );
0267 }
0268
0269 /**
0270 CSGOptiX::InitSim
0271 -------------------
0272
0273 Invoked from CSGOptiX::Create
0274 Instanciation of QSim/QEvt requires an SEvt instance
0275
0276 **/
0277
0278 void CSGOptiX::InitSim( SSim* ssim  )
0279 {
0280     LOG(LEVEL) << "[" ; ;
0281
0282     if(SEventConfig::IsRGModeRender()) return ;
0283
0284     LOG_IF(fatal, ssim == nullptr ) << "simulate/simtrace modes require SSim/QSim setup" ;
0285     assert(ssim);
0286
0287
0288     if( ssim->hasTop() == false )
0289     {
0290         ssim->serialize() ;  // SSim::serialize stree::serialize into NPFold  (moved from InitEvt)
0291     }
0292     else
0293     {
0294         LOG(LEVEL) << " NOT calling SSim::serialize : as already done, loaded ? " ;
0295     }
0296
0297     QSim::UploadComponents(ssim);
0298
0299     QSim* qs = QSim::Create() ;
0300
0301     LOG(LEVEL) << "]" << qs->desc() ;
0302 }
0303
0304
0305
0306 /**
0307 CSGOptiX::InitMeta
0308 -------------------
0309
0310 Invoked from CSGOptiX::Create prior to instanciation
0311
0312 **/
0313
0314 void CSGOptiX::InitMeta()
0315 {
0316     std::string gm = SEventConfig::GetGPUMeta() ;     // (QSim) scontext sdevice::brief
0317     SEvt::SetRunMetaString("GPUMeta", gm.c_str() );  // set CUDA_VISIBLE_DEVICES to control
0318
0319     std::string switches = QSim::Switches() ;
0320     SEvt::SetRunMetaString("QSim__Switches", switches.c_str() );
0321
0322 #ifdef WITH_CUSTOM4
0323     std::string c4 = "TBD" ; //C4Version::Version(); // octal version number bug in Custom4 v0.1.8 : so skip the version metadata
0324     SEvt::SetRunMetaString("C4Version", c4.c_str());
0325 #else
0326     SEvt::SetRunMetaString("C4Version", "NOT-WITH_CUSTOM4" );
0327 #endif
0328
0329 }
0330
0331
0332 /**
0333 CSGOptiX::InitGeo
0334 -------------------
0335
0336 Invoked from CSGOptiX::Create
0337
0338 CSGFoundry not const as upload sets device pointers
0339 CSGOptiX::InitGeo currently takes 20s for full JUNO geometry,
0340 where the total gxs.sh running time for one event is 24s.
0341
0342 HMM:that was optimized down to under 1s, by removal of some unused stree.h stuff ?
0343
0344 **/
0345
0346 void CSGOptiX::InitGeo(  CSGFoundry* fd )
0347 {
0348     LOG(LEVEL) << "[" ; ;
0349     fd->upload();
0350     LOG(LEVEL) << "]" ; ;
0351 }
0352
0353
0354
0355 /**
0356 CSGOptiX::Create
0357 --------------------
0358
0359 Canonical invokation from G4CXOpticks::setGeometry_ when a GPU is detected
0360
0361
0362 WIP: static methods cannot be enforced in a protocol, so perhaps just do the below
0363 within the ctor ?
0364
0365 **/
0366
0367 CSGOptiX* CSGOptiX::Create(CSGFoundry* fd )
0368 {
0369     SProf::Add("CSGOptiX__Create_HEAD");
0370     LOG(LEVEL) << "[ fd.descBase " << ( fd ? fd->descBase() : "-" ) ;
0371
0372     SSim* ssim = const_cast<SSim*>(fd->sim) ;
0373
0374
0375     InitEvt(fd);
0376     InitSim(ssim);    // QSim instanciation after uploading SSim arrays
0377     InitMeta();       // recording GPU, switches etc.. into run metadata
0378     InitGeo(fd);      // uploads geometry
0379
0380     CSGOptiX* cx = new CSGOptiX(fd) ;
0381
0382     if(!SEventConfig::IsRGModeRender())
0383     {
0384         QSim* qs = QSim::Get() ;
0385         qs->setLauncher(cx);
0386     }
0387
0388
0389     LOG(LEVEL) << "]" ;
0390     SProf::Add("CSGOptiX__Create_TAIL");
0391     return cx ;
0392 }
0393
0394
0395
0396
0397
0398
0399
0400
0401 Params* CSGOptiX::InitParams( int raygenmode, const SGLM* sglm  ) // static
0402 {
0403     LOG(LEVEL) << "[" ;
0404     return new Params(raygenmode, sglm->Width(), sglm->Height(), 1 ) ;
0405     LOG(LEVEL) << "]" ;
0406 }
0407
0408
0409
0410
0411
0412 CSGOptiX::~CSGOptiX()
0413 {
0414     destroy();
0415 }
0416
0417 /**
0418 CSGOptiX::CSGOptiX
0419 --------------------
0420
0421 Instanciated at near to main level in both running modes:
0422
0423 * pure-Opticks running (no Geant4) eg via cxs_min.sh uses CSGOptiX::SimulateMain
0424   that instanciates via CSGOptiX::Create with event processing CSGOptiX::simulate
0425   called directly from CSGOptiX::SimulateMain
0426
0427 * Geant4 integrated running eg from OJ python "main", instanciation again uses
0428   CSGOptiX::Create that is invoked from G4CXOpticks::setGeometry_
0429   with event processing CSGOptiX::simulate called from
0430   junoSD_PMT_v2_Opticks::EndOfEvent_Simulate
0431
0432 **/
0433
0434
0435 CSGOptiX::CSGOptiX(const CSGFoundry* foundry_)
0436     :
0437     sglm(SGLM::Get()),
0438     flight(SGeoConfig::FlightConfig()),
0439     foundry(foundry_),
0440     outdir(SEventConfig::OutFold()),
0441     _optixpath(std::getenv("CSGOptiX__optixpath")),
0442     optixpath(spath::Resolve(gphox::Config::PtxPath("CSGOptiX7.ptx").c_str())),
0443     tmin_model(ssys::getenvfloat("TMIN",0.1)),    // CAUTION: tmin very different in rendering and simulation
0444     kernel_count(0),
0445     raygenmode(SEventConfig::RGMode()),
0446     params(InitParams(raygenmode,sglm)),
0447     ctx(nullptr),
0448     pip(nullptr),
0449     sbt(nullptr),
0450     framebuf(nullptr),
0451     meta(new SMeta),
0452     kernel_dt(0.),
0453     sctx(nullptr),
0454     sim(QSim::Get()),
0455     qev(sim == nullptr  ? nullptr : sim->qev)   // QEvt
0456 {
0457     init();
0458 }
0459
0460 void CSGOptiX::init()
0461 {
0462     sglm->addlog("CSGOptiX::init", "start");
0463
0464     LOG(LEVEL)
0465         << "["
0466         << " raygenmode " << raygenmode
0467         << " SRG::Name(raygenmode) " << SRG::Name(raygenmode)
0468         << " sim " << sim
0469         << " qev " << qev
0470         ;
0471
0472     assert( outdir && "expecting OUTDIR envvar " );
0473
0474     LOG(LEVEL) << " _optixpath " << _optixpath  ;
0475     LOG(LEVEL) << " optixpath " << optixpath  ;
0476
0477     initMeta();
0478     initCtx();
0479     initPIP();
0480     initSBT();
0481     initCheckSim();
0482     initStack();
0483     initParams();
0484     initGeometry();
0485     initSimulate();
0486
0487 #ifdef WITH_OLD_FRAME
0488     initFrame();
0489 #endif
0490
0491     initRender();
0492     initPIDXYZ();
0493
0494     LOG(LEVEL) << "]" ;
0495 }
0496
0497 /**
0498 CSGOptiX::initMeta
0499 -------------------
0500
0501 Record metadata regarding the *optixpath* kernel source into RunMeta
0502
0503 **/
0504
0505
0506 void CSGOptiX::initMeta()
0507 {
0508     int64_t mtime = spath::last_write_time(optixpath);
0509     std::string str = sstamp::Format(mtime);
0510     int64_t age_secs = sstamp::age_seconds(mtime);
0511     int64_t age_days = sstamp::age_days(mtime);
0512
0513     SEvt::SetRunMetaString("optixpath", optixpath );
0514     SEvt::SetRunMetaString("optixpath_mtime_str", str.c_str() );
0515     SEvt::SetRunMeta<int64_t>("optixpath_mtime", mtime );
0516     SEvt::SetRunMeta<int64_t>("optixpath_age_secs", age_secs );
0517     SEvt::SetRunMeta<int64_t>("optixpath_age_days", age_days );
0518 }
0519
0520
0521 /**
0522 CSGOptiX::initCtx
0523 -------------------
0524
0525 Instanciate the OptixDeviceContext
0526
0527 **/
0528
0529 void CSGOptiX::initCtx()
0530 {
0531     LOG(LEVEL) << "[" ;
0532     ctx = new Ctx ;
0533     LOG(LEVEL) << std::endl << ctx->desc() ;
0534     LOG(LEVEL) << "]" ;
0535 }
0536
0537
0538 /**
0539 CSGOptiX::initPIP
0540 -------------------
0541
0542 Instanciate PIP pipeline
0543
0544 **/
0545
0546 void CSGOptiX::initPIP()
0547 {
0548     LOG(LEVEL) << "["  ;
0549     LOG(LEVEL)
0550         << " optixpath " << ( optixpath ? optixpath : "-" )
0551         ;
0552
0553     pip = new PIP(optixpath, ctx->props ) ;
0554     LOG(LEVEL) << "]" ;
0555 }
0556
0557 /**
0558 CSGOptiX::initSBT
0559 --------------------
0560
0561 Instanciate SBT shader binding table
0562
0563 **/
0564
0565
0566 void CSGOptiX::initSBT()
0567 {
0568     LOG(LEVEL) << "[" ;
0569     sbt = new SBT(pip) ;
0570     LOG(LEVEL) << "]" ;
0571 }
0572
0573
0574
0575 /**
0576 CSGOptiX::initCheckSim
0577 -----------------------
0578
0579 Check (QSim)sim instance for non-render modes
0580
0581 **/
0582
0583
0584 void CSGOptiX::initCheckSim()
0585 {
0586     if(SEventConfig::IsRGModeRender()) return ;
0587     LOG(LEVEL) << " sim " << sim << " qev " << qev ;
0588     LOG_IF(fatal, sim == nullptr) << "simtrace/simulate modes require instanciation of QSim before CSGOptiX " ;
0589     assert(sim);
0590 }
0591
0592
0593 void CSGOptiX::initStack()
0594 {
0595     LOG(LEVEL);
0596     pip->configureStack();
0597 }
0598
0599 void CSGOptiX::initParams()
0600 {
0601     params->device_alloc();
0602 }
0603
0604 /**
0605 CSGOptiX::initGeometry
0606 ------------------------
0607
0608 Notice that the geometry is uploaded to GPU before calling this by CSGOptiX::InitGeo
0609 The SBT::setFoundry kicks off the creation of the NVIDIA OptiX geometry
0610 from the uploaded CSGFoundry with SBT::createGeom.
0611
0612 **/
0613
0614 void CSGOptiX::initGeometry()
0615 {
0616     LOG(LEVEL) << "[" ;
0617     params->node = foundry->d_node ;
0618     params->plan = foundry->d_plan ;
0619     params->tran = nullptr ;
0620     params->itra = foundry->d_itra ;
0621
0622     bool is_uploaded =  params->node != nullptr ;
0623     LOG_IF(fatal, !is_uploaded) << "foundry must be uploaded prior to CSGOptiX::initGeometry " ;
0624     assert( is_uploaded );
0625
0626     LOG(LEVEL) << "[ sbt.setFoundry " ;
0627     sbt->setFoundry(foundry);
0628     params->handle = sbt->getTOPHandle() ;
0629     LOG(LEVEL) << "] sbt.setFoundry " ;
0630     LOG(LEVEL) << "]" ;
0631 }
0632
0633
0634 /**
0635 CSGOptiX::initSimulate
0636 ------------------------
0637
0638 * Once only (not per-event) simulate setup tasks ..  perhaps rename initPhys
0639
0640 Sets device pointers for params.sim params.evt so must be after upload
0641
0642 Q: where are sim and evt uploaded ?
0643 A: QSim::QSim and QSim::init_sim are where sim and evt are populated and uploaded
0644
0645
0646 HMM: get d_sim (qsim.h) now holds d_evt (sevent.h) but this is getting evt again rom QEvt ?
0647 TODO: eliminate params->evt to make more use of the qsim.h encapsulation
0648
0649 **/
0650
0651 void CSGOptiX::initSimulate()
0652 {
0653     LOG(LEVEL) ;
0654     params->sim = sim ? sim->getDevicePtr() : nullptr ;  // qsim<float>*
0655     params->evt = qev ? qev->getDevicePtr() : nullptr ;  // sevent*
0656
0657     params->tmin0 = SEventConfig::PropagateEpsilon0() ;  // epsilon used after step points with flags in below mask
0658     params->PropagateEpsilon0Mask = SEventConfig::PropagateEpsilon0Mask();  // eg from CK|SI|TO|SC|RE
0659
0660     params->PropagateRefine = SEventConfig::PropagateRefine();
0661     params->PropagateRefineDistance = SEventConfig::PropagateRefineDistance();  // approx distance beyond which to refine intersect with 2nd trace
0662
0663     params->tmin = SEventConfig::PropagateEpsilon() ;  // eg 0.1 0.05 to avoid self-intersection off boundaries
0664     params->tmax = 1000000.f ;
0665     params->max_time = SEventConfig::MaxTime() ;
0666
0667
0668 }
0669
0670
0671
0672
0673
0674
0675
0676
0677 #ifdef WITH_OLD_FRAME
0678
0679 /**
0680 CSGOptiX::initFrame (formerly G4CXOpticks::setupFrame)
0681 ---------------------------------------------------------
0682
0683 The frame used depends on envvars INST, MOI, OPTICKS_INPUT_PHOTON_FRAME
0684 it comprises : fr.ce fr.m2w fr.w2m set by CSGTarget::getFrame
0685
0686 Q: why is the frame needed ?
0687 A: cx rendering viewpoint, input photon frame and the simtrace genstep grid
0688    are all based on the frame center, extent and transforms
0689
0690 Q: Given the sframe and SEvt are from sysrap it feels too high level to do this here,
0691    should be at CSG or sysrap level perhaps ?
0692    And then CSGOptix could grab the SEvt frame at its initialization.
0693
0694 TODO: see CSGFoundry::AfterLoadOrCreate for maybe auto frame hookup
0695
0696 **/
0697
0698 void CSGOptiX::initFrame()
0699 {
0700     assert(0);
0701
0702     sframe _fr = foundry->getFrameE() ;   // TODO: migrate to lighweight sfr from stree level
0703     LOG(LEVEL) << _fr ;
0704
0705     SEvt::SetFrame(_fr) ;
0706
0707     sfr _lfr = _fr.spawn_lite();
0708     setFrame(_lfr);
0709 }
0710
0711 #endif
0712
0713
0714
0715 /**
0716 CSGOptiX::initRender
0717 --------------------------
0718
0719 To use externally managed device pixels call this
0720 prior to render/render_launch with the device pixel pointer,
0721 otherwise this is called with nullptr d_pixel that arranges
0722 internally allocated device pixels.
0723
0724 **/
0725
0726 void CSGOptiX::initRender()
0727 {
0728     LOG(LEVEL) << "[" ;
0729     framebuf = new Frame(params->width, params->height, params->depth, nullptr ) ;
0730     LOG(LEVEL) << "]" ;
0731
0732     if(SEventConfig::IsRGModeRender())
0733     {
0734         LOG(LEVEL) << "FORMERLY CALLED CSGOptiX::setFrame FROM HERE" ;
0735         //setFrame(); // MOI
0736     }
0737
0738     params->pixels = framebuf->d_pixel ;
0739     params->isect  = framebuf->d_isect ;
0740 #ifdef WITH_FRAME_PHOTON
0741     params->fphoton = framebuf->d_photon ;
0742 #else
0743     params->fphoton = nullptr ;
0744 #endif
0745 }
0746
0747 void CSGOptiX::initPIDXYZ()
0748 {
0749     qvals(params->pidxyz, "PIDXYZ", "-1:-1:-1", -1 ) ;
0750     const char* PIDXYZ = ssys::getenvvar("PIDXYZ") ;
0751     if(PIDXYZ && strcmp(PIDXYZ,"MIDDLE") == 0 )
0752     {
0753         LOG(info) << " special casing PIDXYZ MIDDLE " ;
0754         params->pidxyz.x = params->width/2 ;
0755         params->pidxyz.y = params->height/2 ;
0756         params->pidxyz.z = params->depth/2 ;
0757     }
0758
0759     LOG(LEVEL) << " params->pidxyz " << params->pidxyz ;
0760 }
0761
0762
0763 void CSGOptiX::setExternalDevicePixels(uchar4* _d_pixel )
0764 {
0765     framebuf->setExternalDevicePixels(_d_pixel) ;
0766     params->pixels = framebuf->d_pixel ;
0767 }
0768
0769
0770 void CSGOptiX::destroy()
0771 {
0772     LOG(LEVEL);
0773     delete sbt ;
0774     delete pip ;
0775 }
0776
0777
0778
0779
0780 /**
0781 CSGOptiX::simulate
0782 --------------------
0783
0784 NB the distinction between this and simulate_launch, this
0785 uses QSim::simulate to do genstep setup prior to calling
0786 CSGOptiX::simulate_launch via the SCSGOptiX.h protocol
0787
0788 The QSim::simulate argument reset:true is used in order
0789 to invoke SEvt::endOfEvent after the save, this is because
0790 at CSGOptiX level there is no need to allow the user to
0791 copy hits or other content from SEvt elsewhere.
0792
0793
0794
0795
0796
0797 **/
0798 double CSGOptiX::simulate(int eventID, bool reset)
0799 {
0800     assert(sim);
0801     double dt = sim->simulate(eventID, reset) ; // (QSim)
0802     return dt ;
0803 }
0804
0805
0806 void CSGOptiX::reset(int eventID)
0807 {
0808     assert(sim);
0809     sim->reset(eventID); // (QSim)
0810 }
0811
0812
0813
0814
0815 /**
0816 CSGOptiX::simulate
0817 -------------------
0818
0819 High level interface used by CSGOptiXService.h
0820
0821 **/
0822
0823 NP* CSGOptiX::simulate(const NP* gs, int eventID)
0824 {
0825     return sim->simulate(gs, eventID);
0826 }
0827
0828
0829
0830 /**
0831 CSGOptiX::simtrace
0832 --------------------
0833
0834 NB the distinction between this and simtrace_launch, this
0835 uses QSim::simtrace to do genstep setup prior to calling
0836 CSGOptiX::simtrace_launch via the SCSGOptiX.h protocol
0837
0838 Simtrace effectively always has reset:true because it
0839 always uses SEvt saving, unlike "simulate" which needs
0840 to support grabbing of hits into Geant4 collections.
0841
0842 **/
0843
0844 double CSGOptiX::simtrace(int eventID)
0845 {
0846     LOG(LEVEL) << "[" ;
0847     assert(sim);
0848     double dt = sim->simtrace(eventID) ;  // (QSim)
0849     LOG(LEVEL) << "] " << dt  ;
0850     return dt ;
0851 }
0852
0853
0854
0855 /**
0856 CSGOptiX::setFrame
0857 --------------------------
0858
0859 The no argument method uses MOI envvar or default of "-1"
0860
0861 For global geometry which typically uses default iidx of 0 there is special
0862 handling of iidx -1/-2/-3 implemented in CSGTarget::getCenterExtent
0863
0864
0865 iidx -2
0866     ordinary xyzw frame calulated by SCenterExtentFrame
0867
0868 iidx -3
0869     rtp tangential frame calulated by SCenterExtentFrame
0870
0871
0872 Setting CE center-extent establishes the coordinate system
0873 via calls to Composition::setCenterExtent which results in the
0874 definition of a model2world 4x4 matrix which becomes the frame of
0875 reference used by the EYE LOOK UP navigation controls.
0876
0877
0878 Q: CSGOptiX::setFrame is clearly needed for render but is it needed for simtrace, simulate ?
0879 A: Currently think that it is just a bookkeeping convenience for simtrace and not needed for simulate.
0880
0881    * not anymore, at SEvt level the frame is used for input photon targetting
0882
0883 **/
0884
0885 void CSGOptiX::setFrame()
0886 {
0887     assert(0 && " DONT USE THIS - DIRECTLY USE sglm INSTEAD"  );
0888     setFrame(ssys::getenvvar("MOI", "-1"));  // TODO: generalize to FRS
0889 }
0890
0891 void CSGOptiX::setFrame(const char* frs)
0892 {
0893     assert(0 && " DONT USE THIS - DIRECTLY USE sglm INSTEAD"  );
0894     LOG(LEVEL) << " frs " << frs ;
0895     sframe fr = foundry->getFrame(frs) ;
0896     sfr lfr = fr.spawn_lite();
0897     setFrame(lfr);
0898
0899 }
0900 void CSGOptiX::setFrame(const float4& ce )
0901 {
0902     assert(0 && " DONT USE THIS - DIRECTLY USE sglm INSTEAD"  );
0903     sfr lfr ;   // m2w w2m default to identity
0904
0905     lfr.ce.x = ce.x ;
0906     lfr.ce.y = ce.y ;
0907     lfr.ce.z = ce.z ;
0908     lfr.ce.w = ce.w ;
0909
0910     setFrame(lfr);
0911 }
0912
0913 /**
0914 CSGOptiX::setFrame into the SGLM.h instance
0915 ----------------------------------------------
0916
0917 Note that SEvt already holds an sframe used for input photon transformation,
0918 the sframe here is used for raytrace rendering.  Could perhaps rehome sglm
0919 into SEvt and use a single sframe for both input photon transformation
0920 and rendering ? But SGLM is for viz, so that is out of place.
0921
0922 Where to keep the frame?
0923
0924 * for simulation the obvious location is SEvt
0925 * for simtrace it needs to be in SEvt for genstep preparation
0926 * for rendering the obvious location is SGLM
0927
0928 Currently reluctant to depend on SEvt for rendering, because it
0929 brings complexity without much utility. So live with two locations
0930 for the frame.
0931
0932 **/
0933
0934 void CSGOptiX::setFrame(const sfr& lfr )
0935 {
0936     assert(0 && " DONT USE THIS - DIRECTLY USE sglm INSTEAD"  );
0937     sglm->set_frame(lfr);   // TODO: aim to remove sframe from sglm ? instead operate at ce (or sometimes m2w w2m level)
0938
0939     LOG(LEVEL) << "sglm.desc:" << std::endl << sglm->desc() ;
0940
0941     LOG(LEVEL) << lfr.desc() ;
0942
0943     LOG(LEVEL)
0944         << " sglm.TMIN " << sglm->TMIN
0945         << " sglm.tmin_abs " << sglm->tmin_abs()
0946         ;
0947
0948     LOG(LEVEL) << "]" ;
0949 }
0950
0951
0952
0953
0954
0955
0956
0957 void CSGOptiX::prepareParamRender()
0958 {
0959     int prepareParamRender_DEBUG = ssys::getenvint(_prepareParamRender_DEBUG, 0) ;
0960
0961     float extent = sglm->fr.ce.w ;
0962     const glm::vec3& eye = sglm->e ;
0963     const glm::vec3& U = sglm->u ;
0964     const glm::vec3& V = sglm->v ;
0965     const glm::vec3& W = sglm->w ;
0966     const glm::vec3& WNORM = sglm->wnorm ;
0967     const glm::vec4& ZPROJ = sglm->zproj ;
0968
0969     float tmin = sglm->get_near_abs() ;
0970     float tmax = sglm->get_far_abs() ;
0971     unsigned vizmask = sglm->vizmask ;
0972     unsigned cameratype = sglm->cam ;
0973     int traceyflip = sglm->traceyflip ;
0974     int rendertype = sglm->rendertype ;
0975     float length = 0.f ;
0976
0977     LOG_IF(info, prepareParamRender_DEBUG > 0 && kernel_count == 0)
0978         << _prepareParamRender_DEBUG << ":" << prepareParamRender_DEBUG
0979         << std::endl
0980         << std::setw(20) << " kernel_count " << kernel_count << std::endl
0981         << std::setw(20) << " extent "     << extent << std::endl
0982         << std::setw(20) << " sglm.fr.ce.w "  << sglm->fr.ce.w << std::endl
0983         << std::setw(20) << " sglm.getGazeLength "  << sglm->getGazeLength()  << std::endl
0984         << std::setw(20) << " comp.length"   << length
0985         << std::endl
0986         << std::setw(20) << " tmin "       << tmin  << std::endl
0987         << std::setw(20) << " tmax "       << tmax  << std::endl
0988         << std::setw(20) << " vizmask "    << vizmask  << std::endl
0989         << std::endl
0990         << std::setw(20) << " sglm.near "  << sglm->near  << std::endl
0991         << std::setw(20) << " sglm.get_near_abs "  << sglm->get_near_abs()  << std::endl
0992         << std::endl
0993         << std::setw(20) << " sglm.far "   << sglm->far  << std::endl
0994         << std::setw(20) << " sglm.get_far_abs "   << sglm->get_far_abs()  << std::endl
0995         << std::endl
0996         << std::setw(25) << " sglm.get_nearfar_basis "  << sglm->get_nearfar_basis()
0997         << std::setw(25) << " sglm.get_nearfar_mode "  << sglm->get_nearfar_mode()
0998         << std::endl
0999         << std::setw(25) << " sglm.get_focal_basis "  << sglm->get_focal_basis()
1000         << std::setw(25) << " sglm.get_focal_mode "  << sglm->get_focal_mode()
1001         << std::endl
1002         << std::setw(20) << " eye ("       << eye.x << " " << eye.y << " " << eye.z << " ) " << std::endl
1003         << std::setw(20) << " U ("         << U.x << " " << U.y << " " << U.z << " ) " << std::endl
1004         << std::setw(20) << " V ("         << V.x << " " << V.y << " " << V.z << " ) " << std::endl
1005         << std::setw(20) << " W ("         << W.x << " " << W.y << " " << W.z << " ) " << std::endl
1006         << std::setw(20) << " WNORM ("     << WNORM.x << " " << WNORM.y << " " << WNORM.z << " ) " << std::endl
1007         << std::endl
1008         << std::setw(20) << " cameratype " << cameratype << " "           << SCAM::Name(cameratype) << std::endl
1009         << std::setw(20) << " traceyflip " << traceyflip << std::endl
1010         << std::setw(20) << " sglm.cam " << sglm->cam << " " << SCAM::Name(sglm->cam) << std::endl
1011         << std::setw(20) << " ZPROJ ("     << ZPROJ.x << " " << ZPROJ.y << " " << ZPROJ.z << " " << ZPROJ.w << " ) " << std::endl
1012         << std::endl
1013         << "SGLM::DescEyeBasis (sglm->e,w,v,w)\n"
1014         << SGLM::DescEyeBasis( sglm->e, sglm->u, sglm->v, sglm->w )
1015         << std::endl
1016         << std::endl
1017         <<  "sglm.descEyeBasis\n"
1018         << sglm->descEyeBasis()
1019         << std::endl
1020         << "Composition basis  SGLM::DescEyeBasis( eye, U, V, W ) \n"
1021         << SGLM::DescEyeBasis( eye, U, V, W )
1022         << std::endl
1023         << "sglm.descELU \n"
1024         << sglm->descELU()
1025         << std::endl
1026         << std::endl
1027         << "sglm.descLog "
1028         << std::endl
1029         << sglm->descLog()
1030         << std::endl
1031         ;
1032
1033
1034
1035     params->setView(eye, U, V, W, WNORM );
1036     params->setCamera(tmin, tmax, cameratype, traceyflip, rendertype, ZPROJ );
1037     params->setVizmask(vizmask);
1038
1039     LOG(level) << std::endl << params->desc() ;
1040
1041     if(flight) return ;
1042
1043     LOG(level)
1044         << "sglm.desc " << std::endl
1045         << sglm->desc()
1046         ;
1047
1048 }
1049
1050
1051 /**
1052 CSGOptiX::prepareParamSimulate
1053 -------------------------------
1054
1055 Per-event simulate setup invoked just prior to optix launch
1056
1057 QSim::get_photon_slot_offset/QEvt::get_photon_slot_offset returns
1058
1059 **/
1060
1061 void CSGOptiX::prepareParamSimulate()
1062 {
1063     LOG(LEVEL);
1064     params->set_photon_slot_offset(sim->get_photon_slot_offset());
1065 }
1066
1067
1068
1069 /**
1070 CSGOptiX::prepareParam and upload
1071 -------------------------------------
1072
1073 This is invoked by CSGOptiX::launch just before the OptiX launch,
1074 depending on raygenmode the simulate/render param are prepared and uploaded.
1075
1076 Q: can Six use the same uploaded params ?
1077 A: not from device code it seems : only by using the hostside params to
1078    populate the pre-7 optix::context
1079
1080 **/
1081
1082 void CSGOptiX::prepareParam()
1083 {
1084     const glm::tvec4<double>& ce = sglm->fr.ce ;
1085
1086     params->setCenterExtent(ce.x, ce.y, ce.z, ce.w);
1087     switch(raygenmode)
1088     {
1089         case SRG_RENDER   : prepareParamRender()   ; break ;
1090         case SRG_SIMTRACE : prepareParamSimulate() ; break ;
1091         case SRG_SIMULATE : prepareParamSimulate() ; break ;
1092     }
1093
1094     params->upload();
1095     LOG_IF(level, !flight) << params->detail();
1096 }
1097
1098
1099 /**
1100 CSGOptiX::launch
1101 -------------------
1102
1103 For what happens next, see CSGOptiX7.cu::__raygen__rg OR CSGOptiX6.cu::raygen
1104 Depending on params.raygenmode the "render" or "simulate" method is called.
1105
1106 Formerly followed an OptiX 7 SDK example, creating a stream for the launch::
1107
1108     CUstream stream ;
1109     CUDA_CHECK( cudaStreamCreate( &stream ) );
1110     OPTIX_CHECK( optixLaunch( pip->pipeline, stream, d_param, sizeof( Params ), &(sbt->sbt), width, height, depth ) );
1111
1112 But that leaks 14kb for every launch, see:
1113
1114 * ~/opticks/notes/issues/okjob_GPU_memory_leak.rst
1115 * ~/opticks/CSGOptiX/cxs_min_igs.sh
1116
1117 Instead using default "stream=0" avoids the leak.
1118 Presumably that means every launch uses the same single default stream.
1119
1120 **/
1121
1122 double CSGOptiX::launch()
1123 {
1124     bool DEBUG_SKIP_LAUNCH = ssys::getenvbool("CSGOptiX__launch_DEBUG_SKIP_LAUNCH") ;
1125
1126     prepareParam();
1127     if(raygenmode != SRG_RENDER) assert(qev) ;
1128
1129     unsigned width = 0 ;
1130     unsigned height = 0 ;
1131     unsigned depth  = 0 ;
1132     switch(raygenmode)
1133     {
1134         case SRG_RENDER:    { width = params->width         ; height = params->height ; depth = params->depth ; } ; break ;
1135         case SRG_SIMTRACE:  { width = qev->getNumSimtrace() ; height = 1              ; depth = 1             ; } ; break ;
1136         case SRG_SIMULATE:  { width = qev->getNumPhoton()   ; height = 1              ; depth = 1             ; } ; break ;
1137     }
1138
1139     bool expect = width > 0 ;
1140     LOG_IF(fatal, !expect)
1141         << " qev.getNumSimtrace " << ( qev ? qev->getNumSimtrace() : -1 )
1142         << " qev.getNumPhoton   " << ( qev ? qev->getNumPhoton() : -1 )
1143         << " width " << width
1144         << " height " << height
1145         << " depth " << depth
1146         << " expect " << ( expect ? "YES" : "NO " )
1147         ;
1148
1149     assert(expect );
1150
1151     typedef std::chrono::time_point<std::chrono::high_resolution_clock> TP ;
1152     typedef std::chrono::duration<double> DT ;
1153     TP _t0 = std::chrono::high_resolution_clock::now();
1154     int64_t t0 = sstamp::Now();
1155
1156     LOG(LEVEL)
1157          << " raygenmode " << raygenmode
1158          << " SRG::Name(raygenmode) " << SRG::Name(raygenmode)
1159          << " width " << width
1160          << " height " << height
1161          << " depth " << depth
1162          << " DEBUG_SKIP_LAUNCH " << ( DEBUG_SKIP_LAUNCH ? "YES" : "NO " )
1163          ;
1164
1165     if(DEBUG_SKIP_LAUNCH == false)
1166     {
1167         CUdeviceptr d_param = (CUdeviceptr)Params::d_param ; ;
1168         assert( d_param && "must alloc and upload params before launch");
1169
1170         //cudaStream_t stream = SMgr::Stream();
1171         cudaStream_t stream = 0 ; // default stream
1172         OPTIX_CHECK( optixLaunch( pip->pipeline, (CUstream)stream, d_param, sizeof( Params ), &(sbt->sbt), width, height, depth ) );
1173
1174         CUDA_SYNC_CHECK();
1175         // see CSG/CUDA_CHECK.h the CUDA_SYNC_CHECK does cudaDeviceSyncronize
1176         // THIS LIKELY HAS LARGE PERFORMANCE IMPLICATIONS : BUT NOT EASY TO AVOID (MULTI-BUFFERING ETC..)
1177         kernel_count += 1 ;
1178     }
1179
1180
1181     TP _t1 = std::chrono::high_resolution_clock::now();
1182     DT _dt = _t1 - _t0;
1183
1184     int64_t t1 = sstamp::Now();
1185     int64_t dt = t1 - t0 ;
1186
1187     kernel_dt = _dt.count() ;   // formerly mis-named launch_dt
1188     kernel_times.push_back(kernel_dt);
1189
1190     kernel_times_.push_back(dt);
1191
1192     LOG(LEVEL)
1193           << " (params.width, params.height, params.depth) ( "
1194           << params->width << "," << params->height << "," << params->depth << ")"
1195           << std::fixed << std::setw(7) << std::setprecision(4) << kernel_dt
1196           ;
1197     return kernel_dt ;
1198 }
1199
1200
1201
1202 /**
1203 CSGOptiX::render_launch CSGOptiX::simtrace_launch CSGOptiX::simulate_launch
1204 --------------------------------------------------------------------------------
1205
1206 All the launch set (double)dt
1207
1208 CAUTION : *simulate_launch* and *simtrace_launch*
1209 MUST BE invoked from QSim::simulate and QSim::simtrace using the SCSGOptiX.h protocol.
1210 This is because genstep preparations are needed prior to launch.
1211
1212 These three methods currently all call *CSGOptiX::launch*
1213 with params.raygenmode switch function inside OptiX7Test.cu:__raygen__rg
1214 As it is likely better to instead have multiple raygen entry points
1215 are retaining the distinct methods up here.
1216
1217 *render* is also still needed to fulfil SRenderer protocol base
1218
1219 **/
1220 double CSGOptiX::render_launch()
1221 {
1222     assert(raygenmode == SRG_RENDER) ;
1223     return launch() ;
1224 }
1225 double CSGOptiX::simtrace_launch()
1226 {
1227     assert(raygenmode == SRG_SIMTRACE) ;
1228     return launch() ;
1229 }
1230 double CSGOptiX::simulate_launch()
1231 {
1232     assert(raygenmode == SRG_SIMULATE) ;
1233     return launch()  ;
1234 }
1235
1236 const CSGFoundry* CSGOptiX::getFoundry() const
1237 {
1238     return foundry ;
1239 }
1240
1241 std::string CSGOptiX::AnnotationTime( double dt, const char* extra )  // static
1242 {
1243     std::stringstream ss ;
1244     ss << std::fixed << std::setw(10) << std::setprecision(4) << dt ;
1245     if(extra) ss << " " << extra << " " ;
1246     std::string str = ss.str();
1247     return str ;
1248 }
1249 std::string CSGOptiX::Annotation( double dt, const char* bot_line, const char* extra )  // static
1250 {
1251     std::stringstream ss ;
1252     ss << AnnotationTime(dt, extra) ;
1253     if(bot_line) ss << std::setw(30) << " " << bot_line ;
1254     std::string str = ss.str();
1255     return str ;
1256 }
1257
1258 const char* CSGOptiX::getDefaultSnapPath() const
1259 {
1260     assert( foundry );
1261     const char* cfbase = foundry->getOriginCFBase();
1262     assert( cfbase );
1263     const char* path = SPath::Resolve(cfbase, "CSGOptiX/snap.jpg" , FILEPATH );
1264     return path ;
1265 }
1266
1267
1268
1269
1270 /**
1271 CSGOptiX::getRenderStemDefault
1272 --------------------------------
1273
1274 Example NAMEPREFIX from cxr_min.sh::
1275
1276    cxr_min__eye_0,0.8,0__zoom_1.0__tmin_0.1_
1277
1278 Example resulting stem::
1279
1280    cxr_min__eye_0,0.8,0__zoom_1.0__tmin_0.1__ALL
1281
1282 WIP: interactive view control makes this approach obsolete
1283 To provide dynamic naming have added::
1284
1285     ssys::setenvctx
1286     ssys::setenvmap
1287
1288 The idea being to generate a map<string,string> of
1289 view params for feeding into the environment that
1290 is used in the resolution of a name pattern that is
1291 specied from bash, ie::
1292
1293     cxr_min_eye_${EYE}__zoom_${ZOOM}__tmin_${TMIN}
1294
1295 HMM: SGLM seems more appropriate place to do this than here
1296
1297 **/
1298 const char* CSGOptiX::getRenderStemDefault() const
1299 {
1300     const std::string& fr_name = sglm->fr.get_name() ;
1301
1302     std::stringstream ss ;
1303     ss << ssys::getenvvar("NAMEPREFIX","nonamepfx") ;
1304     ss << "_" ;
1305     ss << ( fr_name.empty() ? "no_frame_name" : fr_name ) ;
1306
1307     std::string str = ss.str();
1308     return strdup(str.c_str());
1309 }
1310
1311
1312
1313
1314 bool CSGOptiX::handle_snap(int wanted_snap)
1315 {
1316     bool can_handle = wanted_snap == 1 || wanted_snap == 2 ;
1317     if(!can_handle) return false ;
1318     switch(wanted_snap)
1319     {
1320         case 1: render_save()          ; break ;
1321         case 2: render_save_inverted() ; break ;
1322     }
1323     return true ;
1324 }
1325
1326
1327 /**
1328 CSGOptiX::render (formerly render_snap)
1329 -------------------------------------------
1330 **/
1331
1332 double CSGOptiX::render( const char* stem_ )
1333 {
1334     render_launch();
1335     render_save(stem_);
1336     return kernel_dt ;
1337 }
1338
1339
1340
1341 /**
1342 CSGOptiX::render_save
1343 ----------------------
1344
1345 TODO: update file naming impl, currently using old inflexible approach
1346
1347 **/
1348 void CSGOptiX::render_save(const char* stem_)
1349 {
1350     render_save_(stem_, false);
1351 }
1352 void CSGOptiX::render_save_inverted(const char* stem_)
1353 {
1354     render_save_(stem_, true);
1355 }
1356
1357
1358 void CSGOptiX::render_save_(const char* stem_, bool inverted)
1359 {
1360     const char* outdir = SEventConfig::OutDir();
1361     const char* stem = stem_ ? stem_ : getRenderStemDefault() ;  // without ext
1362
1363     bool unique = true ;
1364     const char* outpath = SEventConfig::OutPath(stem, -1, ".jpg", unique );
1365
1366     LOG(LEVEL)
1367           << SEventConfig::DescOutPath(stem, -1, ".jpg", unique );
1368           ;
1369
1370     std::string u_outdir ;
1371     std::string u_stem ;
1372     std::string u_ext ;
1373
1374     [[maybe_unused]] int rc = spath::SplitExt( u_outdir, u_stem, u_ext, outpath )  ;
1375     assert(rc == 0 );
1376
1377     sglm->addlog("CSGOptiX::render_snap", u_stem.c_str() );
1378
1379
1380     const char* topline = ssys::getenvvar("TOPLINE", sproc::ExecutableName() );
1381     std::string _extra = SEventConfig::GetGPUMeta();  // scontext::brief giving GPU name
1382     const char* extra = strdup(_extra.c_str()) ;
1383
1384     const char* botline_ = ssys::getenvvar("BOTLINE", nullptr );
1385     std::string bottom_line = CSGOptiX::Annotation(kernel_dt, botline_, extra );
1386     const char* botline = bottom_line.c_str() ;
1387
1388
1389     LOG(LEVEL)
1390           << " stem " << stem
1391           << " outpath " << outpath
1392           << " outdir " << ( outdir ? outdir : "-" )
1393           << " kernel_dt " << kernel_dt
1394           << " topline [" <<  topline << "]"
1395           << " botline [" <<  botline << "]"
1396           ;
1397
1398     LOG(info) << outpath  << " : " << AnnotationTime(kernel_dt, extra)  ;
1399
1400     unsigned line_height = 24 ;
1401     snap(outpath, botline, topline, line_height, inverted  );
1402
1403
1404     sglm->save( u_outdir.c_str(), u_stem.c_str() );
1405 }
1406
1407
1408 /**
1409 CSGOptiX::snap : Download frame pixels and write to file as jpg.
1410 ------------------------------------------------------------------
1411
1412 WIP: contrast this with SGLFW::snap_local and consider if more consolidation is possible
1413
1414
1415 SGLFW::snap_local
1416     Uses OpenGL glReadPixels to download pixels and write them to file
1417
1418 CSGOptiX::snap
1419     OptiX/CUDA level download ray traced pixels and write to file
1420
1421 **/
1422
1423 void CSGOptiX::snap(const char* path_, const char* bottom_line, const char* top_line, unsigned line_height, bool inverted )
1424 {
1425     const char* path = path_ ? SPath::Resolve(path_, FILEPATH ) : getDefaultSnapPath() ;
1426     LOG(LEVEL) << " path " << path ;
1427
1428 #if OPTIX_VERSION < 70000
1429     const char* top_extra = nullptr ;
1430 #else
1431     const char* top_extra = pip->desc();
1432 #endif
1433     const char* topline = SStr::Concat(top_line, top_extra);
1434
1435     LOG(LEVEL) << " path_ [" << path_ << "]" ;
1436     LOG(LEVEL) << " topline " << topline  ;
1437
1438     LOG(LEVEL) << "[ frame.download " ;
1439     if( inverted == false )
1440     {
1441         framebuf->download();
1442     }
1443     else
1444     {
1445         framebuf->download_inverted();
1446     }
1447     LOG(LEVEL) << "] frame.download " ;
1448
1449     LOG(LEVEL) << "[ frame.annotate " ;
1450     framebuf->annotate( bottom_line, topline, line_height );
1451     LOG(LEVEL) << "] frame.annotate " ;
1452
1453     LOG(LEVEL) << "[ frame.snap " ;
1454     framebuf->snap( path  );
1455     LOG(LEVEL) << "] frame.snap " ;
1456
1457     if(!flight || SStr::Contains(path,"00000"))
1458     {
1459         saveMeta(path);
1460     }
1461 }
1462
1463 #ifdef WITH_FRAME_PHOTON
1464 void CSGOptiX::writeFramePhoton(const char* dir, const char* name)
1465 {
1466 #if OPTIX_VERSION < 70000
1467     assert(0 && "not implemented pre-7");
1468 #else
1469     framebuf->writePhoton(dir, name);
1470 #endif
1471 }
1472 #endif
1473
1474
1475 int CSGOptiX::render_flightpath() // for making mp4 movies
1476 {
1477     LOG(fatal) << "flightpath rendering not yet implemented in now default SGLM branch " ;
1478     return 1 ;
1479 }
1480
1481 void CSGOptiX::saveMeta(const char* jpg_path) const
1482 {
1483     const char* json_path = SStr::ReplaceEnd(jpg_path, ".jpg", ".json");
1484     LOG(LEVEL) << "[ json_path " << json_path  ;
1485
1486     nlohmann::json& js = meta->js ;
1487     js["jpg"] = jpg_path ;
1488     js["emm"] = SGeoConfig::EnabledMergedMesh() ;
1489
1490     if(foundry->hasMeta())
1491     {
1492         js["cfmeta"] = foundry->meta ;
1493     }
1494
1495     std::string extra = SEventConfig::GetGPUMeta();
1496     js["scontext"] = extra.empty() ? "-" : strdup(extra.c_str()) ;
1497
1498     const std::vector<double>& t = kernel_times ;
1499     if( t.size() > 0 )
1500     {
1501         double mn, mx, av ;
1502         SVec<double>::MinMaxAvg(t,mn,mx,av);
1503
1504         js["mn"] = mn ;
1505         js["mx"] = mx ;
1506         js["av"] = av ;
1507     }
1508
1509     meta->save(json_path);
1510     LOG(LEVEL) << "] json_path " << json_path  ;
1511 }
1512
1513
1514
1515 void CSGOptiX::write_Ctx_log(const char* dir) const
1516 {
1517 #if OPTIX_VERSION < 70000
1518 #else
1519     std::string ctxlog = Ctx::GetLOG() ;
1520     spath::Write(ctxlog.c_str() , dir, CTX_LOGNAME  );
1521 #endif
1522 }
1523
1524
1525 /**
1526 CSGOptiX::_OPTIX_VERSION
1527 -------------------------
1528
1529 This depends on the the optix.h header only which provides the OPTIX_VERSION macro
1530 SEventConfig::so it could be done at the lowest level, no need for it to be
1531 up at this "elevation"
1532
1533 TODO: relocate to OKConf or SysRap, BUT this must wait until switch to full proj 7
1534
1535 **/
1536
1537 #define xstr(s) str(s)
1538 #define str(s) #s
1539
1540 int CSGOptiX::_OPTIX_VERSION()   // static
1541 {
1542     char vers[16] ;
1543     snprintf(vers, 16, "%s",xstr(OPTIX_VERSION));
1544     return std::atoi(vers) ;
1545 }