Skip to content

[None][fix] Pin kimiK3AttnRes to the sm_100 family - #17085

Merged
brnguyen2 merged 1 commit into
NVIDIA:feat/kimi_k3from
brnguyen2:fix-kimik3-attnres-sm120f
Jul 30, 2026
Merged

[None][fix] Pin kimiK3AttnRes to the sm_100 family#17085
brnguyen2 merged 1 commit into
NVIDIA:feat/kimi_k3from
brnguyen2:fix-kimik3-attnres-sm120f

Conversation

@brnguyen2

@brnguyen2 brnguyen2 commented Jul 30, 2026

Copy link
Copy Markdown
Collaborator

@coderabbitai summary

Description

The x86_64 wheel build fails for any build whose CUDA architecture list includes sm_120f:

ptxas ... error : Instruction 'tcgen05.st' not supported on .target 'sm_120f'
ptxas ... error : Feature '.32x32b' not supported on .target 'sm_120f'

cpp/tensorrt_llm/kernels/kimiK3AttnRes/attnResFwd.cu is warp-specialized for the sm_100 (datacenter Blackwell) family and uses tcgen05/TMEM PTX. Its code paths are guarded with __CUDA_ARCH__ >= 1000, which also admits sm_120 (__CUDA_ARCH__ == 1200), so those instructions are emitted for an architecture that does not support them. Because the object target inherited the global architecture list rather than declaring its own, every build containing sm_120f hit this. Architectures below sm_100 were unaffected: the guards exclude the tcgen05 paths entirely, so the helpers are never emitted.

This blocks the build for all pull requests targeting feat/kimi_k3.

Two changes:

  1. cpp/tensorrt_llm/kernels/kimiK3AttnRes/CMakeLists.txt — pin the object target with set_cuda_architectures(kimi_k3_attn_res_src 100f), the same mechanism other architecture-specific kernel targets use (for example fp4_gemm_src, marlin_src). On a build with no sm_100-family architecture the helper falls back to its placeholder path, which is safe here because the guards already exclude the tcgen05 code.
  2. cpp/tensorrt_llm/thop/attnResOp.cpp — the kernel binary now exists only for the sm_100 family, so the capability check requires compute capability major == 10 rather than >= 10. An sm_120 device now gets a clear TORCH_CHECK message instead of a missing-kernel-image launch failure. The helper is renamed is_sm100_family() to match its meaning.

Test Coverage

Compiled attnResFwd.cu directly with CUDA 13.2 across an architecture matrix, before and after the change:

Architecture Result Notes
sm_80 compiles tcgen05 paths excluded by the existing guards
sm_90 compiles as above
sm_100f compiles the family this kernel targets, and the one now pinned
sm_120f fails (tcgen05.st / .32x32b rejected) reproduces the reported build failure; this architecture is no longer requested for the target after the change

The compile matrix was run on an aarch64 host; device-side PTX generation and ptxas validation are independent of the host architecture, so this reproduces the failure seen in x86_64 builds. The CI build on this PR exercises the full x86_64 wheel path.

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)

  • If PR introduces API changes, an appropriate PR label is added - either api-compatible or api-breaking. For api-breaking, include BREAKING in the PR title.

  • 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.

attnResFwd.cu guards its tcgen05/TMEM code paths with `__CUDA_ARCH__ >= 1000`,
which also admits sm_120. The target inherited the global architecture list, so
builds that include sm_120f emitted tcgen05 PTX for it and ptxas rejected it:

    Instruction 'tcgen05.st' not supported on .target 'sm_120f'
    Feature '.32x32b' not supported on .target 'sm_120f'

Pin the object target to the sm_100 family with set_cuda_architectures(), the
same pattern other architecture-specific kernel targets use, and tighten the
Torch-op capability check from major >= 10 to major == 10 so unsupported
devices get a clear error instead of a missing-kernel-image launch failure.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Brian Nguyen <brnguyen@nvidia.com>
@brnguyen2
brnguyen2 requested a review from a team as a code owner July 30, 2026 23:03
@brnguyen2
brnguyen2 requested a review from zongfeijing July 30, 2026 23:03
@brnguyen2

Copy link
Copy Markdown
Collaborator Author

/bot run

@brnguyen2
brnguyen2 merged commit 40a867e into NVIDIA:feat/kimi_k3 Jul 30, 2026
8 of 9 checks passed
@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62871 [ ] completed with state FAILURE. Commit: 2baeeef
Not allowed on merged PR

Link to invocation

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants