Spec Sheets Are Not Kernels: An ISA- and Source-Level Audit of INT8 Availability on NVIDIA Blackwell Ultra

Teng-Ruei Chen

Spec Sheets Are Not Kernels: An ISA- and Source-Level Audit of INT8 Availability on NVIDIA Blackwell Ultra

T0 review · 1 major / 4 minor · reviewed 2026-08-16 · deepseek-v4-flash

Challenge this review Re-run · record.json Download PDF Read on arXiv ↗ Review JSON

Pith's one-line read On Blackwell Ultra, INT8 appears on the datasheet but is withdrawn from the ISA, kernel library, and serving engines.

desk verdict A careful, well-pinned audit showing INT8 W8A8 is effectively unavailable on B300 across the whole stack; the ISA absence claim should be tested with a ptxas compile before someone builds on it. read the letter →

arxiv 2608.11693 v2 pith:P7SLECKW submitted 2026-08-12 cs.AR

Teng-Ruei Chen This is my paper · ORCID

classification cs.AR

keywords INT8quantizationW8A8BlackwellUltraB300PTXISAtensorcoresLLMservingkernelavailability

verification ladder T0 review T1 audit T2 compute T3 formal

The pith

A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.

The reading

On NVIDIA Blackwell Ultra (B300), the paper argues, INT8 W8A8 quantization is effectively unavailable by default even though the datasheet lists an INT8 tensor-core peak. The published dense ratio is $\sim 30{:}1$ between FP8 and INT8, where H200 and B200 are $1{:}1$; more importantly, the audit finds the format withdrawn at four layers: the PTX ISA never exposes the fifth-generation integer tensor-core instruction on the B300 architecture target, NVIDIA's own CUTLASS kernel generator skips INT8 UMMA for B300 builds, and the two major open-source LLM serving engines (vLLM and SGLang) ship no Blackwell-capable INT8 GEMM. The practical upshot is that a W8A8 checkpoint that passes every load-time check dies at the first forward pass with a hard runtime error, after the full model has been downloaded and loaded. The paper's larger claim is that format availability is a property of the whole stack, not of the model or the spec sheet, and that locally rational decisions at unrelated layers compose into de facto unavailability.

What carries the argument

The load-bearing object is the PTX instruction tcgen05.mma and its .kind qualifier, the fifth-generation tensor-core matrix-multiply family introduced with Blackwell, specifically the integer kind .kind::i8. The paper's argument turns on the ISA manual's per-kind target list: because sm_103a is absent from that list while being present in the same revision's FP4 extension rows, the fifth-generation integer path is declared nonexistent on B300, leaving only the legacy warp-level IMMA family as an architecturally legal integer tensor-core path. That single absence then explains the software layers: CUTLASS's generator guard, vLLM's null-pointer dispatch, and SGLang's SM90-stopping kernel all mirror the ISA carve-out.

What would settle it

Run ptxas or the equivalent assembler on a minimal kernel containing tcgen05.mma with .kind::i8 targeted at sm_103a; if it assembles and executes with nonzero tensor-core activity, the central absence claim is wrong. A simpler check is to inspect the next PTX ISA release for sm_103a in the .kind::i8 support statement.

Watch

Extended reading notes

Core claim

The central claim is that the integer tensor-core path for INT8 on NVIDIA Blackwell Ultra was never completed. The PTX ISA v9.3 support statement for tcgen05.mma with qualifier .kind::i8 lists sm_100a, sm_101a (renamed to sm_110a), and sm_110a, but not sm_103a, the B300 target; the same instruction family's FP4 kinds were extended to sm_103a in PTX ISA 8.8. CUTLASS's kernel generator accordingly guards INT8 UMMA generation behind the absence of the family targets $\{100f, 101f, 103a\}$, generates FP8 unconditionally, and its only SM103-specific generator is FP4. vLLM compiles no INT8 GEMM for any Blackwell target and dispatches a null pointer for the integer case on SM100+, which emits the hard error "Int8 not supported on SM <N>" at the first forward pass; SGLang's ahead-of-time INT8 GEMM stops at SM90 and raises a generic not-implemented error above it. The authors document one non-default escape hatch in vLLM, an environment variable that reroutes INT8 to a JIT-compiled Triton backend, and leave its attained throughput unmeasured. Taking the four layers together, the paper concludes that a format present on the datasheet is, by default, undeployable on this hardware.

Load-bearing premise

The load-bearing premise is that NVIDIA's PTX ISA manual v9.3 is complete and current in its target list for tcgen05.mma .kind::i8; if the actual compiler or assembler already accepts sm_103a for that kind despite the manual's omission, the ISA-level withdrawal and the downstream skips built on it would not hold.

Editorial extensions

If this is right

A W8A8 INT8 checkpoint that runs on H200 or B200 will, by default, fail on B300 with a hard runtime error at the first forward pass, after the full model has been downloaded and loaded.
FP8 traverses all four layers natively on B300, so the practical default migration for INT8 users is re-quantization to FP8 or the non-default Triton reroute.
The obvious profiler criterion for "native INT8" on this GPU, grepping kernel names for tcgen05 or SM103 kernels, produces a false negative; the valid criterion is the IMMA tensor-op pipe counter combined with int8 operand dtypes.
vLLM's INT8-versus-FP8 comparison on B300 is by construction also a Triton-versus-native-backend comparison under matched per-tensor scaling recipes, so such measurements must be reported as path-specific bounds.
All repository findings are point-in-time; a single merged pull request can restore INT8 at any layer, so pinned commits and container digests are essential to reproducing the observations.

Reading between the lines

Editorial extensions of the paper, not claims the author makes directly.

A natural extension, not pursued in the paper, would be to run the same four-layer audit on later PTX ISA revisions or on the next architecture generation; if sm_103a is silently added to the .kind::i8 list later, the withdrawal would turn out to be a documentation lag rather than an architectural decision.
The maintenance artifacts used here, CMake gates, tuning-config file inventories, and dispatch tables, could serve as a general early-warning signal for format deprecation on other hardware, since they change before spec sheets or feature tables do.
The paper's suggestion that kernel coverage per (format, architecture) be tracked as a first-class planning property could be operationalized as a public coverage matrix maintained per serving stack, making this class of failure visible before hardware is purchased.
Because INT8 and FP8 are byte-identical in storage, the H200-to-B300 migration transfers memory-footprint planning unchanged while silently changing the compute path; one testable corollary is that capacity planners should treat format availability and compute-path availability as separate dimensions.

Share X Bluesky LinkedIn Reddit HN

Signed reviews

No signed human review yet.

Editorial analysis

A structured set of objections, weighed in public.

Desk editor's note, referee report, and a circularity audit.

Desk Editor's Note

read the letter

Colleague,

Read this one if you care about what actually runs on B300. The paper is a documentary audit: it traces INT8 W8A8 support from NVIDIA spec sheets through PTX ISA, CUTLASS, vLLM, and SGLang, and shows that while the datasheet lists INT8, no default execution path exists on Blackwell Ultra. That claim is well supported. The author pins commits, container digests, quotes source with line-level detail, and is explicit about scope: no performance numbers, no concealment narrative, no TensorRT-LLM claims. That is honest work, and the layered evidence is mutually consistent.

The strongest part is the software-layer audit. vLLM's dispatch passes a null pointer for INT8 on SM100, SGLang's AOT kernels stop at Hopper, and CUTLASS's generator guards INT8 UMMA behind absence of 103a targets while generating FP8 unconditionally. These are public code at pinned revisions and reproduce cleanly. The failure semantics (load-time checks pass, first forward pass dies) are a real service to practitioners.

The soft spots are real but not load-bearing. The ISA-level finding is an absence claim from PTX ISA v9.3: sm_103a is not in the .kind::i8 target list. The stress-test note is right that a single ptxas compile of tcgen05.mma with .kind::i8 for sm_103a would settle whether the manual lags the assembler. The CUTLASS and serving-engine findings are corroborating but not independent of the same documentation or engineering decisions, so the "mutually consistent withdrawal at four layers" framing should be softened to "three software layers plus a documentation absence." That said, the practical conclusion—INT8 W8A8 is undeployable by default on B300—survives even if the ISA manual is wrong, because the shipped software layers don't instantiate the path. The escape hatch was only functionally verified on sm_89, and the profiler false-negative trap is a prediction, not an empirical demonstration; both are minor and the author flags them.

The citation pattern is fine. The paper engages the relevant quantization and hardware literature and doesn't oversell its own novelty. It is a technical report, not a measurement study, and it reads like one.

Verdict: deserves a serious referee. I'd accept after a minor revision that either performs the ptxas compile or explicitly downgrades the ISA claim to "undocumented as of PTX 9.3" with the compile as future work. I'd bring it to our reading group.

Referee Report

1 major / 4 minor

Summary. This technical report audits INT8 W8A8 availability on NVIDIA Blackwell Ultra (B300) across four layers: published specifications, PTX ISA v9.3, CUTLASS's kernel generator, and the vLLM and SGLang serving engines. It reports that the spec sheet lists INT8 at roughly 1/30 of the FP8 dense rate, the PTX ISA does not list sm_103a for tcgen05.mma with .kind::i8, CUTLASS guards INT8 UMMA generation away from 103a builds, vLLM ships no Blackwell-capable INT8 GEMM and fails at the first forward pass after load, and SGLang's ahead-of-time INT8 GEMM kernels stop at SM90. The report also documents a Triton JIT reroute for vLLM, a profiler false-negative trap, and detailed reproducibility pins. It explicitly disclaims performance measurements and states what it does not claim.

Significance. If the central claims hold, this is a valuable and unusually well-documented snapshot of an important ecosystem fact: format availability is a property of the whole software stack, not just the silicon or the datasheet. The audit's strengths are concrete: every finding is traced to pinned commits, container digests, verbatim quoted strings, and explicitly dated sources; the paper is careful about scope and repeatedly states what it does not claim; the SGLang tuning-config asymmetry is a nice natural control; and the failure semantics and profiler guidance are actionable. The audit is self-contained and not circular: no quantity is fitted and no conclusion is derived from the target result. Its main risk is a single load-bearing absence claim at the ISA layer, which is documented rather than empirically tested.

major comments (1)

[Section III, quoted PTX ISA support statement] The central ISA-layer finding is an absence claim: tcgen05.mma with .kind::i8 is not listed for sm_103a in PTX ISA v9.3. The paper then draws a causal conclusion: 'there is no B300 instruction for [the software layers] to target except the legacy path.' That sentence is a claim about the toolchain and hardware, not merely about the manual's text. The corroborating evidence from CUTLASS and the serving engines is not fully independent, since those codebases may encode the same reading of the same documentation or the same engineering decision. Because this is the load-bearing pillar of the 'mutually consistent withdrawal at four layers' framing, the authors should either (a) run ptxas (or an equivalent assembler/driver test) for sm_103a with tcgen05.mma .kind::i8 and report the result, including the ptxas version, or (b) rewrite the causal claim and the abstract's 'never exposes' wording as 'does not document/expose in PTX ISA v9.3' and explicitly flag the possibility of undocumented toolchain support. The same qualification should then be propagated to the profiler guidance in Section VI(1).

minor comments (4)

[Figure 1] The caption explains the color palette and line styles, but the diamond/terminator symbol and the dashed reroute arrow are not listed as legend entries; consider adding a compact legend so the figure is self-contained.
[Section VI(1)] The sentence 'The obvious profiler criterion produces a false negative on sm_103' is worded as an observed fact, but it is a conditional prediction that follows only if the ISA-layer absence claim is correct. Please phrase it as 'if no tcgen05 integer kind exists on sm_103a, then the obvious profiler criterion will produce a false negative.'
[Section I(a) and Section VII] The 'to our knowledge' absence claim about third-party B300 INT8 measurements is well hedged, and the later caveat that an absence of evidence proves nothing is good; consider moving or repeating that caveat at the first occurrence so it is not separated from the claim.
[Appendix] The author-local items (marketplace snapshot and sm_89 functional run) are clearly labeled as such; for full re-executability, please include the exact ptxas version and the exact command used to extract the error string from the release-image .so when adding the requested compile test.

Circularity Check

0 steps flagged · score 0.0 of 10

No significant circularity; the audit is documentary and self-contained against externally pinned sources.

full rationale

The paper is an audit that makes no predictions and fits no parameters. Every load-bearing claim is traced to an external, pinned artifact: the PTX ISA v9.3 manual (Section III), CUTLASS source at commit dcf215a (Section IV), vLLM source and a shipped release binary (Section V-A), and SGLang source at commit b20c375 (Section V-B). The central ISA-layer finding is an absence claim quoted verbatim from the PTX manual ('Qualifier .kind::i8 is supported on following architectures: sm_100a sm_101a ... sm_110a'), with sm_103a absent. This is a documentation-dependence risk, not circularity: the claim is not defined in terms of the paper's own conclusion, and the corroborating CUTLASS and serving-engine evidence are independent public codebases, not restatements of the paper's inputs. The paper explicitly disclaims the causal over-reading by stating it does not claim INT8 is 'removed', reports no performance measurements, and notes that the absence of third-party B300 INT8 measurements 'proves nothing by itself'. No fitted input is renamed as a prediction, no uniqueness theorem is imported from the authors' prior work, and no known result is repackaged under new coordinates. The minimal self-references (a companion study 'in preparation', author-local verification of the Triton reroute on an sm_89 GPU) are explicitly labeled as such and are not load-bearing for the central documentary claims. Therefore no circular step exists, and the appropriate score is 0.

Assumptions & free parameters 0 free parameters · 4 assumptions · 0 invented entities

The audit introduces no free parameters or invented entities. The claim rests on domain assumptions about the accuracy and completeness of NVIDIA's documentation, the correctness of the B300/sm_103a mapping, and the fidelity of the pinned repository state. These are standard assumptions for a documentary audit and are disclosed.

assumptions (4)

domain assumption The PTX ISA manual v9.3 is complete and current in its list of architectures supporting tcgen05.mma .kind::i8.
The central ISA finding (no INT8 kind on sm_103a) is an absence claim based on this manual. If NVIDIA's documentation lags the compiler/assembler, the claim could be wrong. Entered in Section III via the verbatim quote.
domain assumption The mapping between B300 and compute capability 10.3 / sm_103a is correct.
The paper corroborates this via a marketplace API snapshot and CUTLASS's SM103-specific generator name, but NVIDIA's own quoted documents do not explicitly state this mapping in the cited passages. All target-specific claims (ISA, CUTLASS, vLLM, SGLang) depend on it. See Section III footnote and Appendix.
domain assumption The pinned commits and container digests accurately reflect the state of the audited repositories at access time.
Reproducibility rests on the pins in Appendix A. The paper itself notes that vLLM's :latest tag drifted within hours, so the pins are load-bearing for the audit's validity.
domain assumption The quoted source code snippets are genuine and correctly interpreted.
The CUTLASS generator guard and vLLM dispatch null pointer are direct quotes from public repositories at pinned commits; mis-transcription or misinterpretation would undermine the layer-specific findings.

how reviews work

0 comments

Cite this review

Pith. "Pith review of Spec Sheets Are Not Kernels: An ISA- and Source-Level Audit of INT8 Availability on NVIDIA Blackwell Ultra." pith.science (2026). https://pith.science/paper/P7SLECKW

@misc{pith2026260811693,
  author       = {Pith},
  title        = {Pith review of: Spec Sheets Are Not Kernels: An ISA- and Source-Level Audit of INT8 Availability on NVIDIA Blackwell Ultra},
  year         = {2026},
  howpublished = {\url{https://pith.science/paper/P7SLECKW}},
  note         = {Machine review of arXiv:2608.11693}
}

read the original abstract

NVIDIA's published specifications give the Blackwell Ultra GPU (B300) a dense-compute ratio of roughly 30:1 between FP8 and INT8 tensor-core throughput; its predecessors, H200 and B200, both provide 1:1. We audit what this deprioritization means in practice by tracing INT8 W8A8 support through four layers of the stack: the published specifications, the PTX ISA, NVIDIA's CUTLASS kernel library, and the two major open-source LLM serving engines (vLLM and SGLang). We find a consistent, layered withdrawal: (i) the PTX ISA never exposes the fifth-generation tensor-core integer path (tcgen05.mma with .kind::i8) on sm_103a, even though the same PTX revision extends the FP4 kinds to that target, leaving legacy warp-level IMMA as the only architecturally legal integer tensor-core path on B300; (ii) CUTLASS's kernel generator explicitly skips INT8 UMMA generation for any build targeting 103a, while generating FP8 unconditionally; (iii) vLLM ships no INT8 GEMM for Blackwell and fails with a hard runtime error at the first forward pass, after the model has loaded; and (iv) SGLang's ahead-of-time INT8 GEMM stops at Sm90, while its FP8 tuning configurations already cover B200. We document an escape hatch (rerouting vLLM's INT8 path to a JIT-compiled Triton backend via an environment variable), a false-negative trap in the obvious profiler methodology for detecting "native INT8" on sm_103, and the practical failure semantics that make naive testing expensive. Together, these findings show that a quantization format's availability is a property of the whole stack rather than of the model or the spec sheet. Four distinct layers, three of them NVIDIA's own, withdrew INT8 support in mutually consistent ways, and a format that is nominally present on the datasheet is, by default, undeployable on this hardware.

Figures

Figures reproduced from arXiv: 2608.11693 by the authors.

**Figure 1.** The path of a W8A8 checkpoint through the B300 stack. FP8 traverses all four layers natively. INT8 is rate-limited in silicon (Section II); the terminator [PITH_FULL_IMAGE:figures/full_fig_p002_1.png] view at source ↗

**Figure 2.** Published dense per-GPU tensor-core peaks (HGX SKU bin; values as in Table I). INT8 markers are open squares and FP8 markers are filled circles, [PITH_FULL_IMAGE:figures/full_fig_p003_2.png] view at source ↗

Discussion (0). Continue with ORCID to comment.

Pith tools

Reviewed August 16, 2026 · model on record in the stance chip above.