Back to home page

EIC code displayed by LXR
 
 

    

File indexing completed on 2026-04-09 07:49:03 

0001 #include <iostream>
0002 #include <iomanip>
0003 #include <cstring>
0004 #include <csignal>
0005 #include <sstream>
0006 #include <fstream>
0007 
0008 #include <optix.h>
0009 #include <optix_stubs.h>
0010 #include <cuda_runtime.h>
0011 
0012 #include "SGeoConfig.hh"
0013 #include "SSys.hh"
0014 #include "SVec.hh"
0015 #include "scuda.h"
0016 #include "SScene.h"
0017 #include "NPX.h"
0018 
0019 #include "OPTIX_CHECK.h"
0020 #include "CUDA_CHECK.h"
0021 
0022 #include "CSGFoundry.h"
0023 #include "CSGSolid.h"
0024 #include "CSGNode.h"
0025 
0026 #include "Binding.h"
0027 #include "Params.h"
0028 #include "Ctx.h"
0029 
0030 #include "PIP.h"
0031 #include "SBT.h"
0032 #include "Properties.h"
0033 
0034 #include "CU.h"
0035 #include "SLOG.hh"
0036 
0037 #ifdef WITH_SOPTIX_ACCEL
0038 #include "SOPTIX_Accel.h"
0039 #include "SOPTIX_BuildInput_CPA.h"
0040 #include "SOPTIX_BuildInput_IA.h"
0041 #include "SOPTIX_BuildInput_Mesh.h"
0042 #include "SOPTIX_MeshGroup.h"
0043 #else
0044 #include "GAS.h"
0045 #include "GAS_Builder.h"
0046 #include "IAS.h"
0047 #include "IAS_Builder.h"
0048 #endif
0049 
0050 
0051 
0052 /**
0053 SBT
0054 ====
0055 
0056 SBT needs PIP as the packing of SBT record headers requires
0057 access to their corresponding program groups (PGs).
0058 This is one aspect of establishing the connection between the
0059 PGs and their data.
0060 
0061 **/
0062 
0063 const plog::Severity SBT::LEVEL = SLOG::EnvLevel("SBT", "DEBUG");
0064 
0065 
0066 std::string SBT::Desc()  // static
0067 {
0068     std::stringstream ss ;
0069     ss << "SBT::Desc"
0070 #ifdef WITH_SOPTIX_ACCEL
0071        << " WITH_SOPTIX_ACCEL"
0072 #else
0073        << " NOT:WITH_SOPTIX_ACCEL"
0074 #endif
0075        ;
0076     std::string str = ss.str();
0077     return str ;
0078 }
0079 
0080 
0081 
0082 SBT::SBT(const PIP* pip_)
0083     :
0084     emm(SGeoConfig::EnabledMergedMesh()),
0085     pip(pip_),
0086     properties(pip->properties),
0087     raygen(nullptr),
0088     miss(nullptr),
0089     hitgroup(nullptr),
0090     check(nullptr),
0091     foundry(nullptr),
0092     scene(nullptr)
0093 {
0094     init();
0095 }
0096 
0097 SBT::~SBT()
0098 {
0099     destroy();
0100 }
0101 
0102 
0103 void SBT::init()
0104 {
0105     LOG(LEVEL) << "[" ;
0106     createRaygen();
0107     updateRaygen();
0108     createMiss();
0109     updateMiss();
0110     LOG(LEVEL) << "]" ;
0111 }
0112 
0113 
0114 void SBT::destroy()
0115 {
0116     destroyRaygen();
0117     destroyMiss();
0118     destroyHitgroup();
0119 }
0120 
0121 
0122 /**
0123 SBT::createRaygen
0124 ------------------
0125 
0126 Raygen is typedef to SbtRecord<RaygenData>
0127 so this is setting up access to raygen data : but that
0128 is just a placeholder with most everything coming from params
0129 **/
0130 
0131 void SBT::createRaygen()
0132 {
0133     raygen = new Raygen ;
0134     CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_raygen ),   sizeof(Raygen) ) );
0135     sbt.raygenRecord = d_raygen;
0136     OPTIX_CHECK( optixSbtRecordPackHeader( pip->raygen_pg,   raygen ) );
0137 }
0138 
0139 void SBT::destroyRaygen()
0140 {
0141     CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_raygen ) ) );
0142 }
0143 
0144 void SBT::updateRaygen()
0145 {
0146     raygen->data = {};
0147     raygen->data.placeholder = 42.0f ;
0148 
0149     CUDA_CHECK( cudaMemcpy(
0150                 reinterpret_cast<void*>( d_raygen ),
0151                 raygen,
0152                 sizeof( Raygen ),
0153                 cudaMemcpyHostToDevice
0154                 ) );
0155 }
0156 
0157 
0158 /**
0159 SBT::createMiss
0160 --------------------
0161 
0162 NB the records have opaque header and user data
0163 **/
0164 
0165 void SBT::createMiss()
0166 {
0167     miss = new Miss ;
0168     CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_miss ), sizeof(Miss) ) );
0169     sbt.missRecordBase = d_miss;
0170     OPTIX_CHECK( optixSbtRecordPackHeader( pip->miss_pg, miss ) );
0171 
0172     sbt.missRecordStrideInBytes     = sizeof( Miss );
0173     sbt.missRecordCount             = 1;
0174 }
0175 
0176 void SBT::destroyMiss()
0177 {
0178     CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_miss ) ) );
0179 }
0180 
0181 void SBT::updateMiss()
0182 {
0183     //float3 purple = make_float3(0.3f, 0.1f, 0.5f);
0184     //float3 white = make_float3( 1.0f, 1.0f, 1.0f);
0185     //float3 lightgrey = make_float3( 0.9f, 0.9f, 0.9f);
0186     float3 midgrey = make_float3( 0.6f, 0.6f, 0.6f);
0187     const float3& bkg = midgrey  ;
0188 
0189     miss->data.r = bkg.x ;
0190     miss->data.g = bkg.y ;
0191     miss->data.b = bkg.z ;
0192 
0193     CUDA_CHECK( cudaMemcpy(
0194                 reinterpret_cast<void*>( d_miss ),
0195                 miss,
0196                 sizeof(Miss),
0197                 cudaMemcpyHostToDevice
0198                 ) );
0199 }
0200 
0201 
0202 
0203 
0204 
0205 
0206 
0207 
0208 
0209 
0210 
0211 
0212 
0213 
0214 
0215 /**
0216 SBT::setFoundry
0217 ------------------
0218 
0219 Canonical invokation from CSGOptiX::CSGOptiX//CSGOptiX::initGeometry
0220 
0221 1. creates GAS using aabb obtained via geo
0222 2. creates IAS
0223 3. creates Hitgroup SBT records
0224 
0225 **/
0226 
0227 void SBT::setFoundry(const CSGFoundry* foundry_)
0228 {
0229     foundry = foundry_ ;          // analytic
0230     scene = foundry->getScene();  // triangulated
0231 
0232     createGeom();
0233 }
0234 
0235 /**
0236 SBT::createGeom
0237 -----------------
0238 
0239 createGAS
0240     SCSGPrimSpec for each compound solid are converted to GAS and collected into map
0241 createIAS
0242     instance transforms with compound solid references are converted into the IAS
0243 createHitgroup
0244     bringing it all together
0245 
0246 **/
0247 void SBT::createGeom()
0248 {
0249     LOG(LEVEL) << "[" ;
0250     createGAS();
0251     LOG(LEVEL) << "] createGAS " ;
0252     createIAS();
0253     LOG(LEVEL) << "] createIAS " ;
0254     createHitgroup();
0255     LOG(LEVEL) << "] createHitGroup " ;
0256     checkHitgroup();
0257     LOG(LEVEL) << "] checkHitGroup " ;
0258     LOG(LEVEL) << "]" ;
0259 }
0260 
0261 
0262 
0263 
0264 /**
0265 SBT::createGAS
0266 ----------------
0267 
0268 For each compound shape the aabb of each prim (aka layer) is
0269 uploaded to GPU in order to create GAS for each compound shape.
0270 
0271 Note that the prim could be a CSG tree of constituent nodes each
0272 with their own aabb, but only one aabb corresponding to the overall
0273 prim extent is used.
0274 
0275 **/
0276 
0277 void SBT::createGAS()
0278 {
0279     LOG(LEVEL) << SGeoConfig::DescEMM() ;
0280 
0281     unsigned num_solid = foundry->getNumSolid();   // STANDARD_SOLID
0282     for(unsigned i=0 ; i < num_solid ; i++)
0283     {
0284         unsigned gas_idx = i ;
0285 
0286         bool enabled = SGeoConfig::IsEnabledMergedMesh(gas_idx) ;
0287         bool enabled2 = emm & ( 0x1 << gas_idx ) ;
0288         bool enabled_expect = enabled == enabled2 ;
0289         assert( enabled_expect );
0290         if(!enabled_expect) std::raise(SIGINT);
0291 
0292         if( enabled )
0293         {
0294             LOG(LEVEL) << " emm proceed " << gas_idx ;
0295             createGAS(gas_idx);
0296         }
0297         else
0298         {
0299             LOG(LEVEL) << " emm skip " << gas_idx ;
0300         }
0301     }
0302     LOG(LEVEL) << descGAS() ;
0303 }
0304 
0305 
0306 /**
0307 SBT::createGAS
0308 ---------------
0309 
0310 1. gets SCSGPrimSpec for a the *gas_idx* compound solid from foundry
0311 2. converts the SCSGPrimSpec into a GAS, passing in bbox device array pointers
0312 3. inserts gas into vgas map using *gas_idx* key
0313 
0314 **/
0315 
0316 #ifdef WITH_SOPTIX_ACCEL
0317 void SBT::createGAS(unsigned gas_idx)
0318 {
0319     SOPTIX_BuildInput* bi = nullptr ;
0320     SOPTIX_Accel* gas = nullptr ;
0321 
0322     bool trimesh = foundry->isSolidTrimesh(gas_idx); // now based on forced triangulation config
0323 
0324     const std::string& mmlabel = foundry->getSolidMMLabel(gas_idx);
0325 
0326     LOG(LEVEL)
0327         << " WITH_SOPTIX_ACCEL "
0328         << " gas_idx " << gas_idx
0329         << " trimesh " << ( trimesh ? "YES" : "NO " )
0330         << " mmlabel " << mmlabel
0331         ;
0332 
0333     if(trimesh)
0334     {
0335         // note similarity to SOPTIX_Scene::init_GAS
0336         const SMeshGroup* mg = scene->getMeshGroup(gas_idx) ;
0337         LOG_IF(fatal, mg == nullptr)
0338             << " FAILED to SScene::getMeshGroup"
0339             << " gas_idx " << gas_idx
0340             << "\n"
0341             << scene->desc()
0342             ;
0343         assert(mg);
0344 
0345         SOPTIX_MeshGroup* xmg = SOPTIX_MeshGroup::Create( mg ) ;
0346         gas = SOPTIX_Accel::Create(Ctx::context, xmg->bis );
0347         xgas[gas_idx] = xmg ;
0348     }
0349     else
0350     {
0351         // analytic geometry
0352         SCSGPrimSpec ps = foundry->getPrimSpec(gas_idx);
0353         bi = new SOPTIX_BuildInput_CPA(ps) ;
0354         gas = SOPTIX_Accel::Create(Ctx::context, bi );
0355     }
0356     vgas[gas_idx] = gas ;
0357 }
0358 
0359 #else
0360 void SBT::createGAS(unsigned gas_idx)
0361 {
0362     bool trimesh = foundry->isSolidTrimesh(gas_idx);
0363     LOG(fatal, trimesh == true ) << " NOT:WITH_SOPTIX_ACCEL ONLY SUPPORTS ANALYTIC GEOMETRY : INVALID SGeoConfig::Trimesh setting " ;
0364     assert( trimesh == false );
0365 
0366     SCSGPrimSpec ps = foundry->getPrimSpec(gas_idx);
0367     GAS gas = {} ;
0368     GAS_Builder::Build(gas, ps);
0369     vgas[gas_idx] = gas ;
0370 }
0371 #endif
0372 
0373 
0374 
0375 OptixTraversableHandle SBT::getGASHandle(unsigned gas_idx) const
0376 {
0377     unsigned count = vgas.count(gas_idx);
0378     LOG_IF(fatal, count == 0) << " no such gas_idx " << gas_idx ;
0379     assert( count == 1 );
0380 
0381 #ifdef WITH_SOPTIX_ACCEL
0382     SOPTIX_Accel* _gas = vgas.at(gas_idx) ;
0383     OptixTraversableHandle handle = _gas->handle ;
0384 #else
0385     const GAS& gas = vgas.at(gas_idx);
0386     OptixTraversableHandle handle = gas.handle ;
0387 #endif
0388 
0389     return handle ;
0390 }
0391 
0392 
0393 void SBT::createIAS()
0394 {
0395     unsigned num_ias = foundry->getNumUniqueIAS() ;
0396     bool num_ias_expect = num_ias == 1 ;
0397     assert( num_ias_expect );
0398     if(!num_ias_expect) std::raise(SIGINT);
0399 
0400     unsigned ias_idx = 0 ;
0401     createIAS(ias_idx);
0402 }
0403 
0404 /**
0405 SBT::createIAS
0406 ----------------
0407 
0408 Hmm: usually only one IAS.
0409 
0410 2024-04-30 11:08:33.056 INFO  [65240] [SBT::collectInstances@468] ] instances.size 47887
0411 2024-04-30 11:08:33.056 INFO  [65240] [SBT::createIAS@372] SBT::descIAS inst.size 47887 SBT_DUMP_IAS 0
0412  gas_idx          0 num_ins_idx          1 ins_idx_mn          0 ins_idx_mx          0 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)          1
0413  gas_idx          1 num_ins_idx      25600 ins_idx_mn          1 ins_idx_mx      25600 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)      25600
0414  gas_idx          2 num_ins_idx      12615 ins_idx_mn      25601 ins_idx_mx      38215 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)      12615
0415  gas_idx          3 num_ins_idx       4997 ins_idx_mn      38216 ins_idx_mx      43212 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)       4997
0416  gas_idx          4 num_ins_idx       2400 ins_idx_mn      43213 ins_idx_mx      45612 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)       2400
0417  gas_idx          5 num_ins_idx        590 ins_idx_mn      45613 ins_idx_mx      46202 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)        590
0418  gas_idx          6 num_ins_idx        590 ins_idx_mn      46203 ins_idx_mx      46792 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)        590
0419  gas_idx          7 num_ins_idx        590 ins_idx_mn      46793 ins_idx_mx      47382 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)        590
0420  gas_idx          8 num_ins_idx        504 ins_idx_mn      47383 ins_idx_mx      47886 ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2)        504
0421 
0422 
0423 
0424 
0425 **/
0426 
0427 void SBT::createIAS(unsigned ias_idx)
0428 {
0429     unsigned num_inst = foundry->getNumInst();
0430     unsigned num_ias_inst = foundry->getNumInstancesIAS(ias_idx, emm);
0431     LOG(LEVEL)
0432         << " ias_idx " << ias_idx
0433         << " num_inst " << num_inst
0434         << " num_ias_inst(getNumInstancesIAS) " << num_ias_inst
0435         ;
0436 
0437     std::vector<qat4> inst ;
0438     foundry->getInstanceTransformsIAS(inst, ias_idx, emm );
0439     assert( num_ias_inst == inst.size() );
0440 
0441 
0442     collectInstances(inst);
0443 
0444     LOG(LEVEL) << descIAS(inst);
0445 
0446 #ifdef WITH_SOPTIX_ACCEL
0447     SOPTIX_BuildInput* ia = new SOPTIX_BuildInput_IA(instances) ;
0448     SOPTIX_Accel* ias = SOPTIX_Accel::Create(Ctx::context, ia );
0449     vias.push_back(ias);
0450 #else
0451     IAS ias = {} ;
0452     IAS_Builder::Build(ias, instances );
0453     vias.push_back(ias);
0454 #endif
0455 
0456 }
0457 
0458 
0459 
0460 
0461 
0462 /**
0463 SBT::collectInstances
0464 ----------------------
0465 
0466 Converts *ias_inst* a vector of qat4 geometry identity instrumented transforms into
0467 a vector of OptixInstance. The instance.sbtOffset are set using SBT::getOffset
0468 for the gas_idx and with prim_idx:0 indicating the outer prim(aka layer)
0469 of the GAS.
0470 
0471 Canonically invoked during CSGOptiX instanciation, from stack::
0472 
0473     CSGOptiX::CSGOptiX
0474     CSGOptiX::init
0475     CSGOptiX::initGeometry
0476     SBT::setFoundry
0477     SBT::createGeom
0478     SBT::createIAS
0479     SBT::collectInstances
0480 
0481 
0482 Collecting OptixInstance was taking 0.42s for 48477 inst,
0483 as SBT::getOffset was being called for every instance. Instead
0484 of doing this caching the result in the gasIdx_sbtOffset brings
0485 the time down to zero.
0486 
0487 HMM: Could make better use of instanceId, eg with bitpack gas_idx, ias_idx ?
0488 See note in InstanceId.h its not so easy due to bit limits.
0489 But it doesnt matter much as can just do lookups CPU side based
0490 on simple indices from GPU side.
0491 
0492 **/
0493 
0494 
0495 void SBT::collectInstances( const std::vector<qat4>& ias_inst )
0496 {
0497     LOG(LEVEL) << "[ ias_inst.size " << ias_inst.size() ;  // eg 48477
0498 
0499     unsigned num_ias_inst = ias_inst.size() ;
0500     unsigned flags = OPTIX_INSTANCE_FLAG_DISABLE_ANYHIT ;
0501     unsigned prim_idx = 0u ;  // need sbt offset for the outer prim(aka layer) of the GAS
0502 
0503     std::map<unsigned, unsigned> gasIdx_sbtOffset ;
0504 
0505     for(unsigned i=0 ; i < num_ias_inst ; i++)
0506     {
0507         const qat4& q = ias_inst[i] ;
0508         int ins_idx,  gasIdx, sensor_identifier, sensor_index ;
0509         q.getIdentity(ins_idx, gasIdx, sensor_identifier, sensor_index );
0510 
0511         unsigned instanceId = q.get_IAS_OptixInstance_instanceId() ;
0512         assert( int(instanceId) == sensor_identifier );
0513 
0514         bool instanceId_is_allowed = instanceId < properties->limitMaxInstanceId ;
0515         LOG_IF(fatal, !instanceId_is_allowed)
0516             << " instanceId " << instanceId
0517             << " sbt->properties->limitMaxInstanceId " << properties->limitMaxInstanceId
0518             << " instanceId_is_allowed " << ( instanceId_is_allowed ? "YES" : "NO " )
0519             ;
0520         assert( instanceId_is_allowed  ) ;
0521 
0522         OptixTraversableHandle handle = getGASHandle(gasIdx);
0523 
0524         bool found = gasIdx_sbtOffset.count(gasIdx) == 1 ;
0525         unsigned sbtOffset = found ? gasIdx_sbtOffset.at(gasIdx) : getOffset(gasIdx, prim_idx ) ;
0526         if(!found)
0527         {
0528             gasIdx_sbtOffset[gasIdx] = sbtOffset ;
0529             LOG(LEVEL)
0530                 << " i " << std::setw(7) << i
0531                 << " gasIdx " << std::setw(3) << gasIdx
0532                 << " sbtOffset " << std::setw(6) << sbtOffset
0533                 << " gasIdx_sbtOffset.size " << std::setw(3) << gasIdx_sbtOffset.size()
0534                 << " instanceId " << instanceId
0535                 ;
0536         }
0537 
0538         //unsigned visibilityMask = 255;  // cf SOPTIX_Scene::init_Instances
0539         unsigned visibilityMask = properties->visibilityMask(gasIdx);
0540 
0541         OptixInstance instance = {} ;
0542         q.copy_columns_3x4( instance.transform );
0543         instance.instanceId = instanceId ;
0544         instance.sbtOffset = sbtOffset ;
0545         instance.visibilityMask = visibilityMask ;
0546 
0547         instance.flags = flags ;
0548         instance.traversableHandle = handle ;
0549 
0550         instances.push_back(instance);
0551     }
0552     LOG(LEVEL) << "] instances.size " << instances.size() ;
0553 }
0554 
0555 NP* SBT::serializeInstances() const
0556 {
0557     return NPX::ArrayFromVec<unsigned, OptixInstance>(instances) ;
0558 }
0559 
0560 
0561 /**
0562 SBT::descIAS (actually descINST would be more appropriate)
0563 ------------------------------------------------------------
0564 
0565 1. traverse over *inst* collecting *ins_idx* for each gas into a map keyed on gas_idx *ins_idx_per_gas*
0566 2. emit description of that map
0567 
0568 **/
0569 
0570 std::string SBT::descIAS(const std::vector<qat4>& inst ) const
0571 {
0572     std::stringstream ss ;
0573     bool sbt_dump_ias = SSys::getenvbool("SBT_DUMP_IAS") ;
0574     ss
0575         << "SBT::descIAS"
0576         << " inst.size " << inst.size()
0577         << " SBT_DUMP_IAS " << sbt_dump_ias
0578         << std::endl
0579         ;
0580 
0581     typedef std::map<int, std::vector<int>> MUV ;
0582     MUV ins_idx_per_gas ;
0583 
0584     for(unsigned i=0 ; i < inst.size() ; i++)
0585     {
0586         const qat4& q = inst[i] ;
0587         int ins_idx,  gas_idx, sensor_identifier, sensor_index ;
0588         q.getIdentity(ins_idx,  gas_idx, sensor_identifier, sensor_index );
0589 
0590         ins_idx_per_gas[gas_idx].push_back(ins_idx);
0591 
0592         if(sbt_dump_ias) ss
0593            << " i "       << std::setw(10) << i
0594            << " ins_idx " << std::setw(10) << ins_idx
0595            << " gas_idx " << std::setw(10) << gas_idx
0596            << " sensor_identifier " << std::setw(10) << sensor_identifier
0597            << " sensor_index " << std::setw(10) << sensor_index
0598            << std::endl
0599            ;
0600     }
0601 
0602     MUV::const_iterator b = ins_idx_per_gas.begin();
0603     MUV::const_iterator e = ins_idx_per_gas.end();
0604     MUV::const_iterator i ;
0605 
0606     for( i=b ; i != e ; i++)
0607     {
0608         int gas_idx = i->first ;
0609         const std::vector<int>& v = i->second ;
0610         int num_ins_idx = int(v.size()) ;
0611 
0612         int ins_idx_mn, ins_idx_mx ;
0613         SVec<int>::MinMax(v, ins_idx_mn, ins_idx_mx)  ;
0614         int num_ins_idx2 = ins_idx_mx - ins_idx_mn + 1 ;
0615 
0616         ss
0617             << " gas_idx " << std::setw(10) <<  gas_idx
0618             << " num_ins_idx " << std::setw(10) << num_ins_idx
0619             << " ins_idx_mn " << std::setw(10) << ins_idx_mn
0620             << " ins_idx_mx " << std::setw(10) << ins_idx_mx
0621             << " ins_idx_mx - ins_idx_mx + 1 (num_ins_idx2) " << std::setw(10) << num_ins_idx2
0622             << std::endl
0623             ;
0624 
0625         assert( num_ins_idx == num_ins_idx2 );
0626     }
0627     std::string s = ss.str();
0628     return s ;
0629 }
0630 
0631 
0632 OptixTraversableHandle SBT::getIASHandle(unsigned ias_idx) const
0633 {
0634     assert( ias_idx < vias.size() );
0635 
0636 #ifdef WITH_SOPTIX_ACCEL
0637     SOPTIX_Accel* _ias = vias[ias_idx] ;
0638     OptixTraversableHandle handle = _ias->handle ;
0639 #else
0640     const IAS& ias = vias[ias_idx];
0641     OptixTraversableHandle handle = ias.handle ;
0642 #endif
0643     return handle ;
0644 }
0645 
0646 
0647 OptixTraversableHandle SBT::getTOPHandle() const
0648 {
0649     return getIASHandle(0);
0650 }
0651 
0652 
0653 
0654 /**
0655 SBT::getOffset
0656 ----------------
0657 
0658 Canonically invoked from both::
0659 
0660    SBT::collectInstances
0661    SBT::createHitgroup
0662 
0663 The q_layer_idx is 0-based within the q_gas_idx composite shape,
0664 ie it is local to the solid.
0665 
0666 **/
0667 
0668 int SBT::getOffset(unsigned q_gas_idx, unsigned q_layer_idx ) const
0669 {
0670     int offset_sbt = _getOffset(q_gas_idx, q_layer_idx );
0671 
0672     LOG_IF(LEVEL, q_layer_idx < 10)
0673         << " q_gas_idx " << q_gas_idx
0674         << " q_layer_idx " << q_layer_idx
0675         << " offset_sbt " << offset_sbt
0676         ;
0677 
0678     assert( offset_sbt > -1 );
0679     return offset_sbt ;
0680 }
0681 
0682 /**
0683 SBT::_getOffset
0684 ----------------
0685 
0686 Implemented as an inner method avoiding "goto"
0687 to break out of multiple for loops.
0688 
0689 Iterates over vgas GAS map in *gas_idx* key order 0,1,2,.. and within
0690 each GAS iterates over the "layers" (aka CSGPrim of the CSGSolid)
0691 counting the number of *sbt* records encountered until reach (solid_idx_, layer_idx_)
0692 at which point returns *offset_sbt*.
0693 
0694 This assumes(implies) that only enabled mergedmesh have vgas entries.
0695 
0696 
0697 **/
0698 int SBT::_getOffset(unsigned q_gas_idx , unsigned q_layer_idx ) const
0699 {
0700     int offset_sbt = 0 ;
0701 
0702     for(IT it=vgas.begin() ; it !=vgas.end() ; it++)
0703     {
0704         unsigned gas_idx = it->first ;
0705         bool trimesh = foundry->isSolidTrimesh(gas_idx);
0706         const std::string& mmlabel = foundry->getSolidMMLabel(gas_idx);
0707         const CSGSolid* so = foundry->getSolid(gas_idx) ;
0708         int numPrim = so->numPrim ;
0709 
0710 #ifdef WITH_SOPTIX_ACCEL
0711         SOPTIX_Accel* gas = it->second ;
0712 #else
0713         const GAS* gas = &(it->second) ;
0714 #endif
0715 
0716         int num_bi = gas->bis.size();
0717         LOG(debug)
0718             << " gas_idx " << gas_idx
0719             << " num_bi " << num_bi
0720             << " trimesh " << ( trimesh ? "YES" : "NO " )
0721             << " mmlabel " << mmlabel
0722             ;
0723 
0724         if(!trimesh) assert(num_bi == 1);
0725         if(trimesh)
0726         {
0727             bool are_equal = num_bi == numPrim ;
0728             LOG_IF(fatal, !are_equal )
0729                 << " UNEXPECTED trimesh with  "
0730                 << " UNEQUAL: "
0731                 << " num_bi " << num_bi
0732                 << " numPrim " << numPrim
0733                 << " gas_idx " << gas_idx
0734                 << " mmlabel " << mmlabel
0735                 ;
0736             assert(are_equal );
0737         }
0738 
0739         for(int j=0 ; j < num_bi ; j++)
0740         {
0741 
0742 #ifdef WITH_SOPTIX_ACCEL
0743             const SOPTIX_BuildInput* bi = gas->bis[j] ;
0744             int num_sbt = bi->numSbtRecords() ;
0745             if(!trimesh) assert( bi->is_BuildInputCustomPrimitiveArray() && num_sbt == numPrim );
0746             if(trimesh)  assert( bi->is_BuildInputTriangleArray() && num_sbt == 1 );
0747 #else
0748             assert( trimesh == false );
0749             const BI& bi = gas->bis[j] ;
0750             const OptixBuildInputCustomPrimitiveArray& buildInputCPA = bi.getBuildInputCPA() ;
0751             unsigned num_sbt = buildInputCPA.numSbtRecords ;  // <-- corresponding to bbox of the GAS
0752 #endif
0753             LOG(debug)
0754                  << " gas_idx " << gas_idx
0755                  << " num_bi " << num_bi
0756                  << " j " << j
0757                  << " num_sbt " << num_sbt
0758                  ;
0759 
0760             for( int k=0 ; k < num_sbt ; k++)
0761             {
0762                 unsigned localPrimIdx = trimesh ? j : k ;
0763                 if( q_gas_idx == gas_idx && q_layer_idx == localPrimIdx ) return offset_sbt ;
0764                 offset_sbt += 1 ;
0765             }
0766         }
0767     }
0768     LOG(error)
0769         << "did not find targetted shape "
0770         << " vgas.size " << vgas.size()
0771         << " q_gas_idx_ " << q_gas_idx
0772         << " q_layer_idx " << q_layer_idx
0773         << " offset_sbt " << offset_sbt
0774         ;
0775     return -1 ;
0776 }
0777 
0778 /**
0779 SBT::getTotalRec
0780 ------------------
0781 
0782 Returns the total number of SBT records for all layers (aka CSGPrim)
0783 of all GAS in the map.
0784 
0785 Corresponds to the total number of enabled Prim in all enabled solids.
0786 
0787 **/
0788 
0789 unsigned SBT::getTotalRec() const
0790 {
0791     unsigned tot_bi = 0 ;
0792     unsigned tot_sbt = 0 ;
0793 
0794     for(IT it=vgas.begin() ; it !=vgas.end() ; it++)
0795     {
0796         unsigned gas_idx = it->first ;
0797         bool trimesh = foundry->isSolidTrimesh(gas_idx);
0798         const std::string& mmlabel = foundry->getSolidMMLabel(gas_idx);
0799 
0800         bool enabled = SGeoConfig::IsEnabledMergedMesh(gas_idx)  ;
0801         LOG_IF(error, !enabled) << "gas_idx " << gas_idx << " enabled " << enabled ;
0802 
0803 
0804 #ifdef WITH_SOPTIX_ACCEL
0805         SOPTIX_Accel* gas = it->second ;
0806 #else
0807         const GAS* gas = &(it->second) ;
0808 #endif
0809         unsigned num_bi = gas->bis.size();
0810         tot_bi += num_bi ;
0811 
0812         LOG(LEVEL)
0813             << " gas_idx " << gas_idx
0814             << " num_bi " << num_bi
0815             << " trimesh " << ( trimesh ? "YES" : "NO " )
0816             << " mmlabel " << mmlabel
0817             ;
0818 
0819         for(unsigned j=0 ; j < num_bi ; j++)
0820         {
0821 #ifdef WITH_SOPTIX_ACCEL
0822             const SOPTIX_BuildInput* bi = gas->bis[j] ;
0823             unsigned num_sbt = bi->numSbtRecords() ;
0824 #else
0825             assert( trimesh == false );
0826             const BI& bi = gas->bis[j] ;
0827             const OptixBuildInputCustomPrimitiveArray& buildInputCPA = bi.getBuildInputCPA() ;
0828             unsigned num_sbt = buildInputCPA.numSbtRecords ;
0829 #endif
0830             tot_sbt += num_sbt ;
0831 
0832             LOG(LEVEL)
0833                 << " gas_idx " << gas_idx
0834                 << " num_bi " << num_bi
0835                 << " j " << j
0836                 << " num_sbt " << num_sbt
0837                 ;
0838 
0839         }
0840     }
0841     assert( tot_bi > 0 && tot_sbt > 0 );
0842     return tot_sbt ;
0843 }
0844 
0845 
0846 /**
0847 SBT::descGAS
0848 --------------
0849 
0850 Description of the sbt record counts per GAS, which corresponds
0851 to the number of prim per solid for all solids.
0852 This is meaningful after createGAS.
0853 
0854 **/
0855 
0856 std::string SBT::descGAS() const
0857 {
0858     unsigned tot_sbt = 0 ;
0859     unsigned tot_bi = 0 ;
0860     std::stringstream ss ;
0861     ss
0862         << "SBT::descGAS"
0863         << " num_gas " << vgas.size()
0864         << " bi.numRec ( "
0865         ;
0866 
0867     for(IT it=vgas.begin() ; it !=vgas.end() ; it++)
0868     {
0869         unsigned gas_idx = it->first ;
0870         bool trimesh = foundry->isSolidTrimesh(gas_idx);
0871         const std::string& mmlabel = foundry->getSolidMMLabel(gas_idx);
0872 
0873 #ifdef WITH_SOPTIX_ACCEL
0874         SOPTIX_Accel* gas = it->second ;
0875 #else
0876         assert( trimesh == false );
0877         const GAS* gas = &(it->second) ;
0878 #endif
0879 
0880         bool enabled = SGeoConfig::IsEnabledMergedMesh(gas_idx)  ;
0881         LOG_IF(error, !enabled)
0882              << " gas_idx " << gas_idx
0883              << " enabled " << enabled
0884              << " trimesh " << ( trimesh ? "YES" : "NO " )
0885              << " mmlabel " << mmlabel
0886              ;
0887 
0888         unsigned num_bi = gas->bis.size();
0889         tot_bi += num_bi ;
0890         for(unsigned j=0 ; j < num_bi ; j++)
0891         {
0892 
0893 #ifdef WITH_SOPTIX_ACCEL
0894             const SOPTIX_BuildInput* bi = gas->bis[j] ;
0895             unsigned num_sbt = bi->numSbtRecords() ;
0896 #else
0897             const BI& bi = gas->bis[j] ;
0898             const OptixBuildInputCustomPrimitiveArray& buildInputCPA = bi.getBuildInputCPA() ;
0899             unsigned num_sbt = buildInputCPA.numSbtRecords ;
0900 #endif
0901             ss << num_sbt << " " ;
0902             tot_sbt += num_sbt ;
0903         }
0904     }
0905 
0906     ss << " ) "
0907        << " tot_sbt " << tot_sbt
0908        << " tot_bi " << tot_bi
0909        ;
0910 
0911     std::string str = ss.str();
0912     return str ;
0913 }
0914 
0915 
0916 
0917 /**
0918 SBT::createHitgroup
0919 ---------------------
0920 
0921 Analytic case
0922 ~~~~~~~~~~~~~~
0923 
0924 The hitgroup array has records for all active Prims of all active Solid.
0925 The records hold (numNode, nodeOffset) of all those active Prim.
0926 
0927 For analytic geom all HitGroup SBT records have the same hitgroup_pg,
0928 different shapes are distinguished by program data not program code
0929 
0930 Prim Selection
0931 ~~~~~~~~~~~~~~~~
0932 
0933 Thoughts on how to implement Prim selection with CSGPrim::MakeSpec
0934 
0935 Q: is there a bi for each node ?
0936 A: NO, roughly speaking the bi hold the bbox references for all CSGPrim of the solid(=GAS)
0937 
0938 
0939 How to do this when each solid can be tri/ana ?
0940 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
0941 
0942 Q: Still one hitgroup_pg (PIP.cc) ?
0943 
0944 
0945 
0946 Ideas for simplification
0947 ~~~~~~~~~~~~~~~~~~~~~~~~~
0948 
0949 1. Could collect a vector like SOPTIX_SBT.h avoiding the need for getTotalRec
0950 
0951 
0952 Note tri/ana structural difference
0953 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
0954 
0955 +----------------+-----------------------------------------+
0956 |  Geom type     | Structure (n-CSGPrim for each CSGSolid) |
0957 +================+=========================================+
0958 | Analytic       |   1 GAS : 1 BI : n-SBT records          |
0959 +----------------+-----------------------------------------+
0960 | Triangulated   |   1 GAS : n-BI : 1-SBT record           |
0961 +----------------+-----------------------------------------+
0962 
0963 **/
0964 
0965 
0966 #ifdef WITH_SOPTIX_ACCEL
0967 void SBT::createHitgroup()
0968 {
0969     unsigned num_solid = foundry->getNumSolid();
0970     unsigned num_gas = vgas.size();
0971     unsigned tot_rec = getTotalRec();   // corresponds to the total number of enabled Prim in all enabled solids
0972 
0973     LOG(LEVEL)
0974         << " WITH_SOPTIX_ACCEL "
0975         << " num_solid " << num_solid
0976         << " num_gas " << num_gas
0977         << " tot_rec " << tot_rec
0978         ;
0979 
0980     hitgroup = new HitGroup[tot_rec] ;
0981     HitGroup* hg = hitgroup ;
0982 
0983     for(unsigned i=0 ; i < tot_rec ; i++)   // pack headers CPU side
0984          OPTIX_CHECK( optixSbtRecordPackHeader( pip->hitgroup_pg, hitgroup + i ) );
0985 
0986     unsigned sbt_offset = 0 ;
0987 
0988     for(IT it=vgas.begin() ; it !=vgas.end() ; it++)
0989     {
0990         unsigned gas_idx = it->first ;
0991         SOPTIX_Accel* gas = it->second ;
0992 
0993 
0994         int num_bi = gas->bis.size();
0995 
0996         bool trimesh = foundry->isSolidTrimesh(gas_idx);
0997         const std::string& mmlabel = foundry->getSolidMMLabel(gas_idx);
0998 
0999         const SOPTIX_MeshGroup* xmg = trimesh ? xgas.at(gas_idx) : nullptr ;
1000         const SCUDA_MeshGroup* cmg = xmg ? xmg->cmg : nullptr ;
1001 
1002 
1003         const CSGSolid* so = foundry->getSolid(gas_idx) ;
1004         int numPrim = so->numPrim ;
1005         int primOffset = so->primOffset ;
1006 
1007         LOG(LEVEL)
1008             << " WITH_SOPTIX_ACCEL "
1009             << " gas_idx " << gas_idx
1010             << " trimesh " << ( trimesh ? "YES" : "NO " )
1011             << " num_bi " << num_bi
1012             << " mmlabel " << mmlabel
1013             << " so.numPrim " << numPrim
1014             << " so.primOffset " << primOffset
1015             ;
1016 
1017         if(!trimesh) assert( num_bi == 1 );
1018         if(trimesh) assert( num_bi == numPrim );
1019 
1020 
1021         // (all CSGPrim "layers" of the compound CSGSolid are ana/tri, with no mixing)
1022         //
1023         // ana: outer j loop is mute [num_bi=1,num_sbt>=1], inner k loop over CSGPrim "layers"
1024         // tri: inner k loop is mute [num_bi>=1,num_sbt=1], outer j loop over CSGPrim "layers"
1025 
1026         for(int j=0 ; j < num_bi ; j++)
1027         {
1028             const SOPTIX_BuildInput* bi = gas->bis[j] ;
1029             int num_sbt = bi->numSbtRecords() ;
1030 
1031             if(!trimesh) assert( bi->is_BuildInputCustomPrimitiveArray() && num_sbt == numPrim );
1032             if(trimesh)  assert( bi->is_BuildInputTriangleArray() && num_sbt == 1 );
1033 
1034             LOG(LEVEL)
1035                 << " gas_idx " << gas_idx
1036                 << " num_sbt " << num_sbt << "(ana: num_sbt is num_CSGPrim in CSGSolid, tri: num_sbt is 1)"
1037                 ;
1038 
1039             for( int k=0 ; k < num_sbt ; k++)
1040             {
1041                 unsigned localPrimIdx = trimesh ? j : k ;
1042                 unsigned globalPrimIdx = primOffset + localPrimIdx ;
1043                 const CSGPrim* prim = foundry->getPrim( globalPrimIdx );
1044 
1045                 unsigned globalPrimIdx_1 = prim->globalPrimIdx();
1046 
1047                 bool same_globalPrimIdx = globalPrimIdx == globalPrimIdx_1 ;
1048                 LOG_IF(info, !same_globalPrimIdx)
1049                     << " globalPrimIdx   " << std::setw(5) << globalPrimIdx
1050                     << " globalPrimIdx_1 " << std::setw(5) << globalPrimIdx_1
1051                     << " YOU PROBABLY NEED TO RECREATE THE PERSISTED CSGFoundry GEOMETRY "
1052                     ;
1053 
1054                 assert( globalPrimIdx == globalPrimIdx_1 );
1055 
1056                 int boundary = foundry->getPrimBoundary_(prim);
1057                 assert( boundary > -1 );
1058 
1059                 if( trimesh == false )  // analytic
1060                 {
1061                     setPrimData( hg->data.prim, prim );  // copy numNode, nodeOffset from CSGPrim into hg->data
1062                 }
1063                 else
1064                 {
1065                     setMeshData( hg->data.mesh, cmg, localPrimIdx, boundary, globalPrimIdx );
1066                 }
1067 
1068                 unsigned check_sbt_offset = getOffset(gas_idx, localPrimIdx );
1069 
1070                 bool sbt_offset_expect = check_sbt_offset == sbt_offset ;
1071                 assert( sbt_offset_expect  );
1072                 if(!sbt_offset_expect) std::raise(SIGINT);
1073 
1074                 hg++ ;
1075                 sbt_offset++ ;
1076             }
1077         }
1078     }
1079     UploadHitGroup(sbt, d_hitgroup, hitgroup, tot_rec );
1080 }
1081 #else
1082 
1083 /**
1084 SBT::createHitgroup NOT:WITH_SOPTIX_ACCEL only CustomPrimitiveArray implemented
1085 ----------------------------------------------------------------------------------
1086 
1087 **/
1088 void SBT::createHitgroup()
1089 {
1090     unsigned num_solid = foundry->getNumSolid();
1091     unsigned num_gas = vgas.size();
1092     unsigned tot_rec = getTotalRec();   // corresponds to the total number of enabled Prim in all enabled solids
1093 
1094     LOG(info) << " NOT:WITH_SOPTIX_ACCEL " ;
1095     LOG(LEVEL)
1096         << " num_solid " << num_solid
1097         << " num_gas " << num_gas
1098         << " tot_rec " << tot_rec
1099         ;
1100 
1101     hitgroup = new HitGroup[tot_rec] ;
1102     HitGroup* hg = hitgroup ;
1103 
1104     for(unsigned i=0 ; i < tot_rec ; i++)   // pack headers CPU side
1105          OPTIX_CHECK( optixSbtRecordPackHeader( pip->hitgroup_pg, hitgroup + i ) );
1106 
1107     unsigned sbt_offset = 0 ;
1108 
1109 
1110     for(IT it=vgas.begin() ; it !=vgas.end() ; it++)
1111     {
1112         unsigned gas_idx = it->first ;
1113         const GAS* gas = &(it->second) ;
1114         unsigned num_bi = gas->bis.size();
1115         assert( num_bi == 1 );  // always 1 with analytic, can be more with triangulated SMeshGroup
1116 
1117         bool trimesh = foundry->isSolidTrimesh(gas_idx);
1118         LOG_IF( fatal, trimesh ) << " NOT:WITH_SOPTIX_ACCEL trimesh NOT ALLOWED gas_idx " << gas_idx ;
1119         assert( !trimesh );
1120 
1121         const CSGSolid* so = foundry->getSolid(gas_idx) ;
1122         int numPrim = so->numPrim ;
1123         int primOffset = so->primOffset ;
1124 
1125         LOG(LEVEL) << "gas_idx " << gas_idx << " so.numPrim " << numPrim << " so.primOffset " << primOffset  ;
1126 
1127         for(unsigned j=0 ; j < num_bi ; j++)
1128         {
1129             const BI& bi = gas->bis[j] ;
1130             const OptixBuildInputCustomPrimitiveArray& buildInputCPA = bi.getBuildInputCPA() ;
1131             unsigned num_sbt = buildInputCPA.numSbtRecords ;
1132             assert( num_sbt == unsigned(numPrim) ) ;
1133 
1134             for( unsigned k=0 ; k < num_sbt ; k++)
1135             {
1136                 unsigned localPrimIdx = k ;
1137 
1138                 unsigned globalPrimIdx = primOffset + localPrimIdx ;
1139                 const CSGPrim* prim = foundry->getPrim( globalPrimIdx );
1140                 setPrimData( hg->data.prim, prim, globalPrimIdx );  // copy numNode, nodeOffset from CSGPrim into hg->data
1141                 unsigned check_sbt_offset = getOffset(gas_idx, localPrimIdx );
1142 
1143                 bool sbt_offset_expect = check_sbt_offset == sbt_offset ;
1144                 assert( sbt_offset_expect  );
1145                 if(!sbt_offset_expect) std::raise(SIGINT);
1146 
1147                 hg++ ;
1148                 sbt_offset++ ;
1149             }
1150         }
1151     }
1152     UploadHitGroup(sbt, d_hitgroup, hitgroup, tot_rec );
1153 }
1154 #endif
1155 
1156 
1157 
1158 void SBT::UploadHitGroup(OptixShaderBindingTable& sbt, CUdeviceptr& d_hitgroup, HitGroup* hitgroup, size_t tot_rec )
1159 {
1160     CUDA_CHECK( cudaMalloc(reinterpret_cast<void**>( &d_hitgroup ), sizeof(HitGroup)*tot_rec ));
1161     CUDA_CHECK( cudaMemcpy(reinterpret_cast<void*>( d_hitgroup ), hitgroup, sizeof(HitGroup)*tot_rec, cudaMemcpyHostToDevice ));
1162 
1163     sbt.hitgroupRecordBase  = d_hitgroup;
1164     sbt.hitgroupRecordStrideInBytes = sizeof(HitGroup);
1165     sbt.hitgroupRecordCount = tot_rec ;
1166 }
1167 
1168 void SBT::destroyHitgroup()
1169 {
1170     CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_hitgroup ) ) );
1171 }
1172 
1173 
1174 void SBT::checkHitgroup()
1175 {
1176     unsigned num_solid = foundry->getNumSolid();
1177     unsigned num_prim = foundry->getNumPrim();
1178     unsigned num_sbt = sbt.hitgroupRecordCount ;
1179 
1180     LOG(LEVEL)
1181         << " num_sbt (sbt.hitgroupRecordCount) " << num_sbt
1182         << " num_solid " << num_solid
1183         << " num_prim " << num_prim
1184         ;
1185 }
1186 
1187 
1188 
1189 
1190 /**
1191 SBT::setPrimData
1192 -----------------
1193 
1194 Called from SBT::createHitgroup to populate HitGroupData for analytic geometry.
1195 
1196 **/
1197 
1198 void SBT::setPrimData( CustomPrim& cp, const CSGPrim* prim )
1199 {
1200     cp.numNode = prim->numNode();
1201     cp.nodeOffset = prim->nodeOffset();
1202     cp.globalPrimIdx = prim->globalPrimIdx();
1203 }
1204 
1205 void SBT::checkPrimData( CustomPrim& cp, const CSGPrim* prim)
1206 {
1207     assert( cp.numNode == prim->numNode() );
1208     assert( cp.nodeOffset == prim->nodeOffset() );
1209 
1210 }
1211 void SBT::dumpPrimData( const CustomPrim& cp ) const
1212 {
1213     std::cout
1214         << "SBT::dumpPrimData"
1215         << " cp.numNode " << cp.numNode
1216         << " cp.nodeOffset " << cp.nodeOffset
1217         << std::endl
1218         ;
1219 }
1220 
1221 
1222 #ifdef WITH_SOPTIX_ACCEL
1223 /**
1224 SBT::setMeshData
1225 -------------------
1226 
1227 Note similarity to SOPTIX_SBT::initHitgroup
1228 
1229 **/
1230 
1231 void SBT::setMeshData( TriMesh& tm, const SCUDA_MeshGroup* cmg, int j, int boundary, unsigned globalPrimIdx )
1232 {
1233     tm.boundary = boundary ;
1234     tm.vertex = reinterpret_cast<float3*>( cmg->vtx.pointer(j) );
1235     tm.normal = reinterpret_cast<float3*>( cmg->nrm.pointer(j) );
1236     tm.indice = reinterpret_cast<uint3*>(  cmg->idx.pointer(j) );
1237     tm.globalPrimIdx = globalPrimIdx ;
1238 }
1239 #endif
1240 
1241 
1242 
1243 
1244 
1245
[ Source navigation ] [ Diff markup ] [ Identifier search ] [ general search ]

This page was automatically generated by the 2.3.7 LXR engine.

The LXR team
Valid CSS 2.1! Valid HTML 4.01!