{"id":"0c91e114-ddf6-408d-93da-86c39b7500b4","arxiv_id":"2507.10789","paper_version":2,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":6.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":0,"one_line_summary":"A microbenchmark study maps memory hierarchy, execution pipelines, and FP4/FP6 tensor-core behavior on Nvidia's Blackwell RTX 5080 and compares it with Hopper's H100.","lead":"The paper uses custom microbenchmarks to measure the internals of Nvidia's new Blackwell GPU (RTX 5080) and compares them with the older Hopper GPU (H100). It reports cache behavior, arithmetic and tensor-core speeds, and power use to help software developers tune code for Blackwell.","discovery_kind":"new_application","skeptic_critique":{"model":"deepseek-v4-flash","headline":"FP4/FP6 tensor-core conclusions rest on a QMMA fallback path, not the native tcgen05/OMMA path the paper claims to dissect.","rationale":"The reader's SKU-confounding worry is legitimate and I share it: any cross-GPU comparison between an RTX 5080 and an H100 PCIe conflates clock, SM count, memory technology, TDP, and software maturity with architecture. If the paper were only a comparison of two products, that would be acceptable, but conclusions such as 'Blackwell's warp scheduler is optimized for low-precision, high-ILP workloads' generalize beyond GB203. My more specific concern is in the tensor-core section, where the paper itself limits its own evidence: Table IV says tcgen05 is not supported for sm_120a, and Section V.A reports that FP4 mma.sync unexpectedly emits QMMA rather than the documented OMMA instruction, calling QMMA the fallback. This means the FP4/FP6 latency, throughput, and power numbers do not necessarily measure the native 5th-gen tensor-core datapath; they measure whatever path the current toolchain selected for an old-style mma.sync instruction. The abstract's '5th generation tensor cores supporting FP4 and FP6' is the paper's distinctive contribution, so this issue is load-bearing. A disassembly and native-path rerun would settle it. I would not reject the whole paper: the unified INT32/FP32, FP64, and memory-hierarchy sections are separate and can stand if reproducible. I also note independent arithmetic slips, such as Table VII using runtime in milliseconds without converting in Equation 2, which makes the reported TFLOP/s entries 1000x low; both GPUs are affected equally, so the relative GEMM finding survives, but the paper's claim of precise quantitative characterization is weakened. The tensor-core conclusions should be conditioned on demonstrating the native path or explicitly re-scoping the claims to QMMA compatibility mode.","tokens_in":14353,"tokens_out":10559,"duration_ms":126671,"concrete_test":"Compile the Section V FP4/FP6 mma.sync kernels on the RTX 5080 with the newest CUDA/PTX toolchain that supports sm_120a, disassemble the resulting cubin with nvdisasm, and confirm which SASS instructions (QMMA, OMMA, or tcgen05) actually execute. Then run the same matrix shapes through a native OMMA/tcgen05 path, for example cuBLASLt FP4 or a tcgen05 PTX kernel if available on sm_120a, and compare throughput, latency, and power against the paper's QMMA-based numbers. If no OMMA/tcgen05 instruction appears, or if the native path's results differ materially, the paper's tensor-core characterization must be re-scoped or redone.","verdict_should_be":"CONDITIONAL","load_bearing_attack":"Section V, the paper's headline contribution on 5th-generation tensor cores, is undercut by the paper's own limitations. Table IV states tcgen05 is not supported for sm_120a, and Section V.A reports that the FP4 mma.sync kernel produces QMMA SASS instructions rather than the OMMA instruction that NVIDIA's CUDA Binary Utilities document for FP4; the authors call QMMA 'the fall back for FP4 inputs in the current software.' If that is correct, the FP4/FP6 latency, throughput, and power measurements in Figures 4-5 and Table VI describe a compatibility/fallback execution path, not the native tcgen05/OMMA tensor-core datapath that distinguishes Blackwell. The conclusion that 'Blackwell's warp scheduler is optimized for low-precision, high-ILP workloads' therefore conflates current software behavior with microarchitecture. In addition, the claimed 1.21-cycle completion latency for all mma formats is inconsistent with a dependent-instruction latency metric and suggests a throughput ratio is being mislabeled.","agreement_with_reader":"partial"},"referee_report":{"model":"deepseek-v4-flash","summary":"This paper presents a microbenchmark-based analysis of NVIDIA's Blackwell architecture, using the GeForce RTX 5080 (GB203) as the Blackwell representative and the H100 PCIe (GH100) as the Hopper baseline. The authors measure compute-pipeline latencies and throughputs for INT32/FP32/FP64, warp-scheduler behavior, the memory hierarchy (shared memory, L1, L2, global), and low-precision tensor-core operations (FP4, FP6, FP8), including power consumption. They also include two case studies: a dense FP8 GEMM and a Transformer inference workload. The central claim is that these microbenchmarks reveal subtle architectural differences between Blackwell and Hopper, providing guidance for developers and compiler writers.","tokens_in":14518,"tokens_out":7172,"duration_ms":75254,"significance":"If the measurements are correct, the paper is a useful addition to the empirical GPU microarchitecture literature, especially for the newly released Blackwell architecture. It reports concrete, falsifiable values (e.g., 4-cycle INT32/FP32 true latency, L1 hit latency of 30-40 cycles, L2 latency of 273/358 cycles on H100/RTX 5080) and cross-checks PTX instructions against generated SASS, which strengthens the methodology. The memory-hierarchy characterization and the explicit comparison of two product generations are valuable. However, the tensor-core section, which is a headline contribution, is undermined by the use of a software fallback path for FP4/FP6 and by confusion between latency and throughput metrics. The D-GEMM case study also contains a factor-of-1000 arithmetic error. These issues must be resolved before the paper's conclusions can be accepted.","major_comments":[{"comment":"The FP4/FP6 tensor-core measurements characterize a compatibility/fallback path, not the native tcgen05/OMMA datapath. The paper itself states in Section V.A that the FP4 mma instruction produces QMMA SASS rather than the documented OMMA instruction, and calls QMMA \"the fall back for FP4 inputs in the current software.\" Table IV also notes that tcgen05 is not supported for sm_120a. Consequently, the latency, throughput, and power values in Figures 4-5 and Table VI describe the behavior of the QMMA fallback path under the current CUDA toolchain, not the underlying 5th-generation tensor-core hardware. The conclusion in Section V.D that \"Blackwell's warp scheduler is optimized for low-precision, high-ILP workloads\" is therefore not supported by these data. The authors should either obtain measurements through the native path (e.g., by using the OMMA instructions they observed for FP8 with block scaling, or by using a toolchain that supports tcgen05 for sm_120a) or explicitly reframe the section as a study of current software behavior on Blackwell, with conclusions restricted accordingly.","section":"V.D"},{"comment":"The claimed \"completion latency\" of 1.21094 cycles for all mma formats at ILP=1 and one warp is implausible as a dependent-instruction latency for a matrix multiply accumulate. Per the paper's own definition in Section IV.A, completion latency is measured on a set of independent instructions and reported as cycles per instruction, which is a throughput-oriented quantity. For a single warp issuing a dependent chain of mma instructions, the latency from issue to data-ready is typically tens of cycles on NVIDIA GPUs. The reported value suggests that the measurement captures the issue interval (reciprocal throughput) rather than the completion latency of a dependent chain. Because Figure 5 is labeled \"Latency\" and the text in Section V.D uses this number to compare GB203 and GH100, the paper conflates throughput with latency. The authors should either measure true dependent-instruction latency (e.g., by chaining mma operations through the accumulator) or relabel and reinterpret the results as throughput measurements throughout the tensor-core section.","section":"V.D"},{"comment":"The D-GEMM throughput values in Table VII are wrong by a factor of 1000. Using Equation (2) with the stated runtime of 4.710 ms for the RTX 5080 at M=N=K=8192, the throughput is 2*8192^3 / 0.00471 s = 2.33e14 FLOP/s = 233 TFLOPS, not 0.233 TFLOPS. The same error appears for all entries in the table (e.g., Hopper at 8192^3 should be roughly 887 TFLOPS, not 0.887 TFLOPS). The relative ordering is unchanged, but the absolute values are incorrect and mislead the reader about the actual performance levels. The authors must correct the unit conversion (milliseconds to seconds) and recalculate all values in Table VII and the corresponding text.","section":"VII.A"},{"comment":"The cross-architecture comparison treats one consumer GPU (RTX 5080, GB203) as representative of Blackwell and one datacenter GPU (H100 PCIe, GH100) as representative of Hopper. These two SKUs differ not only in architecture but also in SM count, memory type (GDDR7 vs HBM2e), L2 size and partitioning, clock rates, TDP, and target market, as shown in Tables I and II. Several conclusions attribute observed differences to architectural choices rather than SKU configuration. For example, Section VI.C states that GH100's partitioned L2 is \"optimized for high concurrency\" while GB203's unified L2 \"simplifies hardware complexity,\" and Section VI.D attributes the global-memory latency difference to HBM2e vs GDDR7, a memory-technology distinction rather than a Blackwell-vs-Hopper architectural one. To support these attributions, the authors should either restrict their claims to the specific products tested or explicitly discuss the confounds and justify why particular differences are architectural. Without such a caveat, the paper overgeneralizes from a two-point comparison.","section":"Introduction, Tables I-II, Sections VI.C-D"},{"comment":"The paper states in Section I that \"We are unable to share the code at this time due to the blind-review policy, but we plan to open-source it post the review process and the outcome.\" For an empirical microbenchmark paper whose entire contribution rests on measurements, this is a significant barrier to verification. Even during double-blind review, code can be provided as supplementary material without revealing author identity. The authors should either make the microbenchmarks available for the review process or provide a detailed experimental appendix covering kernel launch configurations, clock-frequency settings, power-measurement methodology (how nvidia-smi readings were aggregated), and the exact PTX/SASS snippets that underlie each key result. Without this, the load-bearing measurements (e.g., tensor-core latencies, memory latencies) cannot be independently checked.","section":"Global"}],"minor_comments":[{"comment":"There are numerous typographical errors, including \"Kelper\" (Section II), \"Turning\" for Turing (Section II), \"fmaand\" (Section IV.B), \"there respective architectures\" (Section V.D), \"stepper power curve\" (Section VII.A), and \"nv f p8 e4m3\" (Section VII.A). A thorough proofread is needed.","section":"Throughout"},{"comment":"The caption says the 8192^3 Blackwell runtime (4.710 ms) was omitted from the graph, but the graph is supposed to show runtime across all sizes. Please explain why the point was omitted and consider including it in an inset or on a logarithmic axis so that the reader can see the full trend.","section":"Figure 11 caption"},{"comment":"The note \"TCGEN05 IS YET TO BE SUPPORTED FOR THE ARCH. SM120A\" appears as a detached all-caps line under the table. Integrate this into the surrounding text and elaborate on what it means for the experiments (e.g., which instructions were available and which were not).","section":"Table IV"},{"comment":"The header \"POWER USAGE (WATTS)/PERFORMANCE PER WATT\" is ambiguous. The body text describes the values as power consumption, so the header should be simply \"Power usage (watts)\". If performance-per-watt is also reported, separate columns are needed.","section":"Table VI"},{"comment":"\"On GB203, without an instruction between the registers, the subtracted value is 1, compared to a value of 2 forGH100\" is unclear. Clarify whether this is the overhead of the %clock64 measurement itself and how it was accounted for in later latency calculations.","section":"Section IV.A"},{"comment":"The sentence \"The maximum ILP level at which sustained throughput is achieved for each precision format, across decreasing warp counts is ILP=5 with 29 active warps, for GH100, and ILP=6 at 25 active warps, for GB203\" is difficult to parse. Please rephrase to state the experimental procedure and the exact meaning of \"maximum ILP level at which sustained throughput is achieved.\"","section":"Section V.D"},{"comment":"The text says \"with a from' 58.8W to 45W in FP8\" in Section VII.B (Transformer inference); this appears to be a typo for \"with a decrease from 58.8W to 45W.\"","section":"Section VII.B"}],"recommendation":"major_revision","confidential_remarks":"This paper is within the scope of the journal and addresses a timely topic. The core memory-hierarchy and compute-pipeline measurements appear plausible and are a useful contribution. However, the tensor-core section, which is presented as a central contribution, currently measures a software fallback path and conflates latency with throughput. The D-GEMM table also contains a clear factor-of-1000 error. These are fixable with a major revision, provided the authors either obtain native-path tensor-core data or substantially weaken their architectural claims. The comparison between two very different SKUs also needs a careful framing to avoid overgeneralization."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"Here's the short version: this is the first detailed public microbenchmark pass at Blackwell's SM subunits and memory hierarchy, and it contains at least one genuinely new observation—FP4 mma.sync currently lowers to QMMA rather than OMMA on sm_120a. But the paper's tensor-core section is shakier than its headline, and the authors' own text provides the evidence. They report that tcgen05 is not yet supported for sm_120a and that FP4 uses the QMMA fallback; that means the FP4/FP6 latency, throughput, and power numbers in Figures 4-5 and Table VI describe the compatibility path, not the native 5th-generation tensor core path they claim to dissect. The claim that 'Blackwell's warp scheduler is optimized for low-precision, high-ILP workloads' is thus a software-behavior claim, not a microarchitecture conclusion. The 1.21-cycle 'completion latency' for all mma formats also looks like a throughput ratio mislabeled as dependent latency.\n\nWhat's good: the memory hierarchy measurements (L1 ~30-40 cycles, L2 273 vs 358 cycles, shared memory and L2 saturation) are plausible and cross-checked against SASS. The unified INT32/FP32 mixed-workload latency results are interesting. The SASS-level QMMA observation is exactly the kind of concrete data point the field needs.\n\nSoft spots, in order: the tensor-core framing; the confirmed arithmetic error in Table VII (8192^3 GEMM should be ~233 TFLOPS, not 0.233; they forgot the milli- or used ms as seconds); power numbers in Table VI look inconsistent with Section VII's GEMM power (16 W vs 114 W peak); code is not available, which matters for reproducibility; and comparing a consumer GB203 to GH100 confounds architecture with SKU and memory type. These are fixable but real.\n\nOverall: it's an honest paper that flags its own limitations (the fallback is in the text), but it overreaches in conclusions and needs major revision. I'd send it to a serious referee—this is the first Blackwell microbenchmark dataset and it warrants scrutiny—but I would not cite it yet, and I'd push hard for code/data release.","headline":"Useful first Blackwell microbenchmark dataset, but the FP4/FP6 tensor-core story is a software fallback path, not native hardware.","tokens_in":15073,"tokens_out":2983,"would_cite":false,"duration_ms":31936,"reading_group":"maybe","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"This paper uses custom PTX and CUDA microbenchmarks to dissect the NVIDIA Blackwell architecture (RTX 5080/GB203) and compare it against Hopper (H100/GH100), quantifying latency, throughput, cache behavior, tensor-core precision support…","keywords":["Blackwell","Hopper","GPU microbenchmarking","tensor cores","FP4","FP6","memory hierarchy","power efficiency"],"falsifier":"Run the same benchmark suite on another Blackwell GPU, such as an RTX 5090 or a datacenter GB100, and on an H100 SXM rather than PCIe; if the FP4 SASS instruction, the FP64 two-unit latency, the L2 partition behavior, or the unified-core mixed-workload latency differs across Blackwell parts, the paper's architecture-level conclusions are SKU-specific. Also re-test the FP4 mma mapping with a newer CUDA version: if FP4 uses OMMA rather than QMMA, the reported 'fallback' is a transient software state, not an architectural fact.","tokens_in":14142,"feed_emoji":"⚡","tokens_out":5212,"duration_ms":54663,"temperature":0.7,"pith_summary":"The paper tries to establish a detailed, measurement-based picture of NVIDIA's Blackwell GPU architecture by probing it with purpose-built PTX/CUDA microbenchmarks and comparing results with the previous Hopper generation. It claims that Blackwell's unified INT32/FP32 execution units handle mixed integer/floating-point workloads with lower latency than Hopper's separate pipelines, while its FP64 units are so weak (two per SM) that FP64 is effectively supported but meant to be emulated via other precisions. On tensor cores, it reports that FP4, FP6, and FP8 mma instructions map to SASS instructions in a particular way (QMMA for FP8/FP6 and also for FP4 as a fallback; OMMA appears with block scaling), and that Blackwell achieves higher throughput and lower power than Hopper for low-precision formats. The paper also maps the memory hierarchy's latency boundaries, L1/L2 capacities, and global memory bandwidth, and finds that in a real cuBLASLt FP8 D-GEMM, Hopper still outperforms Blackwell on this consumer SKU. The intended value is actionable tuning guidance for developers targeting Blackwell, filling the gap of published microarchitectural data.","feed_headline":"RTX 5080's FP4/FP6 tensor path mapped by microbenchmarks","feed_subtitle":"A PTX-level study compares Blackwell's caches, schedulers, and FP4/FP6 tensor throughput against Hopper, with power data.","key_machinery":"The carrying mechanism is a suite of PTX-level microbenchmarks written as separate kernel files to block compile-time optimizations, with SASS inspection to verify instruction mapping. Clock cycles are captured with the %clock64 special register; latency is measured as true serialized dependent-instruction latency and completion latency with independent instructions; throughput is measured as instructions per cycle per SM. Memory behavior is probed with a pointer-chase benchmark that reveals cache boundaries, plus warp- and stride-swept shared-memory, L1, and L2 access tests. This setup lets the authors attribute observed behavior to specific SM sub-cores, scheduler behavior, and cache partitions.","core_discovery":"The central claim is that the RTX 5080 (GB203) reveals a consistent set of architectural choices in Blackwell: unified INT32/FP32 cores that reduce mixed-workload latency but can create hazards; only two FP64 units per SM, making FP64 execution a compatibility feature rather than a throughput path; fifth-generation tensor cores whose low-precision mma instructions are currently implemented in SASS as QMMA for FP8 and FP6 (with FP4 falling back to QMMA as well in the tested software), delivering sustained throughput above 11 TFLOP/s at ILP=6; a memory hierarchy with 128 KB L1/SM, about 99 KB configurable shared memory, a monolithic 65 MB L2, and GDDR7 global memory with roughly 876.7-cycle latency; and a warp scheduler that prefers higher per-thread instruction-level parallelism over Hopper's bulk concurrency. The paper asserts these measurements are accurate characterizations of the architecture under stress and form a basis for portable optimization guidance.","pith_inferences":["If Blackwell's SASS mapping for FP4 is indeed a software fallback, a future CUDA release may switch FP4 to OMMA; the paper's power and throughput numbers for FP4 would then need re-measurement and could improve.","The conclusion that GB203's scheduler favors high ILP suggests a testable porting rule: Blackwell kernels should be restructured to expose more independent mma operations per thread rather than relying on large active warp counts, and similar gains may appear on other Blackwell SKUs if the scheduler design is shared.","The GB203/H100 comparison conflates architecture with memory technology (GDDR7 vs HBM2e) and SKU class; extending the same benchmarks to a datacenter Blackwell part (e.g., GB100) would separate what is Blackwell from what is consumer-GPU packaging.","For compiler writers, the observed 1.21-cycle completion latency for all low-precision mma formats on GB203 suggests a single unified low-precision pipeline; instruction schedulers could treat FP4/FP6/FP8 mma identically for issue purposes."],"forward_implications":["On GB203, mixed INT32/FP32 workloads should see lower latency than on GH100, so developers can interleave integer and floating-point instructions without the penalty seen on older split-pipeline architectures.","FP64 on Blackwell consumer GPUs is not a viable high-throughput path; scientific codes that need double precision should expect to emulate FP64 through FP32 or tensor-core paths, or run on Hopper-class hardware.","Low-precision tensor-core users on Blackwell should be aware that FP4/FP6 inputs currently execute through QMMA (with FP4 as fallback) and that throughput is maximized with high ILP (around ILP=6) rather than large warp counts.","Memory-bound kernels on GB203 should avoid strided access patterns that trigger bank conflicts in the smaller 128 KB L1/shared partition, while Hopper tolerates higher warp pressure in shared memory.","For FP8 D-GEMM through cuBLASLt on the RTX 5080, current software yields lower throughput than Hopper, so developers should validate kernel selection or expect practical FP8 gains to appear only after software matures."],"supporting_citations":[{"why":"Supplies the microbenchmarking methodology for dissecting SM execution units that this paper adapts to Blackwell.","marker":"[7]"},{"why":"Extends the dissection approach to Turing tensor cores, informing the tensor-core benchmark design here.","marker":"[8]"},{"why":"Provides the recent Hopper microarchitectural measurements against which the H100 results are framed.","marker":"[9]"},{"why":"Establishes the pointer-chase cache-boundary approach used to map the memory hierarchy.","marker":"[5]"},{"why":"Defines %clock64 and the PTX instruction syntax used for all timing and mma kernels.","marker":"[26]"},{"why":"Documents the QMMA and OMMA SASS instructions used to interpret tensor-core instruction mapping.","marker":"[28]"},{"why":"Provides the basis for excluding first-run warm-up effects from reported latencies.","marker":"[27]"},{"why":"Provides the TensorRT inference engine used for the transformer power-efficiency case study.","marker":"[29]"},{"why":"Supplies the GPT-NeoX model exercised in the transformer inference benchmarks.","marker":"[30]"}],"fun_headline_variants":["Blackwell microbenchmarks reveal FP4/FP6 tensor core design","RTX 5080 vs H100: Blackwell's hidden architectural trade-offs","Inside Blackwell: L1, L2, and scheduler secrets from microbenchmarks","Microbenchmarks dissect NVIDIA's Blackwell GPU architecture","Blackwell's FP4/FP6 tensor paths exposed by PTX benchmarks"],"cache_read_input_tokens":3200,"weakest_assumption_plain":"The conclusions treat a single consumer GPU, the RTX 5080 (GB203), as representative of the entire Blackwell architecture and a single H100 PCIe (GH100) as representative of Hopper, so any observed difference could also come from memory technology, clock speed, TDP, or SKU configuration rather than from the architecture itself.","fun_headline_variants_meta":{"raw":{"variants":["Blackwell microbenchmarks reveal FP4/FP6 tensor core design","RTX 5080 vs H100: Blackwell's hidden architectural trade-offs","Inside Blackwell: L1, L2, and scheduler secrets from microbenchmarks","Microbenchmarks dissect NVIDIA's Blackwell GPU architecture","Blackwell's FP4/FP6 tensor paths exposed by PTX benchmarks"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.00043,"raw_usage":{"total_tokens":2212,"prompt_tokens":973,"completion_tokens":1239,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":589,"completion_tokens_details":{"reasoning_tokens":1143}},"tokens_in":589,"tokens_out":1239,"duration_ms":10310,"temperature":1.0,"reasoning_tokens":1143,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-06T17:25:53.262906+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Run the same benchmark suite on another Blackwell GPU, such as an RTX 5090 or a datacenter GB100, and on an H100 SXM rather than PCIe; if the FP4 SASS instruction, the FP64 two-unit latency, the L2 partition behavior, or the unified-core mixed-workload latency differs across Blackwell parts, the paper's architecture-level conclusions are SKU-specific. Also re-test the FP4 mma mapping with a newer CUDA version: if FP4 uses OMMA rather than QMMA, the reported 'fallback' is a transient software state, not an architectural fact.","supporting_citations":[{"cited_title":"Demystifying gpu microarchitecture through microbenchmarking,","cited_arxiv_id":null,"evidence_quote":"Establishes the pointer-chase cache-boundary approach used to map the memory hierarchy."},{"cited_title":"[Online]","cited_arxiv_id":null,"evidence_quote":"Defines %clock64 and the PTX instruction syntax used for all timing and mma kernels."},{"cited_title":"[Online]","cited_arxiv_id":null,"evidence_quote":"Documents the QMMA and OMMA SASS instructions used to interpret tensor-core instruction mapping."},{"cited_title":"A performance model for gpus with caches,","cited_arxiv_id":null,"evidence_quote":"Provides the basis for excluding first-run warm-up effects from reported latencies."},{"cited_title":"[Online]","cited_arxiv_id":null,"evidence_quote":"Provides the TensorRT inference engine used for the transformer power-efficiency case study."}],"review_version":1}