Skip to content

[#12230][fix] Add bounds checking in autotuner _find_nearest_profile for SM121 - #12310

Merged
karljang merged 31 commits into
NVIDIA:mainfrom
mihai-chiorean:fix/autotuner-indexerror-sm121
Jun 11, 2026
Merged

[#12230][fix] Add bounds checking in autotuner _find_nearest_profile for SM121#12310
karljang merged 31 commits into
NVIDIA:mainfrom
mihai-chiorean:fix/autotuner-indexerror-sm121

Conversation

@mihai-chiorean

@mihai-chiorean mihai-chiorean commented Mar 18, 2026

Copy link
Copy Markdown
Contributor

Summary

Fixes #12230 (autotuner crash)

On DGX Spark (SM121), _find_nearest_profile() crashes with IndexError: list assignment index out of range during warmup, preventing the server from starting. This happens because DynamicTensorSpec.input_idx or dim_idx exceeds the base_profile dimensions when ops produce fewer or differently-shaped tensors than the specs expect.

Adds bounds checks before indexing into base_profile for both DynamicTensorSpec and ConstraintSpec loops. Out-of-bounds specs are skipped with a debug log. All in-bounds behavior is unchanged.

Test Coverage

Tested with 8 cases inside the Docker container:

  • Normal operation (in-bounds specs) — unchanged behavior
  • input_idx out of range for DynamicTensorSpec — skipped cleanly
  • dim_idx out of range for DynamicTensorSpec — skipped cleanly
  • input_idx/dim_idx out of range for ConstraintSpec — skipped cleanly
  • Mixed valid + out-of-bounds specs — valid ones still applied correctly
  • tune_max_num_tokens with out-of-bounds spec — handled
  • Empty shapes with non-empty specs — handled

Test plan

  • 8 unit test cases pass (bounds checking, mixed valid/invalid, edge cases)
  • Pre-commit checks pass

Note: Full autotuner integration test requires model serving, deferred to NVIDIA CI.

PR Checklist

Please review the following before submitting your PR:

  • PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.

  • PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.

  • Test cases are provided for new code paths (see test instructions)

  • Any new dependencies have been scanned for license and vulnerabilities

  • CODEOWNERS updated if ownership changes

  • Documentation updated as needed

  • Update tava architecture diagram if there is a significant design change in PR.

  • The reviewers assigned automatically/manually are appropriate for the PR.

  • Please check this after reviewing the above items as appropriate for this PR.

Summary by CodeRabbit

  • Bug Fixes
    • Enhanced autotuner robustness with defensive bounds validation to prevent out-of-bounds access errors during profile matching operations.

@mihai-chiorean
mihai-chiorean requested a review from a team as a code owner March 18, 2026 06:08
@coderabbitai

coderabbitai Bot commented Mar 18, 2026

Copy link
Copy Markdown
Contributor
📝 Walkthrough

Walkthrough

Defensive bounds checks added to _find_nearest_profile in the autotuner to prevent IndexError exceptions when accessing out-of-bounds indices in DynamicTensorSpec and ConstraintSpec entries during profile matching. Skipped specs are logged and processing continues.

Changes

Cohort / File(s) Summary
Autotuner Bounds Checks
tensorrt_llm/_torch/autotuner.py
Added bounds validation for input_idx and dim_idx in both DynamicTensorSpec and ConstraintSpec processing within _find_nearest_profile to prevent out-of-bounds list access. Out-of-range specs are logged at debug level and skipped via early continue statements.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

🚥 Pre-merge checks | ✅ 5
✅ Passed checks (5 passed)
Check name Status Explanation
Title check ✅ Passed The title accurately identifies the primary fix: adding bounds checking to the _find_nearest_profile method to resolve the SM121 autotuner crash, matching the issue #12230.
Description check ✅ Passed The description comprehensively covers the issue (IndexError in autotuner), root cause (out-of-bounds indexing), solution (bounds checks), and test coverage with 8 test cases. All checklist items are addressed.
Linked Issues check ✅ Passed The code changes directly address the primary objective from issue #12230: preventing the autotuner IndexError by adding bounds checks before indexing base_profile in _find_nearest_profile, allowing graceful skipping of out-of-bounds specs.
Out of Scope Changes check ✅ Passed All changes are scoped to _find_nearest_profile bounds checking; no unrelated modifications to other components are present. The PR stays focused on the IndexError fix objective.
Docstring Coverage ✅ Passed Docstring coverage is 100.00% which is sufficient. The required threshold is 80.00%.

✏️ Tip: You can configure your own custom pre-merge checks in the settings.

✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests
📝 Coding Plan
  • Generate coding plan for human review comments

Comment @coderabbitai help to get the list of available commands and usage tips.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Actionable comments posted: 1

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)
tensorrt_llm/_torch/autotuner.py (1)

1406-1435: ⚠️ Potential issue | 🟡 Minor

Formatting check is failing in this hunk (YAPF).

Release Checks report this region was reformatted by the hook. Please run and commit formatter output for this file before merge.

🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.

In `@tensorrt_llm/_torch/autotuner.py` around lines 1406 - 1435, The formatting in
tensorrt_llm/_torch/autotuner.py (around the loop handling DynamicTensorSpec and
constraint_specs, i.e., symbols like DynamicTensorSpec,
apply_map_to_tuning_buckets, tune_max_num_tokens, and constraint_specs) was
changed by the pre-commit hook (YAPF); run the repository's YAPF formatter (or
the configured formatter) on this file, review the resulting
whitespace/line-wrapping for the shown hunk, stage and commit the formatted file
so the reformatting is included in the PR.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.

Inline comments:
In `@tensorrt_llm/_torch/autotuner.py`:
- Around line 1398-1410: The _optimization_profiles function still indexes specs
without bounds checks and can raise IndexError when specs are malformed; update
_optimization_profiles to perform the same input/dimension guards used
elsewhere: before any access to base_profile[spec.input_idx] or
base_profile[spec.input_idx][spec.dim_idx] (and similarly for
spec.dim_idx+offset), verify spec.input_idx < len(base_profile) and spec.dim_idx
< len(base_profile[spec.input_idx]) (and handle negative indices if relevant),
log a debug message referencing DynamicTensorSpec/spec.input_idx/spec.dim_idx
and skip that spec (continue) when out of range, ensuring all spots where
base_profile is indexed (around the existing accesses you noted) are protected.

---

Outside diff comments:
In `@tensorrt_llm/_torch/autotuner.py`:
- Around line 1406-1435: The formatting in tensorrt_llm/_torch/autotuner.py
(around the loop handling DynamicTensorSpec and constraint_specs, i.e., symbols
like DynamicTensorSpec, apply_map_to_tuning_buckets, tune_max_num_tokens, and
constraint_specs) was changed by the pre-commit hook (YAPF); run the
repository's YAPF formatter (or the configured formatter) on this file, review
the resulting whitespace/line-wrapping for the shown hunk, stage and commit the
formatted file so the reformatting is included in the PR.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

Run ID: 17232aea-4e81-47a4-b8c0-3d24a43abe55

📥 Commits

Reviewing files that changed from the base of the PR and between bf90b54 and 38c5479.

📒 Files selected for processing (1)
  • tensorrt_llm/_torch/autotuner.py

Comment thread tensorrt_llm/_torch/autotuner.py
@svc-trtllm-gh-bot svc-trtllm-gh-bot added the Community want to contribute PRs initiated from Community label Mar 18, 2026
@mihai-chiorean
mihai-chiorean force-pushed the fix/autotuner-indexerror-sm121 branch 2 times, most recently from f5af0fe to 775662c Compare March 23, 2026 23:03
@pengbowang-nv

Copy link
Copy Markdown
Collaborator

Hi @pamelap-nvidia , Could you take a look at this and see if it can fix #12230? Also @hyukn do you think we need such guard for autotuner? Thanks!

@mihai-chiorean
mihai-chiorean force-pushed the fix/autotuner-indexerror-sm121 branch from 775662c to a894178 Compare March 25, 2026 15:35
@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

@pengbowang-nv Yes — this fixes #12230. The bounds checks cover both _find_nearest_profile (runtime cache-key path) and _optimization_profiles (tuning-time path), which was the CodeRabbit flag. 8 bounds guards total across both functions.

Tested on DGX Spark SM121 with nvidia/Qwen3-30B-A3B-NVFP4 (30B MoE, 256 experts, NVFP4). Without this fix, model loading crashes with IndexError in _find_nearest_profile when CuteDslFusedMoENvfp4 creates a ConstraintSpec referencing an input index that doesn't exist in the base profile. With the fix, the out-of-bounds spec is skipped with a debug log and the model loads successfully (83 MoE unit tests pass, 0 fail).

Unit tests for the bounds checking are in commit 2 (tests/unittest/_torch/test_autotuner_bounds.py — 7 test cases covering both functions).

Will also test with nvidia/Qwen3-Next-80B-A3B-Thinking-NVFP4 (the model from #12230) to confirm the fix covers that specific case.

…ofile for SM121

On DGX Spark (SM121), _find_nearest_profile() crashes with IndexError
when spec.input_idx or spec.dim_idx exceeds the base_profile dimensions.
This happens because some ops produce fewer or differently-shaped tensors
than the specs expect.

Add bounds checks before indexing into base_profile for both
DynamicTensorSpec and ConstraintSpec loops. Out-of-bounds specs are
skipped with a debug log message. All in-bounds behavior is unchanged.

Tested with 8 cases covering normal operation, out-of-bounds indices,
mixed valid/invalid specs, and edge cases.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
…M121

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@mihai-chiorean
mihai-chiorean force-pushed the fix/autotuner-indexerror-sm121 branch from a894178 to 6d9c094 Compare March 25, 2026 19:34
@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

Follow-up on nvidia/Qwen3-Next-80B-A3B-Thinking-NVFP4 testing: the model fits on DGX Spark (50.8 GB) but modeling_qwen3_next.py uses a custom Triton kernel (fused_qkvzba_split_reshape_cat_kernel) that hits an illegal instruction on aarch64/SM121 (Bug 2 in #12230). There's no non-Triton fallback in the model, so it can't run end-to-end on Spark right now.

That said, the autotuner IndexError (Bug 1 in #12230) is hit during model initialization before the Triton kernel executes — so this PR would let users get past the autotuner crash. The stack trace in #12230 (_find_nearest_profile, line 1406, base_profile[spec.input_idx][spec.dim_idx] = -1) is exactly the code path our bounds checks guard.

Validated with nvidia/Qwen3-30B-A3B-NVFP4 which exercises the same autotuner code path without the Triton dependency.

@johnnynunez

johnnynunez commented Mar 27, 2026

Copy link
Copy Markdown

While running nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-NVFP4 on DGX Spark with vLLM (built with TORCH_CUDA_ARCH_LIST=12.1a, CUDA 13.2), the FlashInfer autotuner logs repeated warnings during warmup:

[Autotuner]: Skipping tactic ... due to failure while profiling:
[TensorRT-LLM][ERROR] Assertion failed: Failed to initialize cutlass TMA WS grouped gemm.
Error: Error Internal (cutlass_kernel_file_gemm_grouped_sm120_M128_BS_group2.generated.cu:60)

The failing kernels are TRT-LLM's CUTLASS TMA warp-specialized MoE grouped GEMM kernels compiled for cutlass::arch::Sm120, specifically:

  • gemm_grouped_sm120_M128_BS_group2 (CTA shape 128x256x128)
  • gemm_grouped_sm120_M256_BS_group0 (CTA shape 256x128x128)

Root cause

SM121 (DGX Spark / GB10) has less shared memory than SM120 (RTX 5090). These larger tile configurations (128x256x128B, 256x128x128B) exceed SM121's ~99KB SMEM limit for grouped GEMM TMA initialization. The smaller tile CtaShape128x128x128B would fit, but it's currently commented out in TRT-LLM due to an old CUTLASS bug that has since been fixed.

Impact

Non-blocking. The FlashInfer autotuner correctly skips the failing tactics and selects working ones. The model loads and serves successfully. These are warnings, not crashes.

cc @eugr log

@pengbowang-nv

Copy link
Copy Markdown
Collaborator

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #40638 [ run ] triggered by Bot. Commit: 416474c Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #40638 [ run ] completed with state SUCCESS. Commit: 416474c
/LLM/main/L0_MergeRequest_PR pipeline #31677 completed with status: 'SUCCESS'

CI Report

Link to invocation

@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

@pengbowang-nv I think this is good to go. I see you tagged @pamelap-nvidia and @hyukn , any chance you all can take a look? happy to address any feedback.

Comment thread tests/unittest/_torch/test_autotuner_bounds.py Outdated
Comment thread tensorrt_llm/_torch/autotuner.py Outdated
…om test docstring

- Change all 8 bounds-check log messages from debug to warning level
  so users are notified when DynamicTensorSpec/ConstraintSpec entries
  are skipped due to out-of-bounds input_idx or dim_idx.
- Remove SM121-specific reference from test docstring since the bounds
  checking is not architecture-specific.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

@hyukn @pamelap-nvidia I fixed the suggestions and a couple of more small things. hopefully we're good this time 🤞

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #51158 [ run ] completed with state SUCCESS. Commit: eb548eb
/LLM/main/L0_MergeRequest_PR pipeline #40592 completed with status: 'SUCCESS'

CI Report

Link to invocation

@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

PR_Github #51158 [ run ] completed with state SUCCESS. Commit: eb548eb /LLM/main/L0_MergeRequest_PR pipeline #40592 completed with status: 'SUCCESS'

CI Report

Link to invocation

@hyukn @yizhang-nv this now has a clean CI pass. any objections to merging it?

…ecks

Per @hyukn review on PR NVIDIA#12310: deduplicate the input_idx / dim_idx
bounds-check pattern between AutoTuner._optimization_profiles and
AutoTuner._find_nearest_profile. Both sites now call a single helper.

No behavioral change.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
Per @hyukn review on PR NVIDIA#12310: collapse the parametrically-similar
OOB cases into a single parametrized test, keeping the valid-spec case
separate for readability. Same behavioral assertions, fewer test
functions.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
Completes the symmetric matrix of (spec type) x (oob input, oob dim,
negative input, negative dim). One additional parametrize case in
test_oob_spec_skipped; no other changes.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@karljang

karljang commented Jun 9, 2026

Copy link
Copy Markdown
Collaborator

Could you please address the pre-commit failures, @mihai-chiorean ?
Thank you~

Pre-commit yapf reformatted the new pytest.param entries and
_find_nearest_profile calls into compact form. No behavioral change.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

Could you please address the pre-commit failures, @mihai-chiorean ?

Thank you~

Oops. Done!

@hyukn

hyukn commented Jun 10, 2026

Copy link
Copy Markdown
Collaborator

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53262 [ run ] triggered by Bot. Commit: 9993947 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53262 [ run ] completed with state SUCCESS. Commit: 9993947
/LLM/main/L0_MergeRequest_PR pipeline #42454 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

@hyukn hyukn left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM, only some minor changes required. Thanks for the effort!

Comment thread tensorrt_llm/_torch/autotuner.py Outdated
Comment thread tests/unittest/_torch/misc/test_autotuner.py Outdated
Per @hyukn review: drop the standalone spec_kind: str parameter from
_spec_in_bounds and derive the class name from the spec itself. Call
sites are no longer required to pass the literal class name string.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
mihai-chiorean added a commit to mihai-chiorean/TensorRT-LLM that referenced this pull request Jun 10, 2026
…ntry points

Per @hyukn review:
- Drop the redundant 'from PR NVIDIA#12310' header comment.
- Collapse TestFindNearestProfileBounds and TestOptimizationProfilesBounds
  into one TestSpecBoundsChecking class.
- Single parametrized test_oob_spec_skipped covers both _find_nearest_profile
  and _optimization_profiles paths (entry x spec_config matrix).
- Keep test_valid_specs_unaffected as a sanity check that valid specs
  are NOT skipped.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
…ntry points

Per @hyukn review:
- Drop the redundant 'from PR NVIDIA#12310' header comment.
- Collapse TestFindNearestProfileBounds and TestOptimizationProfilesBounds
  into one TestSpecBoundsChecking class.
- Single parametrized test_oob_spec_skipped covers both _find_nearest_profile
  and _optimization_profiles paths (entry x spec_config matrix).
- Keep test_valid_specs_unaffected as a sanity check that valid specs
  are NOT skipped.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@mihai-chiorean
mihai-chiorean force-pushed the fix/autotuner-indexerror-sm121 branch from a83d2f3 to 2711af0 Compare June 10, 2026 17:39
Honor the 'compress to 1 test' request literally. TestSpecBoundsChecking
now has a single parametrized test_oob_spec_skipped covering both
_find_nearest_profile and _optimization_profiles.

Signed-off-by: Mihai Chiorean <mihai.v.chiorean@gmail.com>
@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

LGTM, only some minor changes required. Thanks for the effort!

done

@hyukn

hyukn commented Jun 11, 2026

Copy link
Copy Markdown
Collaborator

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53420 [ run ] triggered by Bot. Commit: 3ab45a7 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53420 [ run ] completed with state FAILURE. Commit: 3ab45a7
/LLM/main/L0_MergeRequest_PR pipeline #42591 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

@hyukn

hyukn commented Jun 11, 2026

Copy link
Copy Markdown
Collaborator

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53534 [ run ] triggered by Bot. Commit: 3ab45a7 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #53534 [ run ] completed with state SUCCESS. Commit: 3ab45a7
/LLM/main/L0_MergeRequest_PR pipeline #42688 completed with status: 'SUCCESS'

CI Report

Link to invocation

@mihai-chiorean

Copy link
Copy Markdown
Contributor Author

thanks @hyukn @pamelap-nvidia @karljang !! really appreciate your patience and reviews! looks like this is now ready.

@karljang

Copy link
Copy Markdown
Collaborator

Yeah, finally! Thank you so much for your contribution!

@karljang
karljang merged commit d3748a3 into NVIDIA:main Jun 11, 2026
7 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Community want to contribute PRs initiated from Community

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Bug]: Qwen3-Next-80B FP8/NVFP4 — Autotuner IndexError and Triton Illegal Instruction on DGX Spark (aarch64 Blackwell, TRT-LLM 1.3.0rc7)

8 participants