Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
57 changes: 56 additions & 1 deletion CSGOptiX/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -56,7 +57,8 @@ target_compile_definitions(${OPTIX_PTX_TARGET}
$<$<CONFIG:Release>:PRODUCTION>
OPTICKS_OKCONF
OPTICKS_QUDARAP
NVCC)
NVCC
$<$<BOOL:${SIMPHONY_RNG_REBUILD}>:SIMPHONY_RNG_REBUILD>)
target_link_libraries(${OPTIX_PTX_TARGET} PUBLIC OptiX::OptiX)

get_filename_component(ptx_stem "${OPTIX_PTX_SRC}" NAME_WE)
Expand All @@ -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}
$<$<NOT:$<BOOL:${SIMPHONY_RNG_REBUILD}>>: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
"$<TARGET_OBJECTS:${RNG_REBUILD_OFF_PTX_TARGET}>"
"${PROJECT_BINARY_DIR}/ptx_rng_rebuild/off/CSGOptiX.ptx"
COMMAND ${CMAKE_COMMAND} -E copy_if_different
"$<TARGET_OBJECTS:${RNG_REBUILD_ON_PTX_TARGET}>"
"${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} )
Expand Down
9 changes: 9 additions & 0 deletions CSGOptiX/CSGOptiX.cu
Original file line number Diff line number Diff line change
Expand Up @@ -414,13 +414,22 @@ 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)
{
case 0u: trace<false>( params.handle, ctx.p.pos, ctx.p.mom, tmin, params.tmax, prd, params.vizmask, params.PropagateRefineDistance ); break ;
case 1u: trace<true>( 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 ) ;

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The Philox overload of skipahead(n) counts scalar elements, not four-element blocks. Its implementation applies n & 3 to STATE, advances the counter by n / 4, handles the intra-block carry, and regenerates output. Therefore 4 * ctr.x + STATE is the scalar position, and reinitializing followed by skipahead(delta) restores arbitrary intra-block positions. The fixed-seed ON/OFF CTest also verifies identical per-event counts and byte-identical hit output.

#endif

if( prd->boundary() == 0xffffu ) break ; // SHOULD ONLY HAPPEN FOR PHOTONS STARTING OUTSIDE WORLD
// propagate can do nothing meaningful without a boundary

Expand Down
10 changes: 10 additions & 0 deletions docs/getting-started.md
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
5 changes: 5 additions & 0 deletions docs/performance-and-debugging.md
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
3 changes: 3 additions & 0 deletions src/g4app.h
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down
104 changes: 104 additions & 0 deletions tests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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=$<TARGET_FILE:simg4ox>
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}")
Expand Down
34 changes: 29 additions & 5 deletions tests/check_simg4ox_multievent.py
Original file line number Diff line number Diff line change
Expand Up @@ -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)}")

Expand All @@ -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}")
Expand All @@ -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}")
Expand Down
20 changes: 14 additions & 6 deletions tests/test_simg4ox_multievent.sh
Original file line number Diff line number Diff line change
Expand Up @@ -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}"
Loading