{"id":"c2db345c-d690-4d39-86eb-1ed838c3270d","arxiv_id":"2508.19087","paper_version":1,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":5.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":2,"one_line_summary":"A bit-sliced Tensor Core scheme with a bipolar integer encoding and adaptive kernel autotuning reports 1.65-3.99x speedups for ultra-low-bit GPTQ-quantized LLM inference on three NVIDIA GPUs.","lead":"This paper presents a GPU kernel system, APT-LLM, that runs quantized LLMs at arbitrary low bit widths (e.g., 1-bit weights with 2-bit activations) by decomposing matrices into bit slices and recombining them on NVIDIA Tensor Cores. On RTX 3090, 4090, and H800 GPUs it reports up to 3.99x speedups over FP16 inference and up to 2.16x over CUTLASS INT4 kernels.","discovery_kind":"new_method","skeptic_critique":{"model":"deepseek-v4-flash","headline":"The bit-wise reconstitution (Sec III-B, Fig 5/8) omits correction terms: if the 1-bit TC MMA returns AND/XOR popcounts rather than bipolar ±1 dot products, the shift-and-add formula does not compute WX.","rationale":"The paper's central claim is that arbitrary-precision MatMul can be reconstituted from 1-bit Tensor Core operations losslessly. The manuscript never pins down the TC's 1-bit semantics: it invokes both AND and XOR, and the reconstruction formula is shift-and-add alone. For 0/1 bits, AND and XOR popcounts are affine transforms of the true bipolar dot product, not equal to it. The missing correction terms depend on K and on per-row/per-column sums of the bit matrices; ignoring them would corrupt every output. This is the load-bearing mathematical step because the memory scheduling and kernel mapping optimize an algorithm whose correctness is asserted rather than shown. I am not accusing the authors of fabricated results: the perplexity table suggests a working system exists, but the write-up omits the derivation. A direct microbenchmark of the MMA primitive and an independent re-derivation of the reconstruction would settle the matter. If the corrections are present in code but omitted in prose, the conditional verdict should stand pending release of the artifact; if the shift-add formula is actually what runs and results still match, that would require documenting the hardware behavior. This is the most load-bearing concern; lack of error bars and artifact availability are important but secondary.","tokens_in":24161,"tokens_out":7059,"duration_ms":74943,"concrete_test":"On an RTX 3090, issue the 1-bit MMA used by CUTLASS (e.g., mma.sync m8n8k32 with b1 operands) with fixed bit matrices A,B having known row and column sums. Record the 32-bit accumulator and compare it against popcount(A&B), popcount(A^B), and sum_k(2A-1)(2B-1). Then form random 2-bit matrices W,X, compute WX using the Fig 5 shift-and-add procedure from the raw accumulators, and compare with exact integer WX. If they differ, add the missing correction terms and re-run the Table V perplexity experiment; if they match exactly, the concern is vacated.","verdict_should_be":"UNCHANGED","load_bearing_attack":"Sec III-B defines bipolar bit values as (2x_i-1)2^i and requires the 1-bit MMA result Y(i,j) to be the dot product of bipolar bit vectors for the Fig 5 shift-and-add reconstruction to close. The paper states only that TCs use 'AND or XOR logic gates' for 1-bit MatMul. If the MMA returns the AND popcount P=sum_k w_k x_k, then the true bipolar dot product is S=sum_k(2w_k-1)(2x_k-1)=4P-2*sum w -2*sum x + K. If it returns the XOR popcount Q=sum_k(w_k xor x_k), then S=K-2Q. Neither equals P or Q, so the published formula Y=sum_{i,j}2^{i+j}Y(i,j) is missing row-sum, column-sum, and K-dependent correction terms. These corrections are not derived anywhere in the paper, and no microbenchmark identifies which primitive the Ampere Tensor Core actually returns. Since every kernel result—and hence the speedups in Tables II-IV and the perplexity matches in Table V—depends on this reconstruction, the manuscript's central correctness argument is incomplete as written.","agreement_with_reader":"agree"},"referee_report":{"model":"deepseek-v4-flash","summary":"The paper proposes APT-LLM, a GPU acceleration scheme for arbitrary-precision quantized LLM inference on NVIDIA Tensor Cores. The scheme has three components: (i) a bipolar-INT data format intended to replace signed INT losslessly and to make bit-level operations uniform; (ii) a bit-wise MatMul reconstitution method that decomposes operands into bit planes, performs 1-bit Tensor-Core MatMuls, and reconstructs the output by shift-and-add; and (iii) a GPU memory scheduling strategy plus an adaptive kernel-mapping framework that selects kernel hyperparameters from a precomputed lookup table. The paper reports speedups over FP16 and CUTLASS INT4/INT8 baselines on RTX 3090, RTX 4090, and H800, and a perplexity comparison at W4A4 showing no degradation for several quantization methods.","tokens_in":24549,"tokens_out":8970,"duration_ms":79848,"significance":"If the central dataflow is correct, the paper addresses a genuine gap: existing Tensor Cores do not natively support the INT2/INT3 formats used by recent ultra-low-bit LLM quantization methods, and the proposed bit-sliced approach could make arbitrary precisions practical on current GPUs. The paper has concrete strengths: extensive kernel and end-to-end LLM benchmarks across three GPU generations; an ablation study separating memory scheduling from kernel mapping; and a W4A4 perplexity sanity check across multiple quantization methods. However, the mathematical description of the bit-wise reconstitution is incomplete, and the correctness of every reported speedup depends on it. The manuscript is therefore not yet self-contained enough to support its central claim as written.","major_comments":[{"comment":"The reconstitution formula in Fig. 5, Y = Σ_{i,j} 2^{i+j} Y(i,j), is only valid if the 1-bit Tensor-Core MatMul returns the exact bipolar dot product S = Σ_k (2A_{i,k}-1)(2C_{j,k}-1). The text states that TCs use 'AND or XOR logic gates' for 1-bit MatMul. If Y(i,j) is the AND popcount P = Σ A·C, then S = 4P - 2ΣA - 2ΣC + K; if it is the XOR popcount Q, then S = K - 2Q. Neither is equal to P or Q, and the required row-sum, column-sum, and K-dependent correction terms are not derived anywhere in the paper. Since every kernel result and downstream speedup depends on this reconstruction, the central correctness argument is incomplete. Please state exactly what Y(i,j) is, add the correction terms, and/or provide a microbenchmark identifying the primitive actually returned by the Ampere/Ada/Hopper Tensor Cores.","section":"Sec. III-B / Fig. 5"},{"comment":"The bipolar-INT conversion relation is stated as 'ˆx′ = 2ˆx + 1' in the text, but the definition of bipolar-INT given above is (x)_D = Σ_{i=0}^{n-1} (2x(i)−1)·2^i. For an individual bit b, the bipolar coefficient is 2b−1, not 2b+1. The 'sign bit flip' description in Fig. 4 is also not obviously consistent with the algebraic definition for two's-complement values. Because the claim 'seamlessly replace the INT format without any loss of accuracy' rests on this conversion, the paper needs to define the mapping unambiguously and correct the equation.","section":"Sec. III-A"},{"comment":"The accuracy evaluation is limited to W4A4. The speedup claims are made for W1A2, W2A2, and W3A4, but no perplexity or downstream-accuracy numbers are reported for these precisions. The paper argues the conversion is mathematically lossless, yet the actual quantized-model behavior at these precisions is not demonstrated. Additionally, latency numbers in Tables II–IV and the perplexity numbers in Table V are reported without error bars or run-to-run variance, which makes it difficult to judge whether the reported differences are significant. Please add accuracy results for the precisions used in the main speedup claims and report variance or at least multiple seeds.","section":"Sec. VI-C / Table V"}],"minor_comments":[{"comment":"The text says W1A2 approaches CUTLASS INT1 in the 64/4k/4k task, 'achieving a 92.5% speedup.' From Table II the intended statement is that APT W1A2 achieves 92.5% of CUTLASS INT1's speedup (6.40× vs. 6.92×), not a 92.5% speedup.","section":"Sec. VI-B1"},{"comment":"The decode-phase paragraph refers to 'the LLAMA3-7B model,' while the rest of the paper uses LLaMA3-8B. This is likely a typo and should be corrected.","section":"Sec. VI-B2"},{"comment":"The term 'Fragment' is capitalised and used as if it were a defined technical term, but it is not defined in the text. Define it explicitly (e.g., fragment as the register-resident tile used by a Tensor Core warp instruction) or use a different word.","section":"Sec. IV-B"},{"comment":"The displayed reconstitution equation in Fig. 5 is not numbered. Numbering this equation and explicitly defining the dimensions and data types of Y(i,j) would help the reader verify the dataflow.","section":"Fig. 5 / Sec. III-B"}],"recommendation":"major_revision","confidential_remarks":"The main mathematical gap is fixable: the authors can either add the missing correction terms and show how the implementation applies them, or demonstrate with a microbenchmark that the Tensor-Core 1-bit MMA returns the bipolar dot product. I would not reject on the current evidence, because the kernel may well be correct in implementation; the paper just does not yet describe it correctly. The accuracy coverage at W1A2/W2A2/W3A4 should also be strengthened before the arbitrary-precision claim is accepted."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"The short version: this is a plausible and fairly thorough systems paper on accelerating ultra-low-bit LLM inference on Tensor Cores, with real measurements across three GPU generations. The soft spot is exactly where the reader puts it: the bit-wise reconstruction formula in Sec III-B is missing the correction terms needed if the 1-bit TC MMA returns AND/XOR popcounts rather than bipolar dot products. As written, Fig 5's shift-and-add does not compute WX.\n\nWhat's actually here: a working kernel stack—bipolar-INT encoding, fragment-level recovery, shared-memory scheduling, adaptive kernel mapping—benchmarked against APNN-TC, BSTC, BTC, CUTLASS, and FP baselines. The kernel mapping ablation is a plus. The authors are also honest about the speedups being largest on Ampere and shrinking on Ada/Hopper. If the reconstruction issue gets sorted, the engineering contribution is real.\n\nThe first soft spot is the reconstruction derivation. They define bipolar-INT as (2x-1)2^i, then say the 1-bit MMA yields Y(i,j) and that summing shifted Y(i,j) recovers the product. That only holds if Y(i,j) is the dot product of ±1 vectors. They themselves say TCs use AND or XOR logic gates. For AND popcount P, the bipolar dot product is 4P - 2 sum w - 2 sum x + K; for XOR popcount Q, it's K - 2Q. Neither is P or Q, and the paper never states these corrections. Since every speedup table depends on the kernel producing correct results, this is load-bearing, not a nit. The authors need to either show the exact primitive they used on each architecture or provide the corrected reconstruction and a microbenchmark verifying it.\n\nSecond, the accuracy check only covers W4A4. The speedup claims are for W3A4, W2A2, W1A2, but Table V only validates perplexity at W4A4. They should report accuracy for the actual low-bit configs used in the kernel benchmarks.\n\nThird, minor: no artifact or code release, and no error bars. Common for systems papers, but still worth noting.\n\nThis paper is for the efficient-inference community. A reader tracking arbitrary-precision TC methods should see it; it's a credible extension of APNN-TC with better memory handling and a broader config space. But the central correctness argument has a hole that must be patched before the speedup numbers can be taken at face value.\n\nI'd send it to peer review—the engineering is substantial and the flaw is fixable—but I would not accept it in current form. The authors need to close the bit-reconstruction derivation and validate accuracy at the precisions actually benchmarked.","headline":"Useful kernel-engineering paper on arbitrary-precision LLM MatMul, but the bit-reconstruction math as written doesn't close; fixable, but central.","tokens_in":24975,"tokens_out":3053,"would_cite":false,"duration_ms":31485,"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":"By decomposing quantized LLM matrices into 1-bit bipolar pieces, APT-LLM runs arbitrary-precision inference on GPU Tensor Cores and reports up to 3.99x speedup over FP16 on RTX 3090.","keywords":["LLM inference acceleration","Tensor Cores","arbitrary-precision quantization","bipolar-INT","bit-wise matrix multiplication","GPU shared memory scheduling","adaptive kernel mapping","low-bit LLM"],"falsifier":"Run the proposed kernel on a small known 2-bit case, such as two 2x2 matrices containing both same-sign and opposite-sign bipolar entries, and compare the shift-and-add reconstructed output against the exact signed integer product. Any mismatch in the low bits, or an error that grows with K, would show that the raw 1-bit MMA output is not algebraic +1/-1 multiplication and that the written reconstruction is incomplete.","tokens_in":1944,"feed_emoji":"⚡","tokens_out":2155,"duration_ms":77406,"temperature":0.7,"pith_summary":"APT-LLM tries to close the gap between ultra-low-bit LLM quantization and what GPU Tensor Cores actually support. Its central device is a new data format, bipolar-INT, in which every bit represents -1 or +1 and conversion from signed INT is a lossless sign-bit flip, plus a bit-wise MatMul that decomposes operands into 1-bit matrices, runs them through Tensor Core 1-bit GEMMs, and reconstructs the exact product by shift-and-add. The paper claims this enables arbitrary precision with better Tensor Core utilization, reporting up to 3.99x speedup over FP16 and 2.16x over NVIDIA CUTLASS INT4 on RTX 3090. It also introduces shared-memory-centric data recovery and adaptive kernel mapping, which the ablation credits for a further 3.22x speedup over a naive implementation.","feed_headline":"Bit-level tensor cores run quantized LLMs up to 3.99x faster","feed_subtitle":"APT-LLM splits quantized matrices into bipolar 1-bit pieces and reassembles exact results at any precision.","key_machinery":"Bipolar-INT is a numeric format where bit i contributes +2^i or -2^i instead of 0 or 2^i; it is obtained from signed INT by flipping the sign bit, so hat-x-prime = 2*hat-x + 1. Because linear quantization W = s*hat-W + z becomes W = (s/2)*hat-W-prime + (z - s/2), the conversion is algebraically lossless after adjusting scale and zero. The carrying object is bit-wise MatMul reconstitution: decompose weight and activation into bit matrices W(i) and X(j), let the Tensor Core execute pairwise 1-bit GEMMs to produce intermediate matrices Y(i,j), then recover Y = sum over i,j of Y(i,j) * 2^(i+j).","core_discovery":"The paper's central claim is that any fixed-point INT quantized LLM can be converted, without retraining or meaningful accuracy loss, to bipolar-INT, where each bit contributes +2^i or -2^i. Because an n-bit value is then a linear combination of bit matrices, a MatMul of two such values is exactly the shift-and-add sum of pairwise 1-bit Tensor Core products. This makes INT2, INT3, and mixed precisions like W3A4 first-class citizens on GPU Tensor Cores that natively support only 1-bit and 4-bit integer GEMM. The authors further claim that performing the shift-and-add recovery inside shared memory or register fragments rather than global memory, and adaptively selecting block and warp tile siz","pith_inferences":["The mathematical correctness of every kernel result depends on an undocumented hardware detail: the 1-bit Tensor Core MMA returns the algebraic product of bipolar +1/-1 bits. If it instead returns bitwise AND or XOR on 0/1 values, the shift-and-add reconstruction needs row and column sum correction terms that the paper never states.","The same bit-decomposition strategy could be combined with native INT2 or INT4 MMA where available, reducing the number of intermediate bit-pair products and lowering recovery cost; the paper does not explore this hybrid path.","The adaptive kernel mapping uses a lookup table for common LLM matrix shapes. A closed-form cost model would generalize it to arbitrary shapes and would also make prediction possible for when the speedup disappears on GPUs with very fast native INT4 paths.","The perplexity tables show small drift after bipolar-INT conversion, which the paper attributes to storing the adjusted scale and zero in FP16/FP32; a numerical analysis of that rounding sensitivity would be a natural stress test of the lossless claim."],"forward_implications":["Ultra-low-bit W1A2, W2A2, and W3A4 quantized LLMs can run on existing Ampere, Ada, and Hopper Tensor Cores without first converting to supported INT4 or INT8 formats, yielding the reported up to 3.99x speedup over FP16 on RTX 3090.","Any standard linear-quantized model can switch to bipolar-INT by updating scale s to s/2 and zero z to z - s/2, with no retraining required.","Performing data recovery in shared memory or fragments substantially reduces latency; the ablation credits memory scheduling with 2.15x and kernel mapping with an additional 1.50x.","On newer GPUs the absolute speedups persist but shrink, reaching up to 2.44x over FP16 and 1.65x over CUTLASS integer baselines, because baseline throughput improved while the non-Tensor-Core recovery phase did not scale as well.","The adaptive kernel mapping shows that no single hyperparameter configuration is optimal for all LLM layers or phases, which motivates the lookup-table search used before inference."],"supporting_citations":[{"why":"The main prior arbitrary-precision Tensor Core MatMul design; APT-LLM extends its precision coverage and is benchmarked against it.","marker":"[54]"},{"why":"NVIDIA's low-bit GEMM kernel collection used as the INT1/INT4/INT8 baseline whose speedups the paper reports.","marker":"[63]"},{"why":"GPTQ quantization is used to produce the low-bit LLM weights for all end-to-end inference speedup measurements.","marker":"[10]"},{"why":"Microbenchmark study of Tensor Core latency, throughput, and numeric behavior that grounds the feasibility of 1-bit and bit-level integer MatMul.","marker":"[21]"},{"why":"Earlier binarized-soft-tensor-core MatMul design used as a comparison baseline for low-bit GEMM throughput.","marker":"[33]"},{"why":"Bit-tensor-core BNN acceleration work used as a comparison baseline on Turing-era Tensor Core behavior.","marker":"[34]"},{"why":"The LLaMA-3 model whose layer MatMul shapes motivate the kernel mapping and which is used in the end-to-end evaluations.","marker":"[2]"},{"why":"OmniQuant is one of the quantized models used in the W4A4 perplexity-preservation check.","marker":"[15]"}],"fun_headline_variants":["Arbitrary-precision tensor cores speed LLMs 3.99x","Bipolar-INT recipe makes any bit-width LLM fast","Shift-and-add tensor cores: 3.99x LLM speedup","From INT1 to INT4: tensor cores break precision barrier","Bitwise tensor cores hit 3.99x LLM speedup"],"cache_read_input_tokens":26752,"weakest_assumption_plain":"The reconstruction assumes that the 1-bit Tensor Core MMA returns the algebraic product of +1/-1 bipolar bits, so the final result is obtained by shift-and-add alone; if the hardware operation is actually bitwise AND, XOR, or equality over 0/1 bits, extra correction terms are needed and are not stated in the paper.","fun_headline_variants_meta":{"raw":{"variants":["Arbitrary-precision tensor cores speed LLMs 3.99x","Bipolar-INT recipe makes any bit-width LLM fast","Shift-and-add tensor cores: 3.99x LLM speedup","From INT1 to INT4: tensor cores break precision barrier","Bitwise tensor cores hit 3.99x LLM speedup"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000767,"raw_usage":{"total_tokens":3304,"prompt_tokens":882,"completion_tokens":2422,"prompt_tokens_details":{"cached_tokens":256},"prompt_cache_hit_tokens":256,"prompt_cache_miss_tokens":626,"completion_tokens_details":{"reasoning_tokens":2328}},"tokens_in":626,"tokens_out":2422,"duration_ms":17263,"temperature":1.0,"reasoning_tokens":2328,"cache_read_input_tokens":256,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-05T15:58:03.951816+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Run the proposed kernel on a small known 2-bit case, such as two 2x2 matrices containing both same-sign and opposite-sign bipolar entries, and compare the shift-and-add reconstructed output against the exact signed integer product. Any mismatch in the low bits, or an error that grows with K, would show that the raw 1-bit MMA output is not algebraic +1/-1 multiplication and that the written reconstruction is incomplete.","supporting_citations":[{"cited_title":"Apnn-tc: Accelerating arbitrary precision neural net- works on ampere gpu tensor cores,","cited_arxiv_id":null,"evidence_quote":"The main prior arbitrary-precision Tensor Core MatMul design; APT-LLM extends its precision coverage and is benchmarked against it."},{"cited_title":"CUTLASS,","cited_arxiv_id":null,"evidence_quote":"NVIDIA's low-bit GEMM kernel collection used as the INT1/INT4/INT8 baseline whose speedups the paper reports."},{"cited_title":"Optq: Accurate quantization for generative pre-trained transformers,","cited_arxiv_id":null,"evidence_quote":"GPTQ quantization is used to produce the low-bit LLM weights for all end-to-end inference speedup measurements."},{"cited_title":"Dissecting tensor cores via microbenchmarks: Latency, throughput and numeric behaviors,","cited_arxiv_id":null,"evidence_quote":"Microbenchmark study of Tensor Core latency, throughput, and numeric behavior that grounds the feasibility of 1-bit and bit-level integer MatMul."},{"cited_title":"Bstc: A novel binarized-soft-tensor-core design for acceler- ating bit-based approximated neural nets,","cited_arxiv_id":null,"evidence_quote":"Earlier binarized-soft-tensor-core MatMul design used as a comparison baseline for low-bit GEMM throughput."},{"cited_title":"Accelerating binarized neural networks via bit-tensor-cores in turing gpus,","cited_arxiv_id":null,"evidence_quote":"Bit-tensor-core BNN acceleration work used as a comparison baseline on Turing-era Tensor Core behavior."},{"cited_title":"Omniquant: Omnidirectionally calibrated quantization for large language models,","cited_arxiv_id":null,"evidence_quote":"OmniQuant is one of the quantized models used in the W4A4 perplexity-preservation check."}],"review_version":1}