Enable GPU replay of prebuilt libcu++ lit tests

[jrhemstad]

Description

Today, libcu++ lit testing is split between build and test jobs. The build job compiles the tests. The test job compiles them again, relying on sccache to avoid actual recompilation. Even with a 100% cache hit rate, there is still cost in requesting every compilation artifact from sccache.

This PR keeps the build job responsible for compilation, then packages and uploads the resulting executables. The test job downloads and runs those executables directly instead of invoking the compiler and relying on sccache.

Motivation

To test this idea, I ran all 3,769 libcu++ lit tests locally using CUDA 13.3, GCC15, and SM86. Even with a 100% cloud sccache hit rate, compiling and running normally took 18m10s on the GPU runner. Precompiling the executables took 15m47s off-GPU, and replaying them took 8m05s on the GPU, reducing GPU runner time by 10m05s (55%).

CI results show the same improvement. All jobs below ran 3,769 tests with -j8, CUDA 13.3, GCC15, and C++20.

Job	Time
L4 replay (this PR)	13m37s
H100 previous job	27m42s
H100 previous job	30m14s
T4 previous job	34m19s
T4 previous job	39m10s

The median normal job took 32m17s. Replay took 13m37s, reducing GPU job time by 18m40s (58%).

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[jrhemstad]

/ok to test af4814c

[jrhemstad]

/ok to test d9c4c14

[jrhemstad]

/ok to test ca40e94

[github-actions]

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Result: 1 regression row(s), 2 improvement row(s) above threshold.

Run	Value
Config	public-headers-gcc13
Baseline	origin/main
Preset	all-dev
Targets	cub.headers.base, thrust.cpp.cuda.headers.base, libcudacxx.test.public_headers
GPU / launch args	rtx2080 / --cuda 13.3 --host gcc13

Artifacts: reports and traces

Direct file processing

-f file-processing exclusive --sort total

🔴 Direct file processing — Regressions

Rank	Regression impact	Selected Δ	Baseline	Current	Event	Matched traces
1	0.926151	0.926151	19.530198	20.456349	Processing Header File: libcudacxx/include/cuda/__device/physical_device.h	94

🟢 Direct file processing — Improvements

Rank	Improvement impact	Selected Δ	Baseline	Current	Event	Matched traces
1	1.103432	-1.103432	9.660159	8.556727	Processing Header File: libcudacxx/include/cuda/std/__cccl/prologue.h	548
2	0.427560	-0.427560	3.386191	2.958631	Processing Header File: libcudacxx/include/cuda/std/__cccl/epilogue.h	548

[miscco]

I believe this is too complicated

In lit there is a final test step that removes the build artifacts: inside of _evaluate_pass_test

We can simply add a flag to disable the removal when running in CI

[jrhemstad]

I believe this is too complicated

In lit there is a final test step that removes the build artifacts: inside of _evaluate_pass_test

We can simply add a flag to disable the removal when running in CI

Disabling removal of the executables is exactly what this PR already does.

But that is far from sufficient. The build and test jobs run on different runners. We need to package those executables as an artifact to download and run them on the GPU runners without having to rebuild them.

[coderabbitai]

📝 Walkthrough

Summary by CodeRabbit

• CI Improvements

  • GitHub Actions can now build, package, and reuse libcudacxx test artifacts across jobs.
  • Linux and Windows test workflows download prebuilt artifacts in CI instead of rebuilding them.
  • Local testing continues to build components as needed.

• Testing

  • Added normal, build, and replay test execution modes.
  • Preserved test executables for reuse in GPU and follow-up test jobs.
  • Improved validation and reporting for precompiled test execution.

Walkthrough

Changes

The test harness now supports build and replay modes for preserved libcudacxx executables. Linux and Windows CI producer jobs upload artifacts when consumers exist. Consumer jobs download the artifacts and run replay tests.

libcudacxx CI artifact flow

Layer / File(s)	Summary
Test executable modes libcudacxx/test/utils/libcudacxx/test/config.py, libcudacxx/test/utils/libcudacxx/test/format.py, libcudacxx/test/libcudacxx/CMakeLists.txt	The test configuration validates normal, build, and replay modes. Replay requires existing executables, build and replay preserve executables, and normal mode cleans them up.
Linux artifact production and replay ci/build_libcudacxx.sh, ci/upload_libcudacxx_test_artifacts.sh, ci/test_libcudacxx.sh	Linux producer jobs stage and upload build metadata and test binaries when consumers exist. Consumer jobs download the artifact and run lit in replay mode.
Windows artifact production and replay ci/windows/build_libcudacxx.ps1, ci/windows/test_libcudacxx.ps1	Windows producer jobs conditionally package test artifacts. CI consumer jobs download them and use CI-specific CTest and lit replay presets; local runs continue to build locally.

Suggested reviewers: fbusato, pciolkosz, wmaxey

✨ Finishing Touches 💡 1

🛠️ Fix failing CI checks 💡

• Create stacked PR
• Commit on current branch

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 2

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: b68987f0-08b5-4a46-ae82-11f3a70fe61b

📥 Commits

Reviewing files that changed from the base of the PR and between bee9acf and ca40e94.

📒 Files selected for processing (8)

• ci/build_libcudacxx.sh
• ci/test_libcudacxx.sh
• ci/upload_libcudacxx_test_artifacts.sh
• ci/windows/build_libcudacxx.ps1
• ci/windows/test_libcudacxx.ps1
• libcudacxx/test/libcudacxx/CMakeLists.txt
• libcudacxx/test/utils/libcudacxx/test/config.py
• libcudacxx/test/utils/libcudacxx/test/format.py

[wmaxey]

I believe this is too complicated

In lit there is a final test step that removes the build artifacts: inside of _evaluate_pass_test

We can simply add a flag to disable the removal when running in CI

Yeah, I'm fine with the lit changes including the mode flag. It's the external logic to package the artifacts up that needs to stick around.

[github-actions]

😬 CI Workflow Results

🟥 Finished in 11h 25m: Pass: 99%/504 | Total: 4d 14h | Max: 1h 03m | Hits: 80%/917673

See results here.

AI failure analysis

1. NVRTC program destruction segfaults in custom segmented-reduce test · 1 job

Explanation: The custom-type segmented-reduce test crashes while destroying the NVRTC program after successfully compiling and retrieving its LTO IR. This prevents the Windows CTK 13.3/MSVC 14.50 HostJIT test suite from completing, while the remaining tests continue to pass.

Evidence:

CCCL C Parallel v2 (HostJIT) nvcc MSVC / [CTK13.3 MSVC14.50] Test(amd64, RTX2080), step 4

8/20 Test  #8: cccl.c.parallel.v2.test.segmented_reduce ....................................***Exception: SegFault  4.95 sec
C:\cccl\c\parallel\test\test_util.h(110): FAILED:
  SIGSEGV - Segmentation violation signal

Root cause: The decisive failure occurs at the `nvrtcDestroyProgram(&prog)` call after `nvrtcGetLTOIR` succeeds. The logs contain no native stack trace, so they cannot distinguish an NVRTC 13.3/MSVC interaction from earlier memory corruption; the PR diff only changes libcudacxx test-artifact replay and does not touch this test path, making an unrelated toolchain/test flake likely. Sources: c/parallel/test/test_util.h:110, c/parallel/test/test_segmented_reduce.cpp:359, c/parallel/test/test_segmented_reduce.cpp:405.

Suggested next steps: First rerun only this test on the same Windows CTK 13.3/MSVC 14.50 configuration with `ctest --test-dir build/cccl-c-parallel-v2 -R "^cccl.c.parallel.v2.test.segmented_reduce$" --output-on-failure`. If reproducible, capture a native crash stack and reduce the operation source to a standalone NVRTC create/compile/get-LTOIR/destroy sequence before changing CCCL code.

Copy this prompt into a coding agent

Repository: https://github.com/NVIDIA/cccl
Workflow run: https://github.com/NVIDIA/cccl/actions/runs/30832248360
Failure group: NVRTC program destruction segfaults in custom segmented-reduce test
Affected jobs:
- CCCL C Parallel v2 (HostJIT) nvcc MSVC / [CTK13.3 MSVC14.50] Test(amd64, RTX2080): https://github.com/NVIDIA/cccl/actions/runs/30832248360/job/91748951807

Investigate the Windows CTK 13.3/MSVC 14.50 crash in `cccl.c.parallel.v2.test.segmented_reduce`, specifically the custom-types case beginning at `c/parallel/test/test_segmented_reduce.cpp:359`. Reproduce narrowly with `ctest --test-dir build/cccl-c-parallel-v2 -R "^cccl.c.parallel.v2.test.segmented_reduce$" --output-on-failure`; verify whether the crash consistently occurs at `nvrtcDestroyProgram` in `c/parallel/test/test_util.h:110` after LTO IR retrieval. Obtain a native stack trace or construct a minimal NVRTC reproducer to distinguish an NVRTC/toolchain defect from memory corruption. Implement the smallest justified fix—prefer correcting an identified lifetime or ownership bug, but do not suppress cleanup or quarantine the test without evidence—and run the focused segmented-reduce test for validation.

Jobs:

• CCCL C Parallel v2 (HostJIT) nvcc MSVC / [CTK13.3 MSVC14.50] Test(amd64, RTX2080)