Skip to content

Commit 22e5d45

Browse files
authored
fix(qudarap): exclude fake_propagate CUDA path from PRODUCTION builds (#255)
Guard the test-only `_QSim_fake_propagate` launcher and its host entry point with `!defined(PRODUCTION)`. In PRODUCTION builds, make `QSim::fake_propagate` fail fast on the host side instead of compiling the CUDA path. This avoids the CUDA 12.5.x `cicc` compile-time blow-up in `QSim.cu` Release builds while leaving the normal runtime propagation code unchanged and preserving fake-propagate behavior in non-PRODUCTION builds. Verified with: - direct Release `nvcc` compile of `QSim.cu` - `cmake --build /opt/eic-opticks/build --parallel --target install` An actual fix for #57
1 parent bc79bc6 commit 22e5d45

File tree

3 files changed

+13
-17
lines changed

3 files changed

+13
-17
lines changed

qudarap/QSim.cc

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1519,9 +1519,9 @@ quad2* QSim::UploadFakePRD(const NP* ip, const NP* prd) // static
15191519
return d_prd ;
15201520
}
15211521

1522-
1523-
1522+
#if !defined(PRODUCTION)
15241523
extern void QSim_fake_propagate_launch(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim, quad2* prd );
1524+
#endif
15251525

15261526
/**
15271527
@@ -1545,6 +1545,12 @@ using common QEvt functionality
15451545

15461546
void QSim::fake_propagate( const NP* prd, unsigned type )
15471547
{
1548+
#if defined(PRODUCTION)
1549+
(void)prd;
1550+
(void)type;
1551+
LOG(fatal) << "QSim::fake_propagate is disabled in PRODUCTION builds";
1552+
std::raise(SIGINT);
1553+
#else
15481554
const NP* ip = sev->getInputPhoton();
15491555
int num_ip = ip ? ip->shape[0] : 0 ;
15501556
assert( num_ip > 0 );
@@ -1584,6 +1590,7 @@ void QSim::fake_propagate( const NP* prd, unsigned type )
15841590

15851591

15861592
LOG(LEVEL) << "]" ;
1593+
#endif
15871594
}
15881595

15891596

qudarap/QSim.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -648,6 +648,7 @@ extern void QSim_photon_launch(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim,
648648
}
649649

650650

651+
#if !defined(PRODUCTION)
651652
/**
652653
_QSim_fake_propagate
653654
-----------------------
@@ -686,8 +687,7 @@ extern void QSim_fake_propagate_launch(dim3 numBlocks, dim3 threadsPerBlock, qsi
686687
_QSim_fake_propagate<<<numBlocks,threadsPerBlock>>>( sim, prd );
687688
}
688689

689-
690-
690+
#endif
691691

692692
__global__ void _QSim_boundary_lookup_all(qsim* sim, quad* lookup, unsigned width, unsigned height )
693693
{

qudarap/qsim.h

Lines changed: 2 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -24,10 +24,8 @@ Canonical use is from CSGOptiX/CSGOptiX7.cu:simulate
2424

2525
#if defined(__CUDACC__) || defined(__CUDABE__)
2626
#define QSIM_METHOD __device__
27-
#define QSIM_NOINLINE __noinline__
2827
#else
2928
#define QSIM_METHOD
30-
#define QSIM_NOINLINE
3129
#endif
3230

3331
#include "OpticksGenstep.h"
@@ -70,11 +68,6 @@ Canonical use is from CSGOptiX/CSGOptiX7.cu:simulate
7068

7169
struct qcerenkov ;
7270

73-
static QSIM_METHOD QSIM_NOINLINE float qsim_rng_uniform(RNG *rng)
74-
{
75-
return curand_uniform(rng);
76-
}
77-
7871
struct qsim
7972
{
8073
qbase* base ;
@@ -1093,13 +1086,9 @@ inline QSIM_METHOD int qsim::propagate_at_boundary(unsigned& flag, RNG& rng, sct
10931086

10941087

10951088
#if !defined(PRODUCTION) && defined(DEBUG_TAG)
1096-
const float u_boundary_burn = qsim_rng_uniform(&rng) ; // needed for random consumption alignment with Geant4 G4OpBoundaryProcess::PostStepDoIt
1089+
const float u_boundary_burn = curand_uniform(&rng) ; // needed for random consumption alignment with Geant4 G4OpBoundaryProcess::PostStepDoIt
10971090
#endif
1098-
// Keep the scalar Philox draw behind a noinline helper. CUDA 12.5.1 can
1099-
// spend pathological compile time when this curand state-machine is fully
1100-
// inlined into this large function at high optimization. CUDA 13.0.2
1101-
// compiles cleanly either way, but this preserves the original RNG stream.
1102-
const float u_reflect = qsim_rng_uniform(&rng);
1091+
const float u_reflect = curand_uniform(&rng) ;
11031092
bool reflect = u_reflect > TransCoeff ;
11041093

11051094
#if !defined(PRODUCTION) && defined(DEBUG_TAG)

0 commit comments

Comments
 (0)