{"id":"9c6a0107-f956-4ea3-b800-579f293c36f7","arxiv_id":"2608.11693","paper_version":1,"verdict":"ACCEPT","confidence":"HIGH","novelty_score":7.0,"correctness_risk":"low","formal_verification":"none","parameter_count":0,"one_line_summary":"On NVIDIA Blackwell Ultra, INT8 W8A8 is undeployable by default because the PTX ISA, CUTLASS, vLLM, and SGLang all lack a fifth-generation INT8 tensor-core path, despite the datasheet listing INT8 support.","lead":"An audit of NVIDIA's Blackwell Ultra GPU finds that INT8 tensor-core compute is listed on the datasheet but does not work by default in the PTX ISA, CUTLASS, vLLM, or SGLang. If you serve an INT8-quantized LLM on a B300, the model loads and then crashes at the first forward pass.","discovery_kind":"new_application","skeptic_critique":{"model":"deepseek-v4-flash","headline":"The ISA-layer finding is an untested documentation-absence claim; a single ptxas compile of tcgen05.mma .kind::i8 for sm_103a would settle whether the manual is accurate or the 'no instruction to target' explanation is overclaimed.","rationale":"The reader's weakest_assumption is the completeness of PTX ISA v9.3 for sm_103a, and my stress-test agrees that this is the most load-bearing point. The reason it is load-bearing is not that the paper's practical conclusion is fragile: even if ptxas accepted the instruction, vLLM and SGLang would still ship no Blackwell INT8 GEMM and the default deployment would still fail. The load-bearing issue is the paper's stronger, explanatory claim that the ISA never exposed the path and therefore there was no instruction for the software layers to target. That claim converts a documentation silence into an architectural fact, and the paper never tests the assembler that defines what users can actually compile. The CUTLASS guard and vLLM null-dispatch are real and correctly documented, but they are not independent proof of the ISA-level absence because those codebases likely share the same assumption or the same upstream toolchain; a user-level ptxas test is the missing independent arbiter. I recommend CONDITIONAL rather than REJECT or UNCHANGED: the software-stack findings are strong and should stand, but the ISA-layer pillar and the title-level claim should be accepted only after the ptxas check is run. If the check rejects the instruction, the paper's current argument is confirmed and the verdict should revert to ACCEPT; if it accepts, Section III and the abstract's 'never exposes' wording need revision, and the paper's scientific contribution becomes a documentation-gap audit rather than an architectural finding.","tokens_in":13342,"tokens_out":12126,"duration_ms":122906,"concrete_test":"Use the CUDA toolkit corresponding to PTX ISA 9.3 (CUDA 13.x), write a minimal PTX kernel invoking tcgen05.mma with .kind::i8 using an accumulator form from the PTX manual, set .target sm_103a, and run ptxas -v. A control kernel targeting sm_100a with the same instruction should be assembled in the same harness to validate the test. If ptxas rejects the sm_103a version with an unsupported-kind or unsupported-target error, the manual's absence is confirmed at the assembler level. If it assembles, disassemble with cuobjdump and check for tcgen05/UMMA SASS; success falsifies the ISA-layer claim and requires softening the 'no B300 instruction to target' explanation. Report the exact ptxas version and error/warning output.","verdict_should_be":"CONDITIONAL","load_bearing_attack":"The paper's headline claim rests on the Section III assertion that tcgen05.mma .kind::i8 is never exposed on sm_103a. That assertion is a negative claim derived solely from PTX ISA v9.3 text. The paper then draws a causal inference: 'there is no B300 instruction for [the software layers] to target except the legacy path.' The corroborating evidence from CUTLASS (generator guard) and vLLM (null INT8 dispatch) is not independent: both are NVIDIA-adjacent codebases that plausibly encode the same reading of the same PTX documentation or the same engineering decision, so they do not test what a user could assemble with ptxas against the actual hardware. If the PTX manual lags the assembler, or if a newer PTX revision adds .kind::i8 for sm_103a, the ISA pillar of the audit is wrong: the fifth-generation integer path exists but is undocumented and unused. The software-stack conclusion ('undeployable by default') would survive, but the 'mutually consistent withdrawal at four layers' framing and the 'explains, rather than merely describes' causal claim would need to be walked back to a documentation-gap finding.","agreement_with_reader":"agree"},"referee_report":{"model":"deepseek-v4-flash","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.","tokens_in":13515,"tokens_out":7850,"duration_ms":83269,"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":[{"comment":"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).","section":"Section III, quoted PTX ISA support statement"}],"minor_comments":[{"comment":"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":"Figure 1"},{"comment":"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":"Section VI(1)"},{"comment":"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.","section":"Section I(a) and Section VII"},{"comment":"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.","section":"Appendix"}],"recommendation":"major_revision","confidential_remarks":"The manuscript is well within the scope of cs.AR and is likely to be a frequently cited snapshot. The single ISA-absence issue is easy to resolve: either run a ptxas compile check for sm_103a or scope the causal claim to documented ISA exposure. With that change, I would support acceptance. I see no circularity or fitting-related concern; the audit is self-contained and properly pinned."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"Colleague,\n\nRead 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.\n\nThe 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.\n\nThe 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.\n\nThe 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.\n\nVerdict: 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.","headline":"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.","tokens_in":14099,"tokens_out":1943,"would_cite":true,"duration_ms":19282,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"On Blackwell Ultra, INT8 appears on the datasheet but is withdrawn from the ISA, kernel library, and serving engines.","keywords":["INT8 quantization","W8A8","Blackwell Ultra","B300","PTX ISA","tensor cores","LLM serving","kernel availability"],"falsifier":"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.","tokens_in":13088,"feed_emoji":"📉","tokens_out":10468,"duration_ms":90273,"temperature":0.7,"pith_summary":"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.","feed_headline":"INT8 is on the spec sheet, gone from the stack","feed_subtitle":"Four layers (PTX ISA, CUTLASS, vLLM, SGLang) all removed INT8; only a Triton reroute remains.","key_machinery":"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.","core_discovery":"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.","pith_inferences":["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."],"forward_implications":["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."],"supporting_citations":[{"why":"PTX ISA v9.3: its .kind::i8 target list omits sm_103a, establishing the ISA-level absence that the entire audit hangs on; it also supplies the FP4-extension contrast and the repeated 'except .kind::i8' carve-outs.","marker":"[20]"},{"why":"CUTLASS at commit dcf215a: generator.py's guard around GenerateSM100_TensorOp_int8_UMMA_gemm shows INT8 UMMA being skipped for 103a builds while FP8 is generated unconditionally.","marker":"[22]"},{"why":"vLLM at commit 6c95a641: CMakeLists.txt and scaled_mm_c3x_sm100.cu show no Blackwell INT8 GEMM compiled and a null pointer passed for the integer case, producing the hard runtime error.","marker":"[23]"},{"why":"SGLang at commit b20c375: int8_gemm_kernel.cu and tuning-config inventories show INT8 stopping at SM90 while FP8 tuning covers B200.","marker":"[25]"},{"why":"Blackwell Ultra datasheet: lists the INT8 tensor-core peak, the nominal presence that the audit contrasts with actual unavailability.","marker":"[11]"},{"why":"Blackwell Architecture Technical Brief Table 3: gives FP8 4.5 PFLOPS and INT8 0.15 POPS dense for B300, and the 1:1 B200 numbers that locate the discontinuity at B300.","marker":"[12]"},{"why":"CUDA Programming Guide Table 33: lists INT8 as supported at compute capability 10.3, supporting the paper's contention that INT8 is rate-limited rather than removed from the silicon.","marker":"[18]"}],"fun_headline_variants":["Blackwell Ultra: INT8 advertised, not implemented","The INT8 mirage on B300: spec'd, not shipped","B300's INT8: only on the datasheet","No INT8 path on Blackwell Ultra despite datasheet claims","Four layers, zero INT8: the B300 integer gap"],"cache_read_input_tokens":3200,"weakest_assumption_plain":"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.","fun_headline_variants_meta":{"raw":{"variants":["Blackwell Ultra: INT8 advertised, not implemented","The INT8 mirage on B300: spec'd, not shipped","B300's INT8: only on the datasheet","No INT8 path on Blackwell Ultra despite datasheet claims","Four layers, zero INT8: the B300 integer gap"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000878,"raw_usage":{"total_tokens":3954,"prompt_tokens":1258,"completion_tokens":2696,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":874,"completion_tokens_details":{"reasoning_tokens":2612}},"tokens_in":874,"tokens_out":2696,"duration_ms":19809,"temperature":1.0,"reasoning_tokens":2612,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-16T00:30:08.914203+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"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.","supporting_citations":[{"cited_title":"Parallel thread execution ISA version 9.3,","cited_arxiv_id":null,"evidence_quote":"PTX ISA v9.3: its .kind::i8 target list omits sm_103a, establishing the ISA-level absence that the entire audit hangs on; it also supplies the FP4-extension contrast and the repeated 'except .kind::i8' carve-outs."},{"cited_title":"CUTLASS: CUDA templates for linear algebra subroutines,","cited_arxiv_id":null,"evidence_quote":"CUTLASS at commit dcf215a: generator.py's guard around GenerateSM100_TensorOp_int8_UMMA_gemm shows INT8 UMMA being skipped for 103a builds while FP8 is generated unconditionally."},{"cited_title":"vLLM: A high-throughput and memory-efficient in- ference engine for LLMs,","cited_arxiv_id":null,"evidence_quote":"vLLM at commit 6c95a641: CMakeLists.txt and scaled_mm_c3x_sm100.cu show no Blackwell INT8 GEMM compiled and a null pointer passed for the integer case, producing the hard runtime error."},{"cited_title":"SGLang: A fast serving framework for large language models,","cited_arxiv_id":null,"evidence_quote":"SGLang at commit b20c375: int8_gemm_kernel.cu and tuning-config inventories show INT8 stopping at SM90 while FP8 tuning covers B200."},{"cited_title":"NVIDIA Blackwell Ultra datasheet,","cited_arxiv_id":null,"evidence_quote":"Blackwell Ultra datasheet: lists the INT8 tensor-core peak, the nominal presence that the audit contrasts with actual unavailability."},{"cited_title":"NVIDIA Blackwell architecture technical brief,","cited_arxiv_id":null,"evidence_quote":"Blackwell Architecture Technical Brief Table 3: gives FP8 4.5 PFLOPS and INT8 0.15 POPS dense for B300, and the 1:1 B200 numbers that locate the discontinuity at B300."},{"cited_title":"CUDA C++ programming guide: Compute capabilities,","cited_arxiv_id":null,"evidence_quote":"CUDA Programming Guide Table 33: lists INT8 as supported at compute capability 10.3, supporting the paper's contention that INT8 is rate-limited rather than removed from the silicon."}],"review_version":1}