SM103 (B300 / Blackwell Ultra) is missing from five SM allowlists across the build and the attention path

[filipemartinsubrobotics]

Repo: NVIDIA/TensorRT-Edge-LLM · reproduced on main and on v0.10.1 (e8b2952)
Hardware: NVIDIA B300 SXM6, compute capability 10.3, x86_64, CUDA 13.0, TensorRT 10.13.2
Related: #204 / #205 (same toolchain path, verified working on SM89)

Building and loading the INT4-AWQ path for a B300 fails at five separate SM allowlists. Each lists
neighbouring Blackwell parts and omits 103 alone. Across the whole C++ tree, SM103 appears in exactly one
predicate (cuteDslF16MoeRunner.cpp:192) and is absent from every attention and decode gate — so this reads as
SM103 having been added to some tables and not others, rather than as a deliberate exclusion.

1. kernelSrcs/build_cutedsl.py — int4_fp16_gemm

$ python kernelSrcs/build_cutedsl.py --kernels int4_fp16_gemm --gpu_arch sm_103
ValueError: No variants in groups ['int4_fp16_gemm'] support SM103.

All 24 variants excluded, including all 8 int4_fp16_gemv_m* decode kernels. Without the artifact, cmake fails
in cute_dsl_setup and NvInfer_edgellm_plugin cannot be built at all.

supported_sms=[80, 86, 87, 89, 100, 101, 110, 120, 121]     # int4_fp16_gemm — no 103
supported_sms=[100, 101, 103, 110]                          # Blackwell-DC groups elsewhere — has 103

The group's own comment says the kernels are portable: "Ampere instruction floor (cp.async + mma.sync 16x8x16 +
ldmatrix), forward-compatible to SM80 and newer (Ampere / Ada / Hopper / Blackwell)", and the tuning scope is
"Orin (SM87), Thor (SM110), and DGX" — so the list appears to encode which SKUs were tuned, not which are
capable.

Verified: adding 103 compiles all 24 variants cleanly, including every GEMV, producing
libcutedsl_x86_64.a under cpp/kernels/cuteDSLArtifact/x86_64/sm_103. No other source change needed.

2. kernelSrcs/build_cutedsl.py — fmha

Same shape, same omission. Verified: adding 103 compiles the whole group (24 MB archive, all variants).

3, 4, 5. The attention path: three more gates

// cpp/kernels/decodeAttentionKernels/decoderXQAJitCompiler.cpp:254
constexpr std::array<int32_t, 9> kALLOWED_SM_VERSIONS{80, 86, 87, 89, 90, 100, 101, 120, 121};

103 absent, though 100, 101, 120 and 121 are all present. The adjacent comment explains the deliberate choice of
an explicit list over a numeric threshold — which is reasonable, and means new SMs have to be added by hand.

With 1 and 2 patched the plugin builds, registers, and Int4GroupwiseGemmPluginV2 is accepted. Parsing then
reaches AttentionPlugin, which rejects the layer because neither backend claims it:

[attentionPlugin.cpp:753] AttentionPlugin: no XQA decode kernel for Hq=16, Hkv=8, headSize=128 on SM103; skipping XQA JIT.
[attentionPlugin.cpp:613] Cannot implement AttentionPlugin configuration. SM: 103, HeadSize: 128, NumQHeads: 16, NumKVHeads: 8

canCompileXQAKernel returns false on the allowlist above. mCanImplementFMHA is false for the same reason, two
gates further up — selectFMHAKernels (attentionPlugin.cpp) tries the Blackwell runner, then FMHA-V2, and both
reject SM103:

// cpp/kernels/contextAttentionKernels/cuteDslFMHARunner.cpp:41
bool isSupportedBlackwellFmha(int32_t smVersion)
{ return smVersion == 100 || smVersion == 101 || smVersion == 110; }          // no 103

// cpp/kernels/contextAttentionKernels/cuteDslFMHAV2Runner.cpp:64
bool isFMHAV2SM(int32_t smVersion)
{ return smVersion == 80 || smVersion == 86 || smVersion == 87 || smVersion == 89 || smVersion == 90
      || smVersion == 100 || smVersion == 101 || smVersion == 110 || smVersion == 120 || smVersion == 121; }  // no 103

So for headSize=128, Hq=16, Hkv=8 the sequence is: Blackwell FMHA declines (gate 3), FMHA-V2 declines
(gate 4), XQA declines (gate 5), and the constructor throws. Note applyThorSMRenumberWAR immediately above the
selection call shows there is already precedent for mapping one SM onto another's kernels.

Separately: the rejection crashes rather than raising

attentionPlugin.cpp:613 throws std::runtime_error from the plugin constructor. Through the C ABI during ONNX
parsing this becomes:

Segmentation fault (core dumped)

An unsupported configuration should surface as a parser error naming the layer, not a segfault. This is a distinct defect
from the SM coverage above and would be worth fixing regardless — it is what makes the underlying cause hard
to find, exactly as the "Could not create the plugin" masking described in #204.

Suggested change

# build_cutedsl.py, int4_fp16_gemm and fmha
supported_sms=[80, 86, 87, 89, 100, 101, 103, 110, 120, 121]

// decoderXQAJitCompiler.cpp:254
constexpr std::array<int32_t, 10> kALLOWED_SM_VERSIONS{80, 86, 87, 89, 90, 100, 101, 103, 120, 121};

// cuteDslFMHAV2Runner.cpp:64  — add `|| smVersion == 103`
// cuteDslFMHARunner.cpp:41    — add `|| smVersion == 103`, if the Blackwell kernels are in fact 103-capable

If any of these omissions is deliberate — a known correctness or performance problem on Blackwell Ultra — a note
in the comment would help, since the surrounding text reads as an explicit compatibility guarantee.

What we verified, and what we did not

• Verified: with 1 and 2 patched, all INT4 and FMHA variants compile for SM103; the plugin library builds and
  links; its creators register with TensorRT 10.13.2; Int4GroupwiseGemmPluginV2 is accepted during parsing.
• Not verified: numerical correctness or performance of any of these kernels on SM103. The config sweep was
  tuned for SM87 / SM110 / DGX. We are not claiming the kernels are correct or fast on this SKU — only that they
  build and load, and that the allowlists exclude a part the surrounding code describes as compatible.
• Not reached: engine build. Gates 3, 4 and 5 block it.

Incidental: exact pins for an x86_64 + CUDA 13 build

• nvidia-cutlass-dsl: found 4.7.1, need 4.7.0
• CuTe DSL loaded the cu12 compiler backend, but artifact CUDA 13 was requested.
• cupy not found.   Fix: pip install cupy-cuda13x==13.6.0

pip install 'nvidia-cutlass-dsl[cu13]==4.7.0' 'cupy-cuda13x==13.6.0' resolves all three. This message only
appears after variant selection succeeds, so on an unsupported SM you hit the ValueError in item 1 first and
never learn the toolchain is also wrong.

Also worth a line in the build docs: cmake needs -DCUDA_CTK_VERSION=13.0 -DCUDA_DIR=/usr/local/cuda-13.0 -DCUDA_RUNTIME_API_INCLUDE_DIR=/usr/local/cuda-13.0/targets/x86_64-linux/include on CUDA 13, where headers are
under targets/<arch>/include rather than include/. Without them find_path leaves
CUDA_RUNTIME_API_INCLUDE_DIR-NOTFOUND and the build fails much later at generatePluginJitEmbeddedSources with
"embedded source(s) not found", which points nowhere near the cause.

---

Found during the NVIDIA / OpenHackathons / Oracle Open Models Codefest 2026,
while bringing up an INT4-AWQ Cosmos3-Edge reasoner on a DGX B300 for an
offline-first search-and-rescue robotics entry (Team UBR Stack). The production
target is a Jetson Orin Nano; the B300 is a bench machine used for evaluation.