Profile the LBVH build and detect stages on both the CPU and the GPU

[zfergus]

Description

Adds per-stage profiler coverage to both LBVH broad phases, and a benchmark that dumps the breakdown. Split out of #260, which merged before this commit was pushed.

ipc::LBVH's build was already instrumented; ipc::cuda::LBVH had none, and neither broad phase separated its detection into traversal and candidate output.

• ipc::cuda::CudaEventTimer (utils/cuda/cuda_profiler.cuh) is a ProfilePoint timer backed by CUDA events, so a device stage is measured by when the device ran it rather than when the host returned from the launch. IPC_TOOLKIT_PROFILE_BLOCK_CUDA is its block macro.
• ipc::cuda::LBVH gains blocks for the transfers, the box kernels, the domain reduction, and each tree's Morton, sort, hierarchy, and patch stages, plus parent scopes on build() and detect_host() so time not attributed to a stage shows up as a remainder instead of vanishing.
• materialize() separates the host staging allocation, the device-to-host copy, and the Candidate construction. ipc::LBVH::detect_candidates() separates the traversal from the thread-local merge.
• benchmark_lbvh_stages.cu drives both broad phases over the six benchmark scenes and writes the profiler tree as JSON, following the IPC_TOOLKIT_BENCH_OUTPUT convention of the SIMD assembly benchmark.

What it exposes, on Cloth-Ball (5.2M edge-edge candidates): the device wins the descent outright, 8.19 ms to 2.16 ms (3.8x), then spends 11.74 ms of its 16.96 ms (69%) materializing host Candidate objects, against 17.25 ms of 25.46 ms (68%) for the CPU's thread-local merge. The end-to-end detection ratio is set by the candidate output, not by the traversal. On the build side CUB's radix sort is where the win is: 0.54 ms to 0.077 ms per tree, 7.0x.

API changes

• New ipc::cuda::CudaEventTimer and IPC_TOOLKIT_PROFILE_BLOCK_CUDA, both behind IPC_TOOLKIT_WITH_PROFILER and compiled out when it is off.
• No existing signatures change. The only non-profiling edit is in materialize(), where std::vector<int32_t> h_a(count), h_b(count) becomes a default construction plus resize(count) so the allocation can be timed; both allocate and zero the same way.

Performance

The instrumentation is off by default and free when off. With it on, the CUDA event timers wait on each stage's stop event, serializing stages the build overlaps.

Benchmark	Profiler off, vs. before	Profiler on, vs. off
LBVH::build	-3.7% to +1.1%	-3.6% to +3.5%
cuda::LBVH::build	-5.2% to -0.0%	+3.6% to +26.7%
LBVH::detect_edge_edge_candidates	-4.2% to +1.2%	+1.0% to +8.9%
cuda::LBVH::detect_edge_edge_candidates	-1.1% to +0.5%	-0.4% to +1.3%

• Off, the deltas scatter in both directions with no bench trending one way: run-to-run noise, not overhead. The blocks preprocess to ((void)0) without Tracy.
• On, the GPU build is the one real cost, and it scales inversely with scene size (+24.7% on Cloth-Funnel, +3.6% on Puffer-Ball): its 11 stages each become a device wait instead of pipelining into the single cudaDeviceSynchronize() at the end of the build. Short kernels cannot hide the fixed sync cost; large ones can.
• So the profiled GPU build totals run above the fused ones. cuda_profiler.cuh says so at the point of use: take stage shares from a profiling build, absolute times from one without.

Type of change

• Enhancement (non-breaking change which improves existing functionality)
• New feature (non-breaking change which adds functionality)

How Has This Been Tested?

• [lbvh] and [broad_phase] in CUDA Release with the profiler off: all pass, unchanged from before the instrumentation.
• [lbvh_stages] in a CUDA Release + IPC_TOOLKIT_WITH_PROFILER=ON build, over all six scenes. The benchmark asserts the CPU and GPU return the same candidate count, so the stage split is comparing the same work, and stages are credited self time so they sum to the enclosing call.
• Overhead A/B: all four LBVH benchmarks, median of three runs of 20 samples each, on identical sources with only IPC_TOOLKIT_WITH_PROFILER differing (table above).
• clang-format: 0 replacements.

Test Configuration:

• OS and Version: Linux 5.14.0 (RHEL 9.8)
• CPU: Intel Core Ultra 9 285K (24 cores / 24 threads, no SMT), 24 TBB threads
• GPU: NVIDIA RTX 5080 (sm_120)
• Compiler and Version: GCC 11.5.0 with CUDA 12.8 (V12.8.93), -march=native (AVX2+FMA) and -arch=native, Release

Checklist

• I have followed the project style guide
• My code follows the clang-format style guidelines of this project
• I have performed a self-review of my code
• I have commented my code, particularly in hard-to-understand areas
• My changes generate no new warnings
• New and existing unit tests pass locally with my changes