diff --git a/CSGOptiX/CMakeLists.txt b/CSGOptiX/CMakeLists.txt index 41da270d9..f651c1614 100644 --- a/CSGOptiX/CMakeLists.txt +++ b/CSGOptiX/CMakeLists.txt @@ -6,6 +6,7 @@ set(SOURCES) set(OPTIX_PTX_SRC CSGOptiX.cu) set(OPTIX_PTX_TARGET CSGOptiX_ptx) +option(SIMPHONY_RNG_REBUILD "Rebuild the Philox RNG state after each trace instead of carrying it live" OFF) list(APPEND SOURCES Params.cc @@ -56,7 +57,8 @@ target_compile_definitions(${OPTIX_PTX_TARGET} $<$:PRODUCTION> OPTICKS_OKCONF OPTICKS_QUDARAP - NVCC) + NVCC + $<$:SIMPHONY_RNG_REBUILD>) target_link_libraries(${OPTIX_PTX_TARGET} PUBLIC OptiX::OptiX) get_filename_component(ptx_stem "${OPTIX_PTX_SRC}" NAME_WE) @@ -69,6 +71,59 @@ add_custom_target(csgoptix_ptx ) add_dependencies(csgoptix_ptx ${OPTIX_PTX_TARGET}) +if(BUILD_TESTING) + set(RNG_REBUILD_OPPOSITE_PTX_TARGET CSGOptiX_ptx_rng_rebuild_opposite) + add_library(${RNG_REBUILD_OPPOSITE_PTX_TARGET} OBJECT ${OPTIX_PTX_SRC}) + set_target_properties(${RNG_REBUILD_OPPOSITE_PTX_TARGET} + PROPERTIES CUDA_PTX_COMPILATION ON) + + get_target_property(RNG_REBUILD_PTX_INCLUDE_DIRS + ${OPTIX_PTX_TARGET} INCLUDE_DIRECTORIES) + get_target_property(RNG_REBUILD_PTX_COMPILE_DEFINITIONS + ${OPTIX_PTX_TARGET} COMPILE_DEFINITIONS) + list(FILTER RNG_REBUILD_PTX_COMPILE_DEFINITIONS + EXCLUDE REGEX "SIMPHONY_RNG_REBUILD") + + target_include_directories(${RNG_REBUILD_OPPOSITE_PTX_TARGET} PUBLIC + ${RNG_REBUILD_PTX_INCLUDE_DIRS}) + target_compile_definitions(${RNG_REBUILD_OPPOSITE_PTX_TARGET} + PUBLIC ${RNG_REBUILD_PTX_COMPILE_DEFINITIONS} + $<$>:SIMPHONY_RNG_REBUILD>) + target_link_libraries(${RNG_REBUILD_OPPOSITE_PTX_TARGET} PUBLIC OptiX::OptiX) + + if(SIMPHONY_RNG_REBUILD) + set(RNG_REBUILD_OFF_PTX_TARGET ${RNG_REBUILD_OPPOSITE_PTX_TARGET}) + set(RNG_REBUILD_ON_PTX_TARGET ${OPTIX_PTX_TARGET}) + else() + set(RNG_REBUILD_OFF_PTX_TARGET ${OPTIX_PTX_TARGET}) + set(RNG_REBUILD_ON_PTX_TARGET ${RNG_REBUILD_OPPOSITE_PTX_TARGET}) + endif() + + set(SIMPHONY_RNG_REBUILD_PTX_OFF_DIR + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/off" PARENT_SCOPE) + set(SIMPHONY_RNG_REBUILD_PTX_ON_DIR + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/on" PARENT_SCOPE) + + add_custom_target(csgoptix_rng_rebuild_ptx + COMMAND ${CMAKE_COMMAND} -E make_directory + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/off" + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/on" + COMMAND ${CMAKE_COMMAND} -E copy_if_different + "$" + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/off/CSGOptiX.ptx" + COMMAND ${CMAKE_COMMAND} -E copy_if_different + "$" + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/on/CSGOptiX.ptx" + BYPRODUCTS + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/off/CSGOptiX.ptx" + "${PROJECT_BINARY_DIR}/ptx_rng_rebuild/on/CSGOptiX.ptx" + VERBATIM + ) + add_dependencies(csgoptix_rng_rebuild_ptx + ${RNG_REBUILD_OFF_PTX_TARGET} + ${RNG_REBUILD_ON_PTX_TARGET}) +endif() + message(STATUS "Configured PTX build for ${OPTIX_PTX_SRC}") add_library( ${name} SHARED ${SOURCES} ${HEADERS} ) diff --git a/CSGOptiX/CSGOptiX.cu b/CSGOptiX/CSGOptiX.cu index 08a2e0d2c..e062354c0 100644 --- a/CSGOptiX/CSGOptiX.cu +++ b/CSGOptiX/CSGOptiX.cu @@ -414,6 +414,10 @@ static __forceinline__ __device__ void simulate( const uint3& launch_idx, const { float tmin = ( ctx.p.orient_boundary_flag & params.PropagateEpsilon0Mask ) ? params.tmin0 : params.tmin ; +#ifdef SIMPHONY_RNG_REBUILD + unsigned rng_p = ( rng.ctr.x << 2 ) | ( rng.STATE & 3u ) ; +#endif + // intersect query filling (quad2)prd switch(params.PropagateRefine) { @@ -421,6 +425,11 @@ static __forceinline__ __device__ void simulate( const uint3& launch_idx, const case 1u: trace( params.handle, ctx.p.pos, ctx.p.mom, tmin, params.tmax, prd, params.vizmask, params.PropagateRefineDistance ); break ; } +#ifdef SIMPHONY_RNG_REBUILD + sim->rng->init( rng, sim->evt->index, photon_idx ); + skipahead( (unsigned long long)( rng_p - ( ( rng.ctr.x << 2 ) | ( rng.STATE & 3u ) ) ), &rng ) ; +#endif + if( prd->boundary() == 0xffffu ) break ; // SHOULD ONLY HAPPEN FOR PHOTONS STARTING OUTSIDE WORLD // propagate can do nothing meaningful without a boundary diff --git a/docs/getting-started.md b/docs/getting-started.md index abc76da2a..6d0f19c78 100644 --- a/docs/getting-started.md +++ b/docs/getting-started.md @@ -138,6 +138,16 @@ The full suite includes GPU-backed tests. You can build without a GPU, but running those tests requires a compatible NVIDIA driver and GPU access from the container. +### CMake build options + +Set supported build options when configuring the CMake build tree. Re-run the +configure command after changing an option, then rebuild the affected targets. + +| Option | Default | Scope | Description | +|---|---|---|---| +| `SIMPHONY_RNG_REBUILD` | `OFF` | CSGOptiX PTX | Reconstruct the Philox RNG state after each OptiX trace instead of carrying it live. | + + ## Build directly on the host If you prefer to manage the toolchain yourself, install: diff --git a/docs/performance-and-debugging.md b/docs/performance-and-debugging.md index 9c4d92aa6..83c61de3d 100644 --- a/docs/performance-and-debugging.md +++ b/docs/performance-and-debugging.md @@ -26,6 +26,11 @@ docker run --rm -t -v /tmp/out:/tmp/out simphony:release \ run-performance -g tests/geom/opticks_raindrop.gdml -o /tmp/out/release ``` +To benchmark the experimental CSGOptiX RNG-state rebuild variant, see the +[`SIMPHONY_RNG_REBUILD` CMake option](getting-started.md#cmake-build-options). +It is a compile-time choice, so reconfigure and rebuild the PTX target before +measuring each variant. + ### Interpreting `simg4ox` MT timings `simg4ox --threads N` parallelizes Geant4 CPU tracking of its configured torch diff --git a/src/g4app.h b/src/g4app.h index fa6a09dad..9e269fc88 100644 --- a/src/g4app.h +++ b/src/g4app.h @@ -423,6 +423,9 @@ struct EventAction : G4UserEventAction cudaDeviceSynchronize(); SEvt* sev_gpu = SEvt::Get_EGPU(); + const auto num_gpu_photons = sev_gpu->getNumPhotonCollected(); + G4cout << "EventAction::SimulateOnGPU: Event " << event_id + << ": Simulated GPU photons: " << num_gpu_photons << G4endl; auto gpu_hits = CollectGPUHits(sev_gpu); gx->reset(event_id); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 2997c615d..f5edca90a 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -52,6 +52,110 @@ set_tests_properties(Integration.simg4ox_multievent PROPERTIES TIMEOUT 120 ) +add_dependencies(simg4ox csgoptix_rng_rebuild_ptx) + +set(SIMPHONY_RNG_REBUILD_TEST_WORKDIR + "${CMAKE_CURRENT_BINARY_DIR}/simg4ox_rng_rebuild") + +function(add_simg4ox_rng_rebuild_run MODE PTX_DIR) + set(TEST_NAME "Integration.simg4ox_rng_rebuild.${MODE}") + set(TEST_WORKDIR "${SIMPHONY_RNG_REBUILD_TEST_WORKDIR}/${MODE}") + set(TEST_PTX "${PTX_DIR}/CSGOptiX.ptx") + file(MAKE_DIRECTORY "${TEST_WORKDIR}") + + add_test( + NAME "${TEST_NAME}" + COMMAND ${CMAKE_COMMAND} -E env + REPO_DIR=${PROJECT_SOURCE_DIR} + PYTHON=${Python3_EXECUTABLE} + SIMG4OX_BIN=$ + SIMPHONY_PTX_DIR=${PTX_DIR} + SIMG4OX_SEED=42 + ${BASH_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/test_simg4ox_multievent.sh + ) + + set_tests_properties("${TEST_NAME}" PROPERTIES + WORKING_DIRECTORY "${TEST_WORKDIR}" + LABELS "integration;rng-rebuild" + REQUIRED_FILES "${TEST_PTX}" + TIMEOUT 120 + ) +endfunction() + +add_simg4ox_rng_rebuild_run(off "${SIMPHONY_RNG_REBUILD_PTX_OFF_DIR}") +add_simg4ox_rng_rebuild_run(on "${SIMPHONY_RNG_REBUILD_PTX_ON_DIR}") + +set(SIMPHONY_RNG_REBUILD_OFF_WORKDIR "${SIMPHONY_RNG_REBUILD_TEST_WORKDIR}/off") +set(SIMPHONY_RNG_REBUILD_ON_WORKDIR "${SIMPHONY_RNG_REBUILD_TEST_WORKDIR}/on") +set(SIMPHONY_RNG_REBUILD_RUN_TESTS + "Integration.simg4ox_rng_rebuild.off;Integration.simg4ox_rng_rebuild.on") +set(SIMPHONY_RNG_REBUILD_OFF_PTX + "${SIMPHONY_RNG_REBUILD_PTX_OFF_DIR}/CSGOptiX.ptx") +set(SIMPHONY_RNG_REBUILD_ON_PTX + "${SIMPHONY_RNG_REBUILD_PTX_ON_DIR}/CSGOptiX.ptx") + +set(SIMPHONY_RNG_REBUILD_COMPARE_PTX [=[ +for file in "$1" "$2" +do + if [[ ! -r "${file}" ]] + then + echo "Missing or unreadable PTX file: ${file}" >&2 + exit 2 + fi +done + +cmp --silent -- "$1" "$2" +status=$? +if [[ "${status}" -eq 0 ]] +then + echo "RNG rebuild ON and OFF unexpectedly produced identical PTX" >&2 + exit 1 +fi +if [[ "${status}" -eq 1 ]] +then + exit 0 +fi + +echo "PTX comparison failed with status ${status}" >&2 +exit "${status}" +]=]) + +add_test( + NAME Integration.simg4ox_rng_rebuild.ptx_differ + COMMAND ${BASH_EXECUTABLE} -c + "${SIMPHONY_RNG_REBUILD_COMPARE_PTX}" + rng-rebuild-ptx-compare + "${SIMPHONY_RNG_REBUILD_OFF_PTX}" + "${SIMPHONY_RNG_REBUILD_ON_PTX}" +) +set_tests_properties(Integration.simg4ox_rng_rebuild.ptx_differ PROPERTIES + DEPENDS "${SIMPHONY_RNG_REBUILD_RUN_TESTS}" + LABELS "integration;rng-rebuild" + REQUIRED_FILES "${SIMPHONY_RNG_REBUILD_OFF_PTX};${SIMPHONY_RNG_REBUILD_ON_PTX}" +) + +add_test( + NAME Integration.simg4ox_rng_rebuild.counts_match + COMMAND ${CMAKE_COMMAND} -E compare_files + "${SIMPHONY_RNG_REBUILD_OFF_WORKDIR}/simg4ox-multievent-summary.txt" + "${SIMPHONY_RNG_REBUILD_ON_WORKDIR}/simg4ox-multievent-summary.txt" +) +set_tests_properties(Integration.simg4ox_rng_rebuild.counts_match PROPERTIES + DEPENDS "${SIMPHONY_RNG_REBUILD_RUN_TESTS}" + LABELS "integration;rng-rebuild" +) + +add_test( + NAME Integration.simg4ox_rng_rebuild.hits_match + COMMAND ${CMAKE_COMMAND} -E compare_files + "${SIMPHONY_RNG_REBUILD_OFF_WORKDIR}/s_hits.npy" + "${SIMPHONY_RNG_REBUILD_ON_WORKDIR}/s_hits.npy" +) +set_tests_properties(Integration.simg4ox_rng_rebuild.hits_match PROPERTIES + DEPENDS "${SIMPHONY_RNG_REBUILD_RUN_TESTS}" + LABELS "integration;rng-rebuild" +) + if(Geant4_multithreaded_FOUND) set(SIMPHONY_SIMG4OX_MULTITHREAD_TEST_WORKDIR "${CMAKE_CURRENT_BINARY_DIR}/simg4ox_multithread") file(MAKE_DIRECTORY "${SIMPHONY_SIMG4OX_MULTITHREAD_TEST_WORKDIR}") diff --git a/tests/check_simg4ox_multievent.py b/tests/check_simg4ox_multievent.py index 533d16650..ee3b2003e 100755 --- a/tests/check_simg4ox_multievent.py +++ b/tests/check_simg4ox_multievent.py @@ -9,21 +9,35 @@ EVENT_PATTERNS = { - "GPU": re.compile(r"Collected GPU hits:\s*(\d+)"), - "G4": re.compile(r"Collected G4\s+hits:\s*(\d+)"), + "GPU": re.compile(r"EventAction::EndOfEventAction: Event (\d+): Collected GPU hits:\s*(\d+)"), + "G4": re.compile(r"EventAction::EndOfEventAction: Event (\d+): Collected G4\s+hits:\s*(\d+)"), } +PHOTON_PATTERN = re.compile( + r"EventAction::SimulateOnGPU: Event (\d+): Simulated GPU photons:\s*(\d+)" +) TOTAL_PATTERNS = { "GPU": re.compile(r"Total GPU hits:\s*(\d+)"), "G4": re.compile(r"Total G4\s+hits:\s*(\d+)"), } +def order_event_counts(matches, label, events): + parsed = [(int(event_id), int(value)) for event_id, value in matches] + event_ids = [event_id for event_id, _ in parsed] + + if len(parsed) != events: + raise AssertionError(f"Expected {events} {label} event counts, found {len(parsed)}") + if sorted(event_ids) != list(range(events)): + raise AssertionError(f"Expected {label} event ids 0..{events - 1}, found {event_ids}") + + counts_by_event = dict(parsed) + return [counts_by_event[event_id] for event_id in range(events)] + + def parse_counts(log_text, label, events): - event_counts = [int(value) for value in EVENT_PATTERNS[label].findall(log_text)] + event_counts = order_event_counts(EVENT_PATTERNS[label].findall(log_text), label, events) total_matches = TOTAL_PATTERNS[label].findall(log_text) - if len(event_counts) != events: - raise AssertionError(f"Expected {events} {label} event counts, found {len(event_counts)}") if len(total_matches) != 1: raise AssertionError(f"Expected one {label} run total, found {len(total_matches)}") @@ -36,6 +50,14 @@ def parse_counts(log_text, label, events): return event_counts, total +def parse_photon_counts(log_text, events): + photon_counts = order_event_counts(PHOTON_PATTERN.findall(log_text), "GPU photon", events) + if any(value == 0 for value in photon_counts): + raise AssertionError(f"Expected non-empty GPU photon counts, found {photon_counts}") + + return photon_counts + + def check_array(path, expected_rows): if not path.is_file(): raise FileNotFoundError(f"Missing run hit array: {path}") @@ -58,12 +80,14 @@ def main(): args = parser.parse_args() log_text = args.log.read_text() + gpu_photon_counts = parse_photon_counts(log_text, args.events) gpu_event_counts, gpu_total = parse_counts(log_text, "GPU", args.events) g4_event_counts, g4_total = parse_counts(log_text, "G4", args.events) gpu_hits = check_array(args.output_dir / "s_hits.npy", gpu_total) g4_hits = check_array(args.output_dir / "g_hits.npy", g4_total) + print(f"GPU_EVENT_PHOTON_COUNTS={gpu_photon_counts}") print(f"GPU_EVENT_COUNTS={gpu_event_counts}") print(f"G4_EVENT_COUNTS={g4_event_counts}") print(f"S_HITS_SHAPE={gpu_hits.shape}") diff --git a/tests/test_simg4ox_multievent.sh b/tests/test_simg4ox_multievent.sh index 8c5596e4c..7be3af253 100755 --- a/tests/test_simg4ox_multievent.sh +++ b/tests/test_simg4ox_multievent.sh @@ -8,19 +8,27 @@ REPO_DIR=${REPO_DIR:-$(cd "${SCRIPT_DIR}/.." && pwd)} SIMG4OX_BIN=${SIMG4OX_BIN:-simg4ox} PYTHON=${PYTHON:-python3} RUN_LOG="${PWD}/simg4ox-multievent.log" +CHECK_SUMMARY="${PWD}/simg4ox-multievent-summary.txt" +SIMG4OX_SEED=${SIMG4OX_SEED:-} export OPTICKS_HOME="${REPO_DIR}" export SIMPHONY_CONFIG_DIR="${SIMPHONY_CONFIG_DIR:-${REPO_DIR}/config}" export PYTHONPATH="${REPO_DIR}${PYTHONPATH:+:${PYTHONPATH}}" -rm -f "${PWD}/g_hits.npy" "${PWD}/s_hits.npy" "${RUN_LOG}" +rm -f "${PWD}/g_hits.npy" "${PWD}/s_hits.npy" "${RUN_LOG}" "${CHECK_SUMMARY}" -"${SIMG4OX_BIN}" \ - -g "${REPO_DIR}/tests/geom/opticks_raindrop.gdml" \ - -m "${REPO_DIR}/tests/run_5evt.mac" \ - -c dev 2>&1 | tee "${RUN_LOG}" +simg4ox_args=( + -g "${REPO_DIR}/tests/geom/opticks_raindrop.gdml" + -m "${REPO_DIR}/tests/run_5evt.mac" + -c dev +) +if [[ -n "${SIMG4OX_SEED}" ]]; then + simg4ox_args+=(-s "${SIMG4OX_SEED}") +fi + +"${SIMG4OX_BIN}" "${simg4ox_args[@]}" 2>&1 | tee "${RUN_LOG}" "${PYTHON}" "${REPO_DIR}/tests/check_simg4ox_multievent.py" \ --log "${RUN_LOG}" \ --output-dir "${PWD}" \ - --events 5 + --events 5 | tee "${CHECK_SUMMARY}"