REVIEW 3 major objections 6 minor 36 references
Accelerating GPU Inference of Large Language Models with Moderately Unstructured Sparse Weight Matrices
T0 review · 3 major / 6 minor · reviewed 2026-07-13 · grok-4.5
Pith's one-line read A three-layer sparse format and pipelined SpMM kernel finally beats dense matrix multiply on HBM GPUs at the ~50% unstructured sparsity that preserves LLM quality.
desk verdict First real win for moderate unstructured SpMM over dense cuBLAS on HBM GPUs, via a clean three-layer format and TC+CUDA pipeline that actually ships. read the letter →
The pith
A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.
The reading
What carries the argument
The three-layer storage format (Sparse-TC + Slot-Filling via parallel differential distance + residual CSR) together with the overlapping SpMM pipeline that co-schedules sparse tensor cores and CUDA cores.
What would settle it
Measure wall-clock SpMM time of the proposed kernel against cuBLAS on an H100 for a representative 50%-sparse weight matrix (e.g., 66k imes16k) with batch size 32; if the sparse kernel is not faster, the central claim fails.
Extended reading notes
Core claim
At the moderate unstructured sparsity levels (~50%) that keep LLM perplexity acceptable, a three-layer matrix format—Sparse-TC for 2:4 tensor-core blocks, Slot-Filling with parallel differential-distance encoding for redistributed non-zeros, and a lightweight residual CSR layer—combined with a pipeline that jointly drives sparse tensor cores and CUDA cores, produces SpMM kernels that finally outperform dense matrix multiplication on HBM GPUs.
Load-bearing premise
After the 2:4 pattern is extracted, the remaining surplus non-zeros stay small enough (~18%) that they fit into empty slots with only 4-bit differential codes and a residual under 1%, so CUDA-core work can be fully hidden behind tensor-core compute and HBM traffic.
Editorial extensions
If this is right
- Pruned LLMs at quality-preserving 50% unstructured sparsity can now be served faster than their dense counterparts on HBM GPUs.
- Decode-phase inference pipelines that already separate prefill from decode gain an immediate throughput boost without further model changes.
- Storage and kernel designs that previously targeted extreme sparsity can be re-targeted to the moderate-sparsity regime that actually matters for LLMs.
- End-to-end token throughput for OPT-scale models rises by up to 41% relative to prior sparse systems while still reducing peak memory versus dense baselines.
Reading between the lines
- If the same redistribution idea were applied to structured 2:4-plus-residual formats already used by other sparse libraries, those libraries might also cross the dense baseline on Hopper.
- The decode-phase focus suggests that systems that keep prefill and decode on separate hardware pools can adopt the kernel without redesigning the prefill path.
- Because the residual layer stays under 1%, the approach may remain viable even if future pruning methods slightly increase irregularity, provided the 2:4 extraction still captures the bulk of non-zeros.
Editorial analysis
A structured set of objections, weighed in public.
Referee Report
Summary. The paper addresses the gap that existing SpMM GPU kernels fail to beat dense cuBLAS at the moderate (~50%) unstructured sparsity levels that preserve LLM quality under pruning (Wanda/RIA). It introduces a three-layer storage format: a Sparse-TC layer that extracts 2:4 patterns for sparse tensor cores, a Slot-Filling layer that redistributes surplus non-zeros (~17.7%) into empty 2:4 slots via 4-bit parallel differential-distance (PDD) encoding (with ~2.13% dummy points), and a residual CSR layer for the remaining <1% of non-zeros. A co-designed SpMM kernel pipelines sparse-tensor-core MMA (wgmma), CUDA-core PDD decoding, and residual work while overlapping HBM loads with on-chip execution. On H100 HBM3 the kernel is reported to be the first to outperform dense matrix multiplication at 50% sparsity, delivering up to 1.64× over SpInfer and up to 1.41× end-to-end tokens/s over FlashLLM on Wanda-pruned OPT-30B/66B (decode phase, batch 8–32).
Significance. If the empirical claims hold, the work is the first practical demonstration that moderate unstructured sparsity can yield net speedups over dense tensor-core GEMM on modern HBM GPUs—the regime that actually preserves LLM quality. The three-layer format and PDD encoding are concrete, implementable contributions; the released source code (https://github.com/moui0/cudac) and the systematic kernel/end-to-end evaluation against SpInfer (EuroSys’25 Best Paper), FlashLLM, SparTA and Sputnik strengthen reproducibility. The honest limitation discussion (high sparsity, large-N prefill) further increases credibility. The result is directly relevant to production LLM serving on Hopper/Ampere-class hardware.
major comments (3)
- §3.2 states that “around 17.7%” of non-zeros fall into the Slot-Filling layer, dummy points are 2.13%, and the residual layer is “typically less than 1%.” These occupancy figures are load-bearing for the claim that CUDA-core decoding and residual work remain fully overlappable (Figs. 5, 7). Please report the measurement protocol (which models/matrices, min/max/mean across layers) and whether the same fractions hold under SparseGPT and RIA, not only Wanda.
- §5.3 and Fig. 10: for OPT-66B the authors use 2 GPUs while FlashLLM requires 4 GPUs “due to its large memory footprint.” Although the caption says “per GPU,” the absolute system throughput and the memory-footprint comparison (Fig. 11) become harder to interpret. Please either (a) report multi-GPU wall-clock tokens/s under identical GPU counts or (b) explicitly normalize by total HBM capacity so that the 1.41× claim remains unambiguous.
- §6 and Fig. 12 correctly note that the kernel is slower than cuBLAS for large N (prefill). Given that the abstract and introduction advertise “GPU inference of LLMs,” a short quantitative statement of the decode-time fraction under realistic request-length distributions (or under disaggregated prefill/decode serving) is needed to bound the practical end-to-end impact of a decode-only optimization.
minor comments (6)
- Fig. 1 caption and §5.1: clarify whether “cuBLAS_TC” is the absolute best dense baseline (e.g., cuBLASLt with TF32/FP16 tensor cores) or a plain cuBLAS call; the “first to outperform dense” claim depends on it.
- Algorithm 1: the 16-bit first-offset + 4-bit deltas encoding is clear, but a one-sentence statement of the maximum supported tile size / maximum distance before dummy insertion would help implementers.
- Fig. 3 storage comparison: add absolute numbers (or a secondary axis) so that the metadata overhead of CSR vs. PDD vs. bitmap is readable without visual estimation.
- §4.2: the choice of 64×128 tiles for four unrolled wgmma operations is stated without a brief sensitivity study; a one-row ablation (or reference to the Hopper programming guide) would strengthen the design rationale.
- Typos / notation: “Splnfer” in Fig. 1 legend; “mixture-of-exports” (§2.1) should be “mixture-of-experts”; “parallel differential distance for matrix compression” is occasionally abbreviated inconsistently as PDD vs. “PDD code.”
- Table 1: include the exact pruning hyperparameters (calibration set size, etc.) so that the quality-vs-sparsity trade-off can be reproduced.
Circularity Check
No significant circularity: purely empirical systems paper whose speedups are measured against external baselines on public hardware.
full rationale
The paper proposes a three-layer sparse storage format (Sparse-TC + Slot-Filling with PDD + Residual CSR) and a co-designed SpMM kernel that pipelines sparse tensor cores with CUDA-core decoding, then reports measured kernel and end-to-end speedups versus cuBLAS, SpInfer, FlashLLM, SparTA and Sputnik on H100 HBM GPUs. All quantitative claims (1.64 imes kernel, 1.41 imes end-to-end, first to beat dense at ~50 % unstructured sparsity) are direct wall-clock or TFLOPS measurements; no parameter is fitted to data and later re-presented as a prediction, no equation reduces a claimed result to an author-defined quantity by construction, and no uniqueness theorem or ansatz is imported via self-citation. The residual-size and overlap assumptions are validated by the same external measurements (Figs. 3, 5, 8–11) rather than by circular definition. Consequently the derivation chain is self-contained and non-circular.
Assumptions & free parameters
free parameters (3)
- PDD differential bit-width =
4 bits
- Tile size for wgmma loop =
64×128
- Number of warp groups / tiles in pipeline =
2
assumptions (4)
- domain assumption NVIDIA Hopper sparse tensor cores accelerate only 2:4 structured sparsity via wgmma and require operands in that layout.
- domain assumption At ~50% unstructured sparsity produced by Wanda/RIA, a non-negligible fraction of non-zeros can be extracted into 2:4 windows while the surplus (~17.7%) and residual (<1%) remain small enough for the Slot-Filling and Residual layers.
- domain assumption Dense activation matrices in the decode phase (small N) fit in L2 cache, allowing hierarchical loading that overlaps with sparse weight traffic from HBM.
- standard math Standard linear-algebra and GPU programming model (CUDA, shared memory, asynchronous execution between CUDA and tensor cores).
invented entities (2)
-
Three-layer sparse matrix storage format (Sparse-TC + Slot-Filling + Residual)
independent evidence
-
Parallel differential distance (PDD) encoding
independent evidence
Cite this review
Pith. "Pith review of Accelerating GPU Inference of Large Language Models with Moderately Unstructured Sparse Weight Matrices." pith.science (2026). https://pith.science/paper/EYVQDWKL
@misc{pith2026260708786,
author = {Pith},
title = {Pith review of: Accelerating GPU Inference of Large Language Models with Moderately Unstructured Sparse Weight Matrices},
year = {2026},
howpublished = {\url{https://pith.science/paper/EYVQDWKL}},
note = {Machine review of arXiv:2607.08786}
}
read the original abstract
With the growing deployment of large language models (LLMs), LLM inference cost has become a key challenge. Pruning techniques that introduce sparsity into weight matrices can accelerate inference. However, maintaining model quality typically limits pruning to moderate unstructured sparsity (around 50\%). At these sparsity levels, none of the existing GPU kernels for sparse matrix multiplication (SpMM) can outperform their dense counterparts. This paper proposes an efficient GPU inference method for LLMs with moderate sparsity. We propose a three-layer matrix storage format comprising: (i) a Sparse-TC layer enabling sparse tensor cores to accelerate SpMM; (ii) a Slot-Filling layer using parallel differential distance for matrix compression while supporting low-cost on-chip decoding; (iii) a lightweight Residual Layer ensuring correct SpMM computation. Building on this format, we design a SpMM kernel that jointly utilizes sparse tensor cores and CUDA cores. This design enables an efficient execution pipeline and overlaps on-chip computation with memory access. Evaluations show that our work is the first to outperform dense matrix multiplication on modern GPUs equipped with high-bandwidth memory (HBM). It achieves up to 1.64x kernel-level speedup over SpInfer (EuroSys'25, Best paper) and up to 1.41x end-to-end speedups over FlashLLM (VLDB'24). Our source code: https://github.com/moui0/cudac.
Figures
Figures from the paper (7 more)
Reference graph
Works this paper leans on
-
[1]
AMD. 2025. AMD CDNA Architecture. https://www.amd.com/en/technologies/ cdna.html
2025
-
[2]
Rewon Child, Scott Gray, Alec Radford, and Ilya Sutskever. 2019. Generating long sequences with sparse transformers.arXiv preprint arXiv:1904.10509(2019)
arXiv 2019
-
[3]
Guohao Dai, Guyue Huang, Shang Yang, Zhongming Yu, Hengrui Zhang, Yufei Ding, Yuan Xie, Huazhong Yang, and Yu Wang. 2022. Heuristic adaptability to input dynamics for spmm on gpus. InProceedings of the 59th ACM/IEEE Design Automation Conference. 595–600
2022
-
[4]
Ruibo Fan, Wei Wang, and Xiaowen Chu. 2023. Fast sparse gpu kernels for accelerated training of graph neural networks. In2023 IEEE International Parallel and Distributed Processing Symposium (IPDPS). IEEE, 501–511
2023
-
[5]
Ruibo Fan, Wei Wang, and Xiaowen Chu. 2024. Dtc-spmm: Bridging the gap in accelerating general sparse matrix multiplication with tensor cores. InProceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 3. 253–267
2024
-
[6]
Ruibo Fan, Xiangrui Yu, Peijie Dong, Zeyu Li, Gu Gong, Qiang Wang, Wei Wang, and Xiaowen Chu. 2025. SpInfer: Leveraging Low-Level Sparsity for Efficient Large Language Model Inference on GPUs. InProceedings of the Twentieth Euro- pean Conference on Computer Systems. 243–260
2025
-
[7]
William Fedus, Barret Zoph, and Noam Shazeer. 2022. Switch transformers: Scaling to trillion parameter models with simple and efficient sparsity.Journal of Machine Learning Research23, 120 (2022), 1–39
2022
-
[8]
Elias Frantar and Dan Alistarh. 2023. Sparsegpt: Massive language models can be accurately pruned in one-shot. InInternational Conference on Machine Learning. PMLR, 10323–10337
2023
Show all 36 references
-
[9]
Trevor Gale, Matei Zaharia, Cliff Young, and Erich Elsen. 2020. Sparse gpu kernels for deep learning. InSC20: International Conference for High Performance Computing, Networking, Storage and Analysis. IEEE, 1–14
2020
-
[10]
Daya Guo, Dejian Yang, Haowei Zhang, Junxiao Song, Ruoyu Zhang, Runxin Xu, Qihao Zhu, Shirong Ma, Peiyi Wang, Xiao Bi, et al . 2025. Deepseek-r1: Incentivizing reasoning capability in llms via reinforcement learning.arXiv preprint arXiv:2501.12948(2025)
2025 arXiv
-
[11]
Ahan Gupta, Yueming Yuan, Devansh Jain, Yuhao Ge, David Aponte, Yanqi Zhou, and Charith Mendis. 2025. SPLAT: A framework for optimised GPU code- generation for SParse reguLar ATtention.Proceedings of the ACM on Programming Languages9, OOPSLA1 (2025), 1632–1660
2025
-
[12]
Changwan Hong, Aravind Sukumaran-Rajam, Israt Nisa, Kunal Singh, and Pon- nuswamy Sadayappan. 2019. Adaptive sparse tiling for sparse matrix multiplica- tion. InProceedings of the 24th Symposium on Principles and Practice of Parallel Programming. 300–314
2019
-
[13]
Jiedong Lang, Zhehao Guo, and Shuyu Huang. 2024. A comprehensive study on quantization techniques for large language models. In2024 4th International Conference on Artificial Intelligence, Robotics, and Communication (ICAIRC). IEEE, 224–231
2024
-
[14]
Shigang Li, Kazuki Osawa, and Torsten Hoefler. 2022. Efficient quantized sparse matrix operations on tensor cores. InSC22: International Conference for High Performance Computing, Networking, Storage and Analysis. IEEE, 1–15
2022
-
[15]
Sharan Narang, Greg Diamos, Shubho Sengupta, and Erich Elsen. 2017. Exploring Sparsity in Recurrent Neural Networks. InInternational Conference on Learning Representations
2017
-
[16]
Nvidia. 2025. Basic Linear Algebra on NVIDIA GPUs. https://developer.nvidia. com/cublas
2025
-
[17]
Nvidia. 2025. A High-Performance CUDA Library for Sparse Matrix-Matrix Multiplication. https://docs.nvidia.com/cuda/cusparselt/
2025
-
[18]
Nvidia. 2025. NVIDIA Ampere Architecture. https://www.nvidia.com/en-sg/data- center/ampere-architecture
2025
-
[19]
Nvidia. 2025. NVIDIA Hopper Architecture. https://www.nvidia.com/en-sg/data- center/technologies/hopper-architecture
2025
-
[20]
Hyungjun Oh, Kihong Kim, Jaemin Kim, Sungkyun Kim, Junyeol Lee, Du-seong Chang, and Jiwon Seo. 2024. Exegpt: Constraint-aware resource scheduling for llm inference. InProceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Op...
2024
-
[21]
Patrik Okanovic, Grzegorz Kwasniewski, Paolo Sylos Labini, Maciej Besta, Flavio Vella, and Torsten Hoefler. 2024. High Performance Unstructured SpMM Compu- tation Using Tensor Cores. InSC24: International Conference for High Performance Computing, Networking, Storage and Analy...
2024
-
[22]
Meng Pang, Xiang Fei, Peng Qu, Youhui Zhang, and Zhaolin Li. 2024. A row decomposition-based approach for sparse matrix multiplication on GPUs. In Proceedings of the 29th ACM SIGPLAN Annual Symposium on Principles and Practice of Parallel Programming. 377–389
2024
-
[23]
Pratyush Patel, Esha Choukse, Chaojie Zhang, Aashaka Shah, Íñigo Goiri, Saeed Maleki, and Ricardo Bianchini. 2024. Splitwise: Efficient generative llm inference using phase splitting. In2024 ACM/IEEE 51st Annual International Symposium on Computer Architecture (ISCA). IEEE, 118–132
2024
-
[24]
Ruoyu Qin, Zheming Li, Weiran He, Mingxing Zhang, Yongwei Wu, Weimin Zheng, and Xinran Xu Mooncake. 2024. Kimi’s kvcache-centric architecture for llm serving.arXiv preprint arXiv:2407.00079(2024)
2024 arXiv
-
[25]
Alec Radford, Jeffrey Wu, Rewon Child, David Luan, Dario Amodei, Ilya Sutskever, et al. 2019. Language models are unsupervised multitask learners.OpenAI blog 1, 8 (2019), 9
2019
-
[26]
2003.Iterative methods for sparse linear systems
Yousef Saad. 2003.Iterative methods for sparse linear systems. SIAM
2003
-
[27]
Mingjie Sun, Zhuang Liu, Anna Bair, and J Zico Kolter. 2023. A Simple and Effec- tive Pruning Approach for Large Language Models. InThe Twelfth International Conference on Learning Representations
2023
-
[28]
Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Łukasz Kaiser, and Illia Polosukhin. 2017. Attention is all you need.Advances in neural information processing systems30 (2017)
2017
-
[29]
Yuke Wang, Boyuan Feng, and Yufei Ding. 2022. QGTC: accelerating quantized graph neural networks via GPU tensor core. InProceedings of the 27th ACM SIGPLAN symposium on principles and practice of parallel programming. 107–119
2022
-
[30]
Yuke Wang, Boyuan Feng, Zheng Wang, Guyue Huang, and Yufei Ding. 2023. {TC-GNN}: Bridging sparse {GNN} computation and dense tensor cores on {GPUs }. In2023 USENIX Annual Technical Conference (USENIX ATC 23). 149–164
2023
-
[31]
Haojun Xia, Zhen Zheng, Yuchao Li, Donglin Zhuang, Zhongzhu Zhou, Xiafei Qiu, Yong Li, Wei Lin, and Shuaiwen Leon Song. 2023. Flash-LLM: Enabling Cost-Effective and Highly-Efficient Large Generative Model Inference with Un- structured Sparsity.Proceedings of the VLDB Endowment...
2023
-
[32]
Zeyu Xue, Mei Wen, Zhaoyun Chen, Yang Shi, Minjin Tang, Jianchao Yang, and Zhongdi Luo. 2023. Releasing the potential of tensor core for unstructured spmm using tiled-csr format. In2023 IEEE 41st International Conference on Computer Design (ICCD). IEEE, 457–464
2023
-
[33]
Zihao Ye, Ruihang Lai, Junru Shao, Tianqi Chen, and Luis Ceze. 2023. Sparsetir: Composable abstractions for sparse compilation in deep learning. InProceedings of the 28th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Vol...
2023
-
[34]
Susan Zhang, Stephen Roller, Naman Goyal, Mikel Artetxe, Moya Chen, Shuohui Chen, Christopher Dewan, Mona Diab, Xian Li, Xi Victoria Lin, et al. 2022. Opt: Open pre-trained transformer language models.arXiv preprint arXiv:2205.01068 (2022)
2022 arXiv
-
[35]
Yingtao Zhang, Haoli Bai, Haokun Lin, Jialin Zhao, Lu Hou, and Carlo Vittorio Cannistraci. 2024. Plug-and-play: An efficient post-training pruning method for large language models. InThe Twelfth International Conference on Learning Representations
2024
-
[36]
Ningxin Zheng, Bin Lin, Quanlu Zhang, Lingxiao Ma, Yuqing Yang, Fan Yang, Yang Wang, Mao Yang, and Lidong Zhou. 2022. {SparTA}:{Deep-Learning} Model Sparsity via {Tensor-with-Sparsity-Attribute}. In16th USENIX Symposium on Operating Systems Design and Implementation (OSDI 22). 213–232
2022
Reviewed July 13, 2026 · model on record in the stance chip above.
Discussion (0). Sign in to comment.