{"id":"288acaa7-ac6d-422c-ae09-c11a61b3d548","arxiv_id":"2512.07004","paper_version":4,"verdict":"CONDITIONAL","confidence":"HIGH","novelty_score":6.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":4,"one_line_summary":"Software models of NVIDIA tensor cores reproduce GPU matrix-multiply results exactly on millions of random test inputs, for fp8/fp16/bf16/tf19 formats on nine GPU generations.","lead":"The authors reverse-engineer the numerical behavior of NVIDIA tensor cores from V100 to B200 and ship MATLAB models that replicate their inner-product results exactly. These models give numerical analysts a way to test mixed-precision algorithms without needing the actual GPUs.","discovery_kind":"extension","skeptic_critique":{"model":"deepseek-v4-flash","headline":"Unverified fourth carry-bit assumption (footnote 3) undermines the unqualified 'bit accurate' claim: the random normal suite cannot trigger the rare carry-into-bit-30 case, so model correctness on that corner is untested.","rationale":"The reader's weakest_assumption identifies the representativeness of normal-distribution sampling, which is a broad concern. My review finds a more specific, internally admitted gap: the unverified fourth carry bit (footnote 3). This gap is load-bearing because the model's internal accumulator width and carry handling directly determine the final rounded/truncated output, and the random normal test suite is extraordinarily unlikely to exercise the rare combination of aligned significands that would set the fourth carry bit. The paper itself states that no numerical test can verify this feature without assuming no intermediate normalization, making it an explicit limitation rather than an overlooked one. The proposed concrete test—a targeted enumeration of boundary exponent combinations—would either expose a mismatch (showing the model is not universally bit-accurate) or demonstrate that the fourth carry bit is functionally unobservable, thereby closing the gap. This does not change the reader's CONDITIONAL verdict: the unqualified 'bit accurate' claim still needs narrowing or additional validation across broader input distributions. The agreement is 'partial' because the reader's weakest_assumption focuses on the general sampling distribution, while my concern centers on a specific unverified hardware feature, though both point toward the same conclusion that the empirical validation is incomplete.","tokens_in":16655,"tokens_out":5829,"duration_ms":54681,"concrete_test":"Run a targeted search on the A100 fp16/bf16 tensor core (k=8) that attempts to force a carry into the fourth carry bit. Enumerate all combinations of exponents in the range [-14, 15] with significands chosen from boundary patterns (e.g., 1.0, max normal, 1+2^-10) for all 8 products, plus a final addend c selected to maximize the aligned sum. Compare hardware output to the MATLAB model for each case. If any mismatch appears, the carry-bit assumption is incorrect and the model must be revised. If all cases match, additionally run a variant of the model with only three carry bits; if both variants match hardware on all enumerated inputs, the fourth bit is functionally unobservable and the concern is resolved.","verdict_should_be":"UNCHANGED","load_bearing_attack":"In Section IV-A2 (A100) and Table II, the model uses a 30-bit accumulator with four carry bits, but the paper admits in footnote 3 that only three extra carry bits can be detected by testing, and the fourth is 'assumed to exist by logic,' with no test vector able to verify it unless one assumes no intermediate normalization. The validation in Section IV-B compares against 10^5 random standard-normal vectors. Such inputs have exponents concentrated near zero, so the aligned significand sum essentially never produces a carry into the fourth carry bit (which requires the sum to approach the top of the 2-integer-bit range). Therefore the 'bit accurate for every run' claim is not established for this corner of the input space. If the hardware actually lacks the fourth carry bit or performs intermediate normalization, there may exist inputs (e.g., maximal-magnitude products with accumulating signs) where the model output would differ from the GPU. This is not merely a generic representativeness concern; it is a specific, internally admitted unverified assumption in the model's arithmetic path that directly affects the central claim.","agreement_with_reader":"partial"},"referee_report":{"model":"deepseek-v4-flash","summary":"The manuscript develops MATLAB simulation models for the inner-product arithmetic of NVIDIA tensor cores on V100, A100, A2, A30, H100, H200, B200, L40S, and Ada RTX 1000 GPUs, covering fp8, fp16, bf16, and tf19 inputs with fp16/fp32 outputs. The modeling pipeline uses generalized numerical feature test vectors (GNFT) to set parameters such as FMA size, alignment bit widths, rounding modes, and accumulator size, then refines the models through an input-space search method (ISSM) that compares model output with GPU output on random inputs; the body reports 30 verification runs of 10^5 vectors (the abstract says 10^7). The paper reports bit-level agreement on all tested random configurations and demonstrates the toolbox for authenticating feature-detection test vectors and for multi-word GEMM emulation.","tokens_in":16993,"tokens_out":9691,"duration_ms":93793,"significance":"If the bit-accuracy result is taken at face value, the toolbox is a useful contribution: it gives the mixed-precision numerical analysis community open, inspectable models of current NVIDIA tensor cores, and the authors demonstrate a realistic application (multi-word GEMM) where architectural differences (integer bits, extra alignment bits, rounding/truncation, interleaving) change error behavior. The iterative GNFT+ISSM methodology and the reproducible GitHub code are concrete strengths. The main reservation is that the validation is empirical and covers only a narrow random input distribution; the models can be wrong in untested corners, and the paper itself flags one such corner (the fourth accumulator carry bit) that it cannot verify.","major_comments":[{"comment":"The unqualified 'bit accurate' claim is not established for the fourth carry bit of the A100-family accumulator. The paper states in footnote 3 that only three extra carry bits can be detected and the fourth is 'assumed to exist by logic' because no numerical test distinguishes final normalization/truncation from an addition with the fourth carry bit set. The randomized validation in §IV-B draws 10^5 standard-normal inputs; such inputs produce aligned sums that almost never approach the top of the 2-integer-bit accumulator, so the fourth carry bit is effectively untested. If the hardware lacks this bit or performs intermediate normalization, there exist plausible inputs where the model output would differ from the GPU. Please provide a targeted test vector for this corner, or restrict the bit-accuracy claim to the tested range.","section":"§IV-A2, footnote 3; §IV-B"},{"comment":"The validation evidence is narrower than the abstract suggests. The abstract says 'randomised input vectors of 10^7 values,' while §III-B and §IV-B say 10^5 random vectors per run. More fundamentally, the standard-normal distribution used by ISSM samples a tiny fraction of the input space (the paper itself notes ~10^77 possible fp8 inputs for k=16) and is concentrated in a small exponent range; it cannot be assumed representative for extreme exponents, subnormal combinations, or rounding boundaries. The term 'semi-exhaustive comparison' is therefore misleading. I recommend either adding targeted boundary/edge cases to the validation or explicitly reporting the verified input distribution in every claim of bit accuracy.","section":"Abstract; §III-B; §IV-B"}],"minor_comments":[{"comment":"Reference [28] is the Intel BFLOAT16 specification, not a source for a V100 tensor-core model; the phrase 'the model presented in [28, Fig. 2]' appears to be a citation error.","section":"§IV-A1"},{"comment":"The author name 'Xi et al.' for reference [11] should be 'Li et al.'.","section":"§IV-A1"},{"comment":"Typo: 'RXT 1000' should be 'RTX 1000'.","section":"Fig. 4 caption"},{"comment":"The instruction string 'wgmma.mma async.aligned.m64nNk32' has inconsistent capitalization for N/K; it should be checked against the PTX ISA naming.","section":"§IV-A5, Fig. 5 caption"},{"comment":"Conclusion: 'simular' should be 'simulator'.","section":"§VI"}],"recommendation":"major_revision","confidential_remarks":"The paper is publishable after revision if the bit-accuracy claims are qualified or the fourth carry bit is tested. The software artifact is a strength. The abstract should not overstate the validation."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"The thing to know: this paper reverse-engineers NVIDIA tensor core inner products and ships MATLAB models that bit-match the hardware on every vector they tested — 30 runs of 10^5 fresh random draws across nine GPU variants. That is real evidence, and the code is public.\n\nWhat is genuinely new is worth reading. On Hopper and Blackwell, fp8 via mma.sync is converted to fp16 and computed in interleaved pairs — an internal quirk that contradicts the obvious reading of the instruction set. V100 and A100 keep the largest product denormalized during alignment, which changes rounding in small but observable ways and corrects earlier published models. The H100 fp16 path is pinned to exactly two extra alignment bits. These are details that only come from careful probing of undocumented hardware, and the iterative GNFT/ISSM loop is a credible method for finding them.\n\nThe soft spots are real but narrow, and mostly about how the results are worded. The abstract's unqualified 'bit accurate' overstates the evidence: the validation distribution is standard normal, which does not exercise extreme exponent ranges, subnormal-heavy sums, or the rare overflow into the fourth carry bit. The authors admit in footnote 3 that the fourth carry bit is assumed, not verified. I think the stress-test note is right that this is a specific corner of the arithmetic path, not a generic sampling complaint. The fix is easy — say 'bit-accurate on the tested distribution' and add targeted adversarial distributions (max-magnitude products, subnormal swarms, saturated accumulations) before claiming more. Also, the abstract says 10^7 random vectors while the body says 10^5; minor, but for a validation-driven paper, consistency matters.\n\nThe circularity worry does not land. The loop refines the model on failing samples and breaks on a clean pass with new draws, so the final agreement is not on the refinement set. Fitting against a hardware oracle and then testing held-out samples is standard.\n\nAudience: mixed-precision algorithm developers, error analysts, and anyone designing test vectors for non-standard arithmetic. The multi-word GEMM demo illustrates how the models get used. This deserves a serious referee — the empirical claims should be checked by someone with access to the same GPUs — but the referee should push for the claim-narrowing and extra corner-case validation. If the authors do that, it becomes a solid reference.","headline":"Solid reverse-engineering of NVIDIA tensor cores with shipped, bit-accurate-on-tested-inputs MATLAB models; the 'bit accurate' claim needs a distribution caveat and the unverifiable fourth carry bit should be stated up front.","tokens_in":17423,"tokens_out":5227,"would_cite":true,"duration_ms":51520,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":["65G50","65Y04"],"pacs":[],"model":"deepseek-v4-flash","headline":"This paper presents software models that reproduce NVIDIA tensor-core inner-product arithmetic exactly at the bit level across nine GPU variants and all supported low-precision input formats.","keywords":["tensor cores","floating-point arithmetic","mixed precision","matrix multiplication","bit-accurate emulation","rounding modes","GPU numerical behavior","low-precision formats"],"falsifier":"Run a directed search on any of the nine GPUs over small k, for example exhaustively enumerating all fp8 input pairs for k=2 or k=4, or applying a stress-test strategy that targets exponent boundaries, subnormals, and exact ties, and compare the real GPU output against the corresponding MATLAB model; a single mismatch would refute the bit-accurate claim.","tokens_in":16567,"feed_emoji":"🧮","tokens_out":7060,"duration_ms":61872,"temperature":0.7,"pith_summary":"The paper sets out to build software models of NVIDIA tensor cores, the hardware matrix multipliers in V100, A100, H100, B200, and related GPUs, that reproduce the hardware's numerical output exactly, down to the last bit. To do this, the authors first pin down microarchitectural features such as how many products are fused into one addition, how many extra bits are kept during significand alignment, and where rounding is applied. They then refine the models by comparing against real GPUs on hundreds of thousands of random inputs, repeating until model and hardware agree bit-for-bit. If the central claim is right, researchers can study and predict tensor-core arithmetic without owning the hardware, validate test vectors offline, and emulate high-precision matrix multiplication through multi-word schemes.","feed_headline":"MATLAB models match NVIDIA tensor cores bit-for-bit","feed_subtitle":"Researchers can test mixed-precision algorithms against verified GPU behavior without needing the hardware.","key_machinery":"The carrying mechanism is the combination of Generalised Numerical Feature Testing (GNFT) and Input Space Search Method (ISSM). GNFT constructs test vectors that trigger bit-level output differences to identify features like FMA size and extra alignment bits; ISSM draws random inputs, compares model against hardware, and uses mismatches to refine the model. The models themselves are built by combining fixed-point arithmetic with a custom floating-point format simulator inside a generalized block-FMA structure that can be parameterized by FMA count, extra alignment bits, rounding mode, sticky-bit handling, and interleaving pattern.","core_discovery":"The central claim is that the numerical behavior of NVIDIA tensor cores can be captured in a parameterized software model that is bit-accurate against hardware. For each GPU and input/output format, the paper determines a small set of numerical features—FMA block size, extra alignment bits, rounding mode, accumulator width, whether products stay denormalized during accumulation, and whether the input vector is interleaved—and encodes them in a MATLAB toolbox. Randomized testing with 10^5 vectors per configuration drives an iterative refinement loop until GPU and model outputs match exactly for every run. The paper also reports specific microarchitectural findings, including that V100 aligns","pith_inferences":["The bit-accuracy validation samples inputs from a standard normal distribution, so the claim may not cover extreme exponent ranges, subnormal boundaries, or rare rounding ties; a directed stress test could reveal gaps the random ensemble misses.","The finding that fp8 mma on Hopper/Blackwell is computed through fp16 tensor cores with interleaved pairs implies that software using fp8 may inherit fp16 rounding behavior in unexpected ways, a connection the paper notes but does not explore.","The same methodology would translate to non-NVIDIA matrix multipliers, and the parameterized model makes cross-vendor numerical comparison straightforward, although such an extension is not attempted here.","If the models are widely adopted, they could serve as an oracle for formal verification of tensor-core arithmetic, but exhaustive proof would require coverage guarantees beyond randomized sampling."],"forward_implications":["Mixed-precision algorithm developers can test and debug numerical behavior without direct access to GPU hardware.","Test vectors designed to probe tensor-core numerical features can be validated against known-correct models before being run on new hardware.","Multi-word emulation of high-precision matrix multiplication can be evaluated across different GPU generations, revealing accuracy differences attributable to FMA size, alignment bits, and rounding modes.","The GNFT-plus-randomized-refinement loop can be applied to future tensor-core designs to build models quickly.","Bit-accurate models support reproducible numerical experiments and can inform ongoing standardization of low-precision arithmetic."],"fun_headline_variants":["Bit-accurate MATLAB models emulate NVIDIA tensor cores","Model tensor core rounding without the hardware","NVIDIA tensor core quirks now reproducible in MATLAB","Software mimics V100 to B200 tensor cores exactly","Test mixed-precision GPU code without a GPU"],"cache_read_input_tokens":2304,"weakest_assumption_plain":"That the roughly 10^5 random test vectors drawn from a standard normal distribution represent all inputs the hardware will ever see; if tensor-core behavior differs for extreme exponents, subnormal combinations, or boundary rounding cases outside this sample, the bit-accuracy claim fails.","fun_headline_variants_meta":{"raw":{"variants":["Bit-accurate MATLAB models emulate NVIDIA tensor cores","Model tensor core rounding without the hardware","NVIDIA tensor core quirks now reproducible in MATLAB","Software mimics V100 to B200 tensor cores exactly","Test mixed-precision GPU code without a GPU"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000159,"raw_usage":{"total_tokens":1088,"prompt_tokens":790,"completion_tokens":298,"prompt_tokens_details":{"cached_tokens":256},"prompt_cache_hit_tokens":256,"prompt_cache_miss_tokens":534,"completion_tokens_details":{"reasoning_tokens":226}},"tokens_in":534,"tokens_out":298,"duration_ms":3339,"temperature":1.0,"reasoning_tokens":226,"cache_read_input_tokens":256,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-03T17:59:48.252175+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Run a directed search on any of the nine GPUs over small k, for example exhaustively enumerating all fp8 input pairs for k=2 or k=4, or applying a stress-test strategy that targets exponent boundaries, subnormals, and exact ties, and compare the real GPU output against the corresponding MATLAB model; a single mismatch would refute the bit-accurate claim.","supporting_citations":[],"review_version":2}