[None][fix] Restore SM107 2x-mmaK acceptance and fine-grained sync PDL path

[farazkh80]

Summary

Restores SM107 fine-grained (Lamport) MoE sync (#17707) to actually dispatch on Rubin, plus a
kernel-hang workaround discovered while validating it.

1. SM107 2x-mmaK acceptance was dropped from the public export

trtllmGen_bmm_export/GemmOptions.h::checkAndUpdateGemmOptions no longer accepts mmaK=64
(FP8/MxFP8) or mmaK=128 (FP4) on SM107a — those checks were guarded internally and stripped by
the export filter. Every SM107a batched-GEMM cubin with these mmaK values therefore fails
validation, so getValidConfigIndices/getDefaultValidConfigIndex throw and the MoE runners
silently fall back to SM100f kernels. This makes use_fine_grained_sync a no-op for
NVFP4/FP8/MxFP8 MoE on SM107 — no SM107a-specific MoE GEMM cubin is ever selected, regardless
of the flag, and the *_fine_grained_sync accuracy tests pass vacuously.

2. Fine-grained consumer launched without PDL

BatchedGemmInterface.h::run's pdlSafe check didn't include
mFineGrainedConsumerA/mFineGrainedConsumerB, so the FC2 consumer always launched without
programmatic dependent launch and could never start before FC1 completes — defeating the
producer/consumer overlap the protocol is for.

3. FC1 validity check used an unrounded valid hidden size

PermuteGemm1::Runner::{getDefaultValidConfigIndex,isValidConfigIndex} validate against
validHiddenSize as passed in, while run() rounds it up to 512 for MxE4m3 activations before
launching. Made the two consistent.

4. SM107 2x-mmaK MxE2m1 padding rule was dropped from the exports

Once (1) is fixed, every MxFp8Fp6Fp4 (MxE4m3 activation) mmaK=64 SM107a batched-GEMM kernel hung
on its first launch, fine-grained sync on or off. Root cause: the export also stripped the Rubin
branch of trtllm/gen/DtypeDecl.h::dtypeNeedsPadding (and KernelTraits.h::getNumSmemBitsPerElt),
so the host built the MxE2m1 weight tensor map in padded-FP4 mode while the cubin, generated with
Rubin features, consumes unpadded data. The TMA load never completes the mbarrier transaction and
every warp spins in SYNCS.PHASECHK.TRANS64.TRYWAIT (cuda-gdb). Proven by dumping the launch
parameters on this branch vs the internal branch for the identical kernel: only the operand-A tensor
map differed, by the padding bit. Restored in both trtllmGen_bmm_export and trtllmGen_gemm_export
(the dense export ships 12 sm107a mmaK=64 MxE2m1 kernels with the same exposure). The interim
skipQuirks family ban from the first revision of this PR is removed.

Test coverage

Verified on a 4x Rubin (SM107) box via tests/microbenchmarks/bench_moe and trtllm-gen's
standalone BatchedGemm FC1→FC2 harness:

• NVFP4 (DeepSeek-V3 shapes, --quant NVFP4, SM100f excluded to force full SM107a coverage):
  all 65 distinct SM107a FP4 cubin variants dispatch across 4/32/128/512 tokens, including the
  fine-grained producer (_fgFV_fgP) / consumer (_fgCB) pair. No hang.
• GPT-OSS w4a8 (--quant W4A8_MXFP4_MXFP8, use_fine_grained_sync=True): completes cleanly
  across 4/32/128 tokens with zero SM107a kernels selected (clean fallback via the skipQuirks
  entry in (4)), matching pre-[TRTLLM-15316][feat] Rubin trtllmgen batchedGemm MoE #17707 behavior — no hang, no regression.
• trtllm-gen standalone repro (added locally, not part of this diff, source at
  kernels/BatchedGemm/BatchedGemmFc1Fc2LamportRepro.cpp in the trtllm-gen tree): 2000 replays of
  the FC1(tileK=512,splitK=2)/FC2(tileK=256) pair with fine-grained sync off completed cleanly;
  the same pair with fine-grained sync on requires internal SASS-patch tooling
  (nvdisasm_internal/nvasm_internal) not available in this environment, so it could not be
  built/run standalone, but the TRT-LLM-level CUDA_LAUNCH_BLOCKING isolation above used the real
  production cubins directly and is conclusive for that case.
• Confirmed kernel selection via TLLM_BATCHED_GEMM_PRINT_NAME=1 logging in all cases.

I don't have a way to run the standard CI test suite from this environment; happy to address CI
feedback.

Dev Engineer Review

• Restores SM107a support for doubled mmaK configurations in GemmOptions.
• Includes fine-grained MoE consumers in the PDL safety check.
• Aligns MXFP4 validity checks with launch-time hidden-size rounding.
• Restores native shared-memory element sizes for applicable MxFP8/MxFP6/MxFP4 kernels in batched-GEMM and GEMM exports.
• Changes missing-runner diagnostics from debug logs to warnings.
• Main risk: SM107a MxFP8 mmaK=64 kernels remain implementation-sensitive and require continued skipQuirks management.
• Standard CI was not run. The validated SM107 Rubin workloads completed with NVFP4 dispatch and GPT-OSS fallback behavior.

QA Engineer Review

No test changes.

Per-File QA Perspective

• cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/BatchedGemmInterface.h: Verify PDL safety decisions for fine-grained producer and consumer combinations.
• cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/GemmOptions.h: Verify SM107a accepts the intended FP8/MxFP8/FP4/MxFP4 mmaK configurations and rejects them on unsupported architectures.
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu: Verify configuration selection matches launch-time hidden-size rounding for MXFP4 weights with E4M3 activations.
• cpp/tensorrt_llm/thop/mxFp4BlockScaleMoe.cpp: Verify missing-runner diagnostics appear at warning level for BF16 and MXE4m3 paths.
• cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/KernelTraits.h: Verify native dtype-width shared-memory calculations for dense mmaK >= 64 and sparse mmaK >= 128.
• cpp/tensorrt_llm/kernels/trtllmGenKernels/gemm/trtllmGen_gemm_export/KernelTraits.h: Verify GEMM shared-memory calculations remain unchanged for non-Rubin and lower-mmaK configurations.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: cb1e1969-8cf6-4922-b99f-02894bd762ff

📥 Commits

Reviewing files that changed from the base of the PR and between f8e1eef and 24e7d79.

⛔ Files ignored due to path filters (2)

• cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/trtllm/gen/DtypeDecl.h is excluded by !**/gen/**
• cpp/tensorrt_llm/kernels/trtllmGenKernels/gemm/trtllmGen_gemm_export/trtllm/gen/DtypeDecl.h is excluded by !**/gen/**

📒 Files selected for processing (2)

• cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/KernelTraits.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/gemm/trtllmGen_gemm_export/KernelTraits.h

Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.

---

Walkthrough

Changes

GEMM runtime updates

Layer / File(s)	Summary
SM107a GEMM validation, sizing, and safety cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/GemmOptions.h, cpp/tensorrt_llm/kernels/trtllmGenKernels/*/trtllmGen_*_export/KernelTraits.h, cpp/tensorrt_llm/kernels/trtllmGenKernels/batchedGemm/trtllmGen_bmm_export/BatchedGemmInterface.h	SM107a supports additional MMA-K configurations. Shared-memory sizing uses native dtype widths for qualifying kernels. PDL safety checks fine-grained A or B consumers.
MXFP4 configuration alignment cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu	MXFP4 with E4M3 activations rounds validHiddenSize to a multiple of 512 before configuration selection and validation.
Runner diagnostic severity cpp/tensorrt_llm/thop/mxFp4BlockScaleMoe.cpp	Missing BF16 and MXE4M3 runners now produce warning-level diagnostics.

Estimated code review effort: 3 (Moderate) | ~20 minutes

Severity of issue fixed: Low

Suggested reviewers: bowenfu

Merge Risk: ⚪ Minimal · up to 24e7d

This restores SM107a GEMM/MoE configuration support, alignment handling, and diagnostics. No current merge-readiness risk remains.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 31.58% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff. Analyzed 19 functions across 7 files.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Title check	✅ Passed	The title clearly identifies the main changes: restoring SM107 2x-mmaK acceptance and enabling the fine-grained sync PDL path.
Description check	✅ Passed	The description clearly explains the problem, solution, affected code paths, and test coverage. It uses a Summary heading instead of Description and does not include the PR Checklist, but the required…

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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

[farazkh80]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #72787 [ run ] triggered by Bot. Commit: cd506e4 Link to invocation