cuTile smoke test: add sm_121 fragment for GB10, document a driver JIT-link gap by zbrad · Pull Request #2568 · NVIDIA/cuvs

zbrad · 2026-09-09T02:08:54Z

Summary

GB10 (compute capability 12.1, DGX Spark) rejects cutile_smoke's only embedded cubin (sm_120) at cudaLibraryGetKernel with cudaErrorNoKernelImageForDevice, despite matching find_compatible_cubin_fragment()'s same-major/lower-minor forward-compatibility rule. Confirmed empirically on real GB10 hardware that this "forward compatible within a major family" assumption doesn't hold here.

Adds an exact-match cutile_arch_12_1 / sm_121 fragment.
sm_121a (the family-conditional target) isn't reachable yet — tileiras in the currently-published cuda-tile doesn't support any -a suffixed --gpu-name target (confirmed by testing sm_90a/sm_100a/sm_120a/sm_121a, all rejected identically); filed as NVIDIA/cutile-python#105.
Corrects find_compatible_cubin_fragment()'s doc comment, which stated cubins are forward compatible across minor revisions within a major family — not true for this device.

A second, independent finding

Fixing the above surfaced a separate problem: actually launching any cuTile-compiled kernel on this GB10 box fails with cudaErrorJitCompilerNotFound. Traced via strace to a real, confirmed cause: libnvidia-gpucomp.so, needed at kernel-launch time to complete a runtime relocation cuTile's cubins carry (.rela.nv.constant4 against ___NV_TILE_LAUNCH..., absent from a plain nvcc-compiled cubin through the identical cudaLibraryLoadData/cudaLibraryGetKernel/cudaLaunchKernelExC sequence — that path works fine), isn't shipped by the driver branch tested against (580.173.02).

Since is_expected_cutile_unavailable() already treats cudaErrorJitCompilerNotFound as an expected-unavailable condition for the load step, LaunchesCompatibleCubin now also tolerates it at launch time (skip, not fail) for consistency. A new CutileSmoke.RequiresJitLinkCapableDriver test deliberately does not tolerate it, so this gap stays visible rather than silently skipped — it's expected to start passing once a driver ships the missing component, which is the signal to fold its guarantee back into LaunchesCompatibleCubin and delete both.

Testing

CutileSmoke.ResolvesEveryEmbeddedArchitecture — updated for the new fragment, verified passing.
CutileSmoke.LaunchesCompatibleCubin — was failing (ASSERT_NE(launcher, nullptr)), now resolves the launcher correctly (confirms the arch fix); skips cleanly on this driver due to the second issue above.
CutileSmoke.RequiresJitLinkCapableDriver (new) — fails as expected/documented on this driver branch.
Rest of the GB10 gtest suite (34/35 binaries, ~8000+ tests) run clean against this change; the one other failure (CLUSTER_TEST, a KMeans++ random-subsample tolerance flake) is unrelated and reproducibly flaky independent of this change.

Assisted-by: Claude

copy-pr-bot · 2026-09-09T02:08:57Z

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.

Copilot

🟡 Changes recommended

It introduces a test explicitly documented as “expected to FAIL” (risking CI instability) and includes a helper that can leak CUDA allocations on ASSERT_* early-exit paths.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

This PR updates the cuTile smoke-test infrastructure to properly support GB10 (SM 12.1) by embedding an exact-match sm_121 cubin fragment, and it adjusts test behavior/documentation to reflect that “same-major/lower-minor” SASS fallback is not reliable for this family.

Changes:

Add cutile_arch_12_1 / sm_121 to the embedded cubin matrix and register it in the smoke-test fragment list/planner.
Update the find_compatible_cubin_fragment() doc comment to qualify the forward-compatibility assumption for SM 12.x (GB10).
Refactor the smoke launch test to centralize kernel execution/verification and add a new driver capability test (currently described as expected to fail on some drivers).

File summaries

File	Description
`cpp/tests/detail/jit_lto/cutile_smoke.cu`	Adds SM 12.1 fragment usage; refactors launch/verify path; adds a new driver JIT-link capability test.
`cpp/src/detail/jit_lto/cutile_smoke/cutile_smoke_matrix.json`	Adds `sm_121` cubin export entry for the smoke kernel.
`cpp/include/cuvs/detail/jit_lto/cutile_module.hpp`	Updates compatibility-selection documentation to note SM 12.1 rejecting SM 12.0 SASS.
`cpp/include/cuvs/detail/jit_lto/cutile_arch_tags.hpp`	Introduces `cutile_arch_12_1` with documentation explaining why it’s required on GB10.

Review details

Files reviewed: 4/4 changed files
Comments generated: 4
Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

+  float* lhs    = nullptr;
+  float* rhs    = nullptr;
+  float* output = nullptr;
+  ASSERT_EQ(cudaMalloc(&lhs, sizeof(host_lhs)), cudaSuccess);
+  ASSERT_EQ(cudaMalloc(&rhs, sizeof(host_rhs)), cudaSuccess);
+  ASSERT_EQ(cudaMalloc(&output, sizeof(host_output)), cudaSuccess);


+// upgrade ships the missing component, this test starts passing on its own,
+// which is the signal to fold its guarantee back into LaunchesCompatibleCubin
+// and delete this test and the skip branch above.
+TEST(CutileSmoke, RequiresJitLinkCapableDriver)


+// Runs the add-two-tiles kernel and checks the result. Throws (via RTCX_CUDA_TRY,
+// see rtcx/macros.hpp) on any CUDA failure -- callers decide whether to tolerate
+// a specific expected failure or let it propagate as a test failure.


+// Deliberately NOT tolerant of cudaErrorJitCompilerNotFound, unlike
+// LaunchesCompatibleCubin above -- this is expected to FAIL on any driver
+// branch missing libnvidia-gpucomp.so (confirmed absent from the pinned
+// 580.173.02 branch via strace; only present in an untested 595.45.04
+// package cached locally, not installed). Excluded from tuned/full_test.sh's


divyegala

Thank you for catching this bug! Could you please re-target to release/26.10 branch?

divyegala · 2026-09-09T04:13:39Z

+// upgrade ships the missing component, this test starts passing on its own,
+// which is the signal to fold its guarantee back into LaunchesCompatibleCubin
+// and delete this test and the skip branch above.
+TEST(CutileSmoke, RequiresJitLinkCapableDriver)


This test needs to be opt-in, otherwise it will fail on machines that have an embedded cubin but their drivers are not JIT capable

GB10 (compute capability 12.1) rejects the existing sm_120 embedded cubin outright (cudaErrorNoKernelImageForDevice) despite matching find_compatible_cubin_fragment()'s same-major/lower-minor forward-compat rule -- confirmed empirically that rule doesn't hold for this device. Add an exact-match cutile_arch_12_1/sm_121 fragment (sm_121a is not reachable: NVIDIA/cutile-python#105) and correct find_compatible_cubin_fragment's doc comment, which overclaimed universal forward compatibility. That surfaced a second, independent problem: actually launching any cuTile-compiled kernel needs libnvidia-gpucomp.so at kernel-launch time to complete a runtime relocation (confirmed via strace), and that library genuinely isn't shipped by the currently-pinned 580.173.02 driver branch (only found, unused, in a cached 595.45.04 package). LaunchesCompatibleCubin now tolerates cudaErrorJitCompilerNotFound as an expected-unavailable skip during launch, mirroring what is_expected_cutile_unavailable() already does for the load step. A new CutileSmoke.DISABLED_RequiresJitLinkCapableDriver canary does NOT tolerate it and is expected to fail until a future driver ships the missing component -- DISABLED_ (GoogleTest's own opt-in convention) so it doesn't run by default, only via --gtest_also_run_disabled_tests; its own passing becomes the signal to fold the guarantee back into LaunchesCompatibleCubin and delete both the canary and the skip branch. Device buffers in the shared launch/verify helper are now held by unique_ptr<float, CudaDeviceDeleter> so they're freed on every exit path (including an early ASSERT_* return, which doesn't unwind via a C++ exception) rather than only on the success path.

zbrad · 2026-09-09T05:11:46Z

Thanks both — addressed everything, force-pushed:

Re-targeted to release/26.10 and reset the branch onto it directly (the diff is a clean 4-file/155-line change again, not the 51-file diff a bare base-branch swap produced against main's divergence).
Made the driver-gap test opt-in: renamed to DISABLED_RequiresJitLinkCapableDriver (GoogleTest's own convention — excluded from any default ctest/gtest run, only runs via --gtest_also_run_disabled_tests). It no longer relies on anything outside this repo to stay non-gating.
Fixed the leak: run_smoke_add_and_verify's three device buffers are now held by unique_ptr<float, CudaDeviceDeleter>, so they're freed on every exit path — including an early ASSERT_* return, which doesn't unwind via a C++ exception and was leaking them before.
Fixed the misleading doc comment: it now accurately describes that host-observable CUDA calls use ASSERT_* (early return, not a throw) while the kernel dispatch call itself throws via the vendored rtcx dependency — and clarifies rtcx/macros.hpp lives in that CPM-fetched dependency, not this repo's own tree.
Removed the tuned/full_test.sh reference — that was fork-only tooling not present here; the comment now just describes the DISABLED_ mechanism itself.

Verified with a real rebuild + compute-sanitizer --tool memcheck on GB10 hardware: default run is clean (DISABLED_RequiresJitLinkCapableDriver correctly excluded, no failures), explicit --gtest_also_run_disabled_tests --gtest_filter=*RequiresJitLinkCapableDriver run still fails exactly as documented, and the RAII change introduces no new sanitizer findings (the only two flagged events are the already-expected cudaErrorJitCompilerNotFound calls the test itself is built around).

Sync the same fixes from the upstream PR after real review feedback (divyegala + Copilot bot): - Fix a real leak: run_smoke_add_and_verify's device buffers are now held by unique_ptr<float, CudaDeviceDeleter>, freed on every exit path (an early ASSERT_* return doesn't unwind via a C++ exception and was leaking them before). - Fix a misleading doc comment claiming the helper "throws... on any CUDA failure" when most of its CUDA calls use ASSERT_* (early return, not a throw) instead. - Rename the driver-gap canary to DISABLED_RequiresJitLinkCapableDriver (GoogleTest's own opt-in convention), which supersedes this repo's custom --gtest_filter exclusion in tuned/full_test.sh -- simplified that script accordingly, and it no longer relies on fork-only tooling to stay non-gating. Verified: default run clean (DISABLED_ test excluded, no failures), --gtest_also_run_disabled_tests run still fails exactly as documented, compute-sanitizer clean (only the two already-expected cudaErrorJitCompilerNotFound events, no new findings from the RAII change).

zbrad requested a review from a team as a code owner September 9, 2026 02:08

Copilot AI lite review requested due to automatic review settings September 9, 2026 02:08

Copilot started reviewing on behalf of zbrad September 9, 2026 02:09 View session

Copilot AI reviewed Sep 9, 2026

View reviewed changes

divyegala added bug Something isn't working non-breaking Introduces a non-breaking change labels Sep 9, 2026

divyegala requested changes Sep 9, 2026

View reviewed changes

zbrad changed the base branch from main to release/26.10 September 9, 2026 05:08

zbrad requested review from a team as code owners September 9, 2026 05:08

zbrad requested a review from bdice September 9, 2026 05:08

zbrad force-pushed the gb10-cutile-sm121-fix branch from a2a79d4 to 30937fb Compare September 9, 2026 05:11

Navigation Menu

Uh oh!

cuTile smoke test: add sm_121 fragment for GB10, document a driver JIT-link gap - #2568

cuTile smoke test: add sm_121 fragment for GB10, document a driver JIT-link gap#2568
zbrad wants to merge 1 commit into
NVIDIA:release/26.10from
zbrad:gb10-cutile-sm121-fix

zbrad commented Sep 9, 2026 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Sep 9, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

divyegala left a comment

Uh oh!

divyegala Sep 9, 2026

Uh oh!

zbrad commented Sep 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

zbrad commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

A second, independent finding

Testing

Uh oh!

copy-pr-bot Bot commented Sep 9, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

divyegala left a comment

Choose a reason for hiding this comment

Uh oh!

divyegala Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

zbrad commented Sep 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

zbrad commented Sep 9, 2026 •

edited

Loading