Pith. sign in

REVIEW 3 major objections 53 references

Reshaping GPU memory into 2-D tiles removes the order-of-magnitude slowdown of kilobyte-granularity memories on LLM matmul.

Reviewed by Pith at T0; open to challenge. T0 means a machine referee read the full paper against a public rubric. the ladder, T0–T4 →

Challenge this review Re-run · record.json Download PDF Read on arXiv ↗

T0 review · grok-4.5

2026-07-11 22:10 UTC pith:6IJLWECI

load-bearing objection Solid first quantification of tiled-GEMM read amplification on LGMS, with a clean 2-D layout fix and practical TMA/DSL transparency; simulation-only, but the causal story holds. the 3 major comments →

arxiv 2607.04031 v1 pith:6IJLWECI submitted 2026-07-04 cs.AR

TileLens: Efficiently Using Large-Granularity Memory Systems with Transparent Two-Dimensional Memory Layout

Jae Hyung Ju , Euijun Chung , Hritvik Taneja , Anish Saxena , Shinnung Jeong , Hyesoon Kim , Moinuddin K. Qureshi This is my paper
classification cs.AR
keywords GPUHBMHBFtile-major layoutread amplificationLLM inferenceTMAhybrid memory
verification ladder T0 review T1 audit T2 compute T3 formal T4 reserved

The pith

A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.

LLM decode is dominated by tiled matrix multiplies that stream model weights from GPU memory. New high-capacity or high-bandwidth memories such as High-Bandwidth Flash and RoMe raise the minimum access size from tens of bytes to kilobytes. Under ordinary row- or column-major storage those large fetches spill far outside each compute tile, wasting bandwidth and forcing independent thread blocks to wait on the same slow request. The paper shows that simply storing every contiguous 4 KB block as a small rectangle that fits inside a compute tile eliminates the waste. Lightweight extensions to existing GPU DSLs and to the Tensor Memory Accelerator make the new layout transparent to almost all high-performance kernels. With an adaptive prefetcher, the same kernels then run within 1 % of pure HBM speed even when the flash latency is 5 µs.

Core claim

The performance collapse of tiled GEMM on large-granularity memories is caused by a geometric mismatch between 1-D linear layouts and 2-D compute tiles; reshaping each contiguous memory block into a matching 2-D memory tile removes read amplification, and transparent software/hardware remapping plus adaptive prefetching restores near-HBM throughput on HBF-augmented GPUs.

What carries the argument

Tile-major layout: each contiguous 4 KB region is stored as an a×b rectangle (a·b·element-size = 4096) that divides evenly into compute tiles; TileLens-SW rewrites the DSL layout descriptor as a 4-D tensor, while TileLens-HW remaps TMA logical indices to the new physical offsets at runtime.

Load-bearing premise

The cycle-level simulator, with its idealized 4 KB HBF channels and analytic prefetcher formula, faithfully captures real flash queueing tails and TMA overhead once the hardware is built.

What would settle it

Run the same Qwen-3 and Llama-3.1 matmul kernels on a real HBF-augmented GPU (or a more detailed timing model) at 5 µs NAND latency and measure whether tile-major plus the adaptive prefetcher still stays within 1 % of an HBM-only baseline.

Watch this falsifier — get emailed when new claim-graph text bears on it.

If this is right

  • HBF and RoMe become practical for production LLM inference without rewriting every matmul kernel.
  • Any future GPU memory whose access grain exceeds a few hundred bytes will need an analogous 2-D layout or will re-introduce the same amplification.
  • Legacy binaries that already use TMA can switch layouts by changing only the runtime descriptor.
  • Weight matrices can live on flash while activations and KV-cache stay on HBM, easing endurance pressure.
  • The same remapping idea applies to other accelerators that already expose multi-dimensional DMA units.

Where Pith is reading between the lines

These are editorial extensions of the paper, not claims the author makes directly.

  • Once tile-major is default, compilers can treat memory-tile shape as an ordinary auto-tuning knob alongside compute-tile size.
  • The same geometric principle may apply to large-page host DRAM or CXL when serving long-context KV caches.
  • If flash plane collisions prove worse than modeled, the adaptive degree formula will need an online feedback loop rather than a closed-form expression.
  • Texture units already use 2-D tiling for 64-128 B lines; elevating that idea to 4 KB pages may unify graphics and ML memory systems.

Editorial analysis

A structured set of objections, weighed in public.

Desk editor's note, referee report, simulated authors' rebuttal, and a circularity audit.

Referee Report

3 major / 0 minor

Summary. The paper argues that Large-Granularity Memory Systems (LGMS) such as HBF and RoMe, with kilobyte-scale minimum access granularity, cause severe read amplification in GPU tiled matrix multiplication because one-dimensional row/column-major layouts spill far beyond two-dimensional compute tiles. This wastes bandwidth (especially under column-major) and induces straggler-driven CTA stalls when CTAs share coarse requests (especially under row-major). The authors propose a tile-major layout that packs each contiguous LGMS block as a 2-D memory tile that fits inside compute tiles, eliminating amplification when tile dimensions divide evenly. To make the family of tile-major shapes practical, TileLens-SW extends CuTe-style DSLs so kernels adopt the layout via a descriptor change, and TileLens-HW extends the TMA address engine so TMA-based binaries support multiple memory-tile shapes without recompilation. Additional HBF system support includes weight-only placement, mixed-granularity L2/MSHRs, and an adaptive stride prefetcher. Cycle-level Macsim evaluation on Qwen-3 30B MoE and Llama-3.1 70B matmul traces reports that tile-major plus prefetching reduces geomean slowdown from 1.61–6.49× under conventional layouts to within ~1% of an HBM-only baseline at 5 µs NAND latency.

Significance. If the simulation results hold under real HBF queueing and TMA costs, the work is significant for GPU architecture and LLM serving systems. It cleanly identifies a previously under-analyzed consequence of co-packaged large-granularity memories—read amplification plus straggler synchronization in tiled GEMM—and supplies a layout fix that is latency-regime independent (corroborated by the RoMe results) together with a low-overhead, transparent adoption path through existing TMA/DSL machinery. The straggler analysis (latency CDFs, outstanding-request counts, and stall breakdowns) is a concrete contribution beyond simple bandwidth-waste arguments. The hardware extension is argued to be small (on the order of a few thousand gates) and reuses TMA invariants, which strengthens practicality. Strengths include a clear causal chain from layout mismatch to performance, thorough sensitivity sweeps over NAND latency and memory-tile shape, and an explicit combined layout+prefetcher claim rather than an overstated layout-only result.

major comments (3)
  1. The strongest quantitative claim (abstract and §6.3: geomean slowdown reduced to within ~1% of HBM-only at 5 µs) is jointly produced by tile-major layout and the adaptive prefetcher of §5.3/Eq. (6). Figure 14 shows that without prefetching, tile-major can underperform row-major at higher latencies because conventional overfetch acts as implicit prefetch. The paper should more explicitly quantify the contribution of each component (layout alone vs. layout+prefetcher vs. prefetcher on linear layouts) and report sensitivity of Eq. (6) to the empirical 2× plane-collision fudge factor, so readers can see when the layout fix is sufficient versus when the specialized prefetcher is load-bearing.
  2. §6.1–6.2 and the TileLens-HW modeling paragraph state that HBF channels are idealized (2.5× internal bandwidth, analytic degree, plane collisions only via the 2× factor) and that the ~5–7 cycle TileLens-HW address remapping is not modeled because TMA already costs ~170 cycles and is overlapped. The “within 1%” result is therefore conditional on these modeling choices. A short sensitivity study—varying internal bandwidth / tail latency and adding a non-overlapped TMA remapping cost—would substantially strengthen confidence that the central claim survives more pessimistic device behavior.
  3. Evaluation (§6.2–6.3) is limited to two fixed-tile matmul kernels (Qwen fused_moe 128×256 BF16; Llama FFN 64×128) under three batch sizes. TileLens-SW/HW are claimed to cover FlashAttention, CUTLASS, cuBLAS, and DeepGEMM, but those kernels are not measured, nor are non-weight tensors if they were ever placed on LGMS. At minimum, the paper should either evaluate one additional TMA/DSL kernel class or clearly bound the claim to weight-side tiled GEMM under the stated placement policy (§5.1).

Circularity Check

0 steps flagged

No circularity: performance claims rest on independent cycle-level simulation of three layouts, not on quantities defined or fitted from the target results.

full rationale

The paper's load-bearing chain is (1) empirical quantification of read amplification under row-/column-major on LGMS (Figs. 1, 4–6, §2.5), (2) definition of tile-major via inter/intra-tile index decomposition (Eqs. 1–2, §3) that eliminates amplification by construction when memory-tile dimensions divide compute-tile dimensions, (3) transparent remapping via DSL layout descriptors and TMA extensions that reuse existing address-generation hardware under stated invariants (§4), and (4) end-to-end Macsim evaluation of SASS traces under the three layouts plus an analytic adaptive prefetcher (Eq. 6). None of these steps reduces a claimed prediction to a fitted input or to a self-citation that itself asserts the result. The prefetcher degree is derived from bandwidth × latency / wave size (with a fixed 2× plane-collision factor), not tuned to the reported slowdowns. Self-citations appear only as background on HBF/RoMe; the central “within 1 % of HBM” numbers are produced by running independent kernel traces. The derivation is therefore self-contained against the simulator’s own baselines.

Axiom & Free-Parameter Ledger

3 free parameters · 5 axioms · 3 invented entities

The central performance claim rests on standard GPU tiling practice, published HBF/RoMe parameters, and a small set of modeling choices (tile size, prefetcher formula, mixed-granularity L2). No new physical constants are fitted; the free parameters are engineering knobs whose sensitivity is reported.

free parameters (3)
  • memory_tile_shape = (64,32) BF16
    Chosen as (64,32) for BF16 / (64,64) for FP8 so that a×b×s = 4 KB and the tile divides common compute tiles; sensitivity shown in Fig. 19.
  • prefetcher_degree_fudge = 2×
    Eq. 6 multiplies the analytic degree by 2× “to compensate for NAND plane collisions”; empirical constant not derived from first principles.
  • HBF_NAND_latency = 5 µs (main result)
    Swept 1–20 µs; main claim uses 5 µs as the representative operating point taken from projected flash literature.
axioms (5)
  • domain assumption GPU matmul kernels partition weight matrices into 2-D compute tiles of 32–512 elements per side and load them via TMA or equivalent.
    Stated in §2.3 and used throughout; standard practice for cuBLAS/CUTLASS.
  • domain assumption HBF (and RoMe) expose a minimum 4 KB access granularity while offering bandwidth comparable to HBM3e.
    Taken from cited HBF/RoMe papers (§2.2); load-bearing for the amplification numbers.
  • domain assumption Weight matrices are static during inference and can be written once in tile-major layout at load time; activations and KV cache stay on HBM.
    §5.1 allocation policy; required for endurance and for the layout to be fixed at allocation.
  • domain assumption Matrix dimensions are multiples of the memory-tile dimensions (or can be padded); base addresses are aligned to 4 KB.
    §3.3; authors checked recent LLMs are multiples of 256.
  • domain assumption TMA tile contiguous dimension is a power of two (for swizzling), enabling the bit-permutation correction of Case 2.
    §4.3 invariants; required for the low-overhead HW extension.
invented entities (3)
  • tile-major layout / memory tile no independent evidence
    purpose: Reshape each contiguous 4 KB region into an a×b rectangle that fits inside a compute tile, eliminating read amplification.
    Core proposed data layout; no prior GPU global-memory use of this exact construct for LGMS.
  • TileLens-SW no independent evidence
    purpose: Extend CuTe/DSL layout descriptors so kernels adopt tile-major by changing only the descriptor.
    Software path for DSL-based kernels (CUTLASS, FlashAttention).
  • TileLens-HW no independent evidence
    purpose: Extend TMA descriptor and address-generation logic for transparent tile-major remapping without kernel changes.
    Hardware path for closed-source TMA kernels (cuBLAS, DeepGEMM); estimated 3–4 k gates.

pith-pipeline@v1.1.0-grok45 · 32242 in / 3317 out tokens · 29630 ms · 2026-07-11T22:10:23.196744+00:00 · methodology

0 comments
read the original abstract

Large Language Model (LLM) inference is bottlenecked by the capacity and bandwidth of GPU High-Bandwidth Memory (HBM). Recent proposals, such as High-Bandwidth Flash (HBF) and RoMe, offer higher capacity or bandwidth than HBM, but require a minimum access granularity of kilobytes. We show that these Large-Granularity Memory Systems (LGMS) can degrade the performance of tiled matrix-multiplication, which is the dominant operation in LLM inference, by up to an order of magnitude. The root cause of the slowdown is read amplification, where memory requests fetch far more data than the tile actually needs. This waste stems from a fundamental mismatch between the two-dimensional nature of compute tiles and the one-dimensional memory layout, leading to each request spilling well beyond the tile boundaries. To mitigate read amplification, we propose to use tile-major layout for LGMS. Rather than storing data as a one-dimensional strip, tile-major layout reshapes each contiguous memory block into a two-dimensional rectangle, aligning memory granularity with tile boundaries. To ease the adoption of tile-major layout on GPUs, we propose TileLens, lightweight software and hardware extensions that collectively cover major classes of GPU kernels. TileLens-SW extends GPU DSLs so that DSL-based kernels can adopt tile-major in global memory by changing only the layout descriptor. TileLens-HW extends the Tensor Memory Accelerator (TMA) for transparent tile-major support in TMA-based kernels without code changes. We evaluate TileLens on a cycle-level simulator using matrix-multiplication kernels from Qwen-3 30B and Llama-3.1 70B. Combining a tile-major layout with an adaptive hardware prefetcher, TileLens achieves near-HBM performance on HBF-augmented GPUs with a 5us HBF NAND read latency, reducing the geomean slowdown from 1.61-6.49x with conventional layouts to within 1% of an HBM-only baseline.

Figures

Figures reproduced from arXiv: 2607.04031 by Anish Saxena, Euijun Chung, Hritvik Taneja, Hyesoon Kim, Jae Hyung Ju, Moinuddin K. Qureshi, Shinnung Jeong.

Figure 1
Figure 1. Figure 1: (a) HBF and RoMe both require a minimum access granularity of 4 KB, 128 [PITH_FULL_IMAGE:figures/full_fig_p002_1.png] view at source ↗
Figure 2
Figure 2. Figure 2: Tiled matrix multiplication on GPUs. Each CTA [PITH_FULL_IMAGE:figures/full_fig_p003_2.png] view at source ↗
Figure 3
Figure 3. Figure 3: Address offsets of an 8 × 8 matrix under (a) column￾major and (b) row-major layout. 2.5 Challenges of Employing LGMS We show that LGMS causes severe read amplification in tiled ma￾trix multiplication. Because the weight matrix is stored in a one￾dimensional layout (row- or column-major), each 4 KB access re￾turns a contiguous strip along a single dimension, most of which falls outside the two-dimensional c… view at source ↗
Figure 4
Figure 4. Figure 4: (a) Log-scale distribution of HBF read latencies [PITH_FULL_IMAGE:figures/full_fig_p004_4.png] view at source ↗
Figure 5
Figure 5. Figure 5: CTA execution timelines under baseline HBM (top) [PITH_FULL_IMAGE:figures/full_fig_p004_5.png] view at source ↗
Figure 7
Figure 7. Figure 7: (a) Address offsets of an 8 × 8 matrix under tile-major layout with 4 × 2 memory tile. (b) Memory tile shapes for a 4 KB region relative to a compute tile (dashed). Memory tiles beyond the compute tile (4096 × 1) cause read amplification; shapes that fit within (e.g., 64 × 64) eliminate it. (c) Tile-major address translation: element (𝑖, 𝑗) is decomposed into an inter-tile index (𝑇𝑖 ,𝑇𝑗) and an intra-tile … view at source ↗
Figure 8
Figure 8. Figure 8: TMA descriptor and TMA load operation [PITH_FULL_IMAGE:figures/full_fig_p006_8.png] view at source ↗
Figure 9
Figure 9. Figure 9: A single 384×128 column-major matrix and the 128×64 tile at logical coordinates (256, 64), viewed as equiva￾lent (a) 2-D, (b) 3-D, and (c) 4-D TMA tensor descriptors. Moreover, except for the contiguous first dimension, the pro￾grammer is free to reorder the remaining dimensions, along with their strides, tile sizes, and coordinates. Kernels routinely use the different views to simplify coordinate computat… view at source ↗
Figure 12
Figure 12. Figure 12: Two cases of tile-major with memory tile of (a, b) [PITH_FULL_IMAGE:figures/full_fig_p007_12.png] view at source ↗
Figure 11
Figure 11. Figure 11: (a) Computation of the tile base address from co [PITH_FULL_IMAGE:figures/full_fig_p007_11.png] view at source ↗
Figure 13
Figure 13. Figure 13: GPU System configurations for evaluation. [PITH_FULL_IMAGE:figures/full_fig_p009_13.png] view at source ↗
Figure 14
Figure 14. Figure 14: Normalized kernel execution time for Qwen MoE and Llama FFN kernels under three global memory layouts. HBF [PITH_FULL_IMAGE:figures/full_fig_p010_14.png] view at source ↗
Figure 16
Figure 16. Figure 16: Average compute tile arrival latency breakdown [PITH_FULL_IMAGE:figures/full_fig_p010_16.png] view at source ↗
Figure 17
Figure 17. Figure 17: Per-request latency distribution (top) and CDF [PITH_FULL_IMAGE:figures/full_fig_p010_17.png] view at source ↗
Figure 19
Figure 19. Figure 19: Normalized kernel execution time for tile-major [PITH_FULL_IMAGE:figures/full_fig_p011_19.png] view at source ↗
Figure 18
Figure 18. Figure 18: Kernel execution time sensitivity to HBF NAND [PITH_FULL_IMAGE:figures/full_fig_p011_18.png] view at source ↗
Figure 20
Figure 20. Figure 20: RoMe evaluation results: execution time and effec [PITH_FULL_IMAGE:figures/full_fig_p011_20.png] view at source ↗

discussion (0)

Sign in with ORCID, Apple, or X to comment. Anyone can read and Pith papers without signing in.

Reference graph

Works this paper leans on

53 extracted references · 11 linked inside Pith

  1. [1]

    Josh Achiam, Steven Adler, Sandhini Agarwal, Lama Ahmad, Ilge Akkaya, Floren- cia Leoni Aleman, Diogo Almeida, Janko Altenschmidt, Sam Altman, Shyamal Anadkat, et al. 2023. Gpt-4 technical report.arXiv preprint arXiv:2303.08774 (2023)

    Pith/arXiv arXiv 2023

  2. [2]

    Sandhini Agarwal, Lama Ahmad, Jason Ai, Sam Altman, Andy Applebaum, Edwin Arbus, Rahul K Arora, Yu Bai, Bowen Baker, Haiming Bao, et al. 2025. gpt-oss-120b & gpt-oss-20b model card.arXiv preprint arXiv:2508.10925(2025)

    Pith/arXiv arXiv 2025

  3. [3]

    Berkin Akin, Franz Franchetti, and James C Hoe. 2015. Data reorganization in memory using 3D-stacked DRAM.ACM SIGARCH Computer Architecture News 43, 3S (2015), 131–143

    2015

  4. [4]

    AMD. 2026. Versal Adaptive SoC AIE-ML Architecture Manual (AM020). https: //docs.amd.com/r/en-US/am020-versal-aie-ml

    2026

  5. [5]

    John Carter, Wilson Hsieh, Leigh Stoller, Mark Swanson, Lixin Zhang, Erik Brun- vand, Al Davis, Chen-Chi Kuo, Ravindra Kuramkote, Michael Parker, et al. 1999. Impulse: Building a smarter memory controller. InProceedings Fifth International Symposium on High-Performance Computer Architecture. IEEE, 70–79

    1999

  6. [6]

    Tri Dao. 2023. Flashattention-2: Faster attention with better parallelism and work partitioning.arXiv preprint arXiv:2307.08691(2023)

    Pith/arXiv arXiv 2023

  7. [7]

    DeepSeek-AI. 2025. DeepGEMM: Clean and Efficient FP8 GEMM Kernels with Fine-Grained Scaling. https://github.com/deepseek-ai/DeepGEMM

    2025

  8. [8]

    Cagdas Dirik and Bruce Jacob. 2009. The performance of PC solid-state disks (SSDs) as a function of bandwidth, concurrency, device architecture, and system organization.ACM SIGARCH Computer Architecture News37, 3 (2009), 279–289

    2009

  9. [9]

    Aaron Grattafiori, Abhimanyu Dubey, Abhinav Jauhri, Abhinav Pandey, Abhishek Kadian, Ahmad Al-Dahle, Aiesha Letman, Akhil Mathur, Alan Schelten, Alex Vaughan, et al. 2024. The llama 3 herd of models.arXiv preprint arXiv:2407.21783 (2024)

    Pith/arXiv arXiv 2024

  10. [10]

    Daya Guo, Dejian Yang, Haowei Zhang, Junxiao Song, Peiyi Wang, Qihao Zhu, Runxin Xu, Ruoyu Zhang, Shirong Ma, Xiao Bi, et al. 2025. Deepseek-r1: Incen- tivizing reasoning capability in llms via reinforcement learning.arXiv preprint arXiv:2501.12948(2025)

    Pith/arXiv arXiv 2025

  11. [11]

    Minho Ha, Euiseok Kim, and Hoshik Kim. 2026. H 3: Hybrid Architecture Using High Bandwidth Memory and High Bandwidth Flash for Cost-Efficient LLM Inference.IEEE Computer Architecture Letters(2026)

    2026

  12. [12]

    Po-Kai Hsu, Weihong Xu, Qunyou Liu, Tajana Rosing, and Shimeng Yu. 2026. HAVEN: High-Bandwidth Flash Augmented Vector Engine for Large-Scale Ap- proximate Nearest-Neighbor Search Acceleration.arXiv preprint arXiv:2603.01175 (2026)

    arXiv 2026

  13. [13]

    Intel. 2025. Optimizing Explicit SIMD Kernels. https://www.intel.com/content/ www/us/en/docs/oneapi/optimization-guide-gpu/2025-2/optimizing-explicit- simd-kernels.html

    2025

  14. [14]

    Adwait Jog, Onur Kayiran, Asit K Mishra, Mahmut T Kandemir, Onur Mutlu, Ravishankar Iyer, and Chita R Das. 2013. Orchestrated scheduling and prefetch- ing for GPGPUs. InProceedings of the 40th Annual International Symposium on Computer Architecture. 332–343

    2013

  15. [15]

    Hyojong Kim, Ramyad Hadidi, Lifeng Nai, Hyesoon Kim, Nuwan Jayasena, Ya- suko Eckert, Onur Kayiran, and Gabriel Loh. 2018. Coda: Enabling co-location of computation and data for multiple gpu systems.ACM Transactions on Architecture and Code Optimization (TACO)15, 3 (2018), 1–23

    2018

  16. [16]

    Hyesoon Kim, Jaekyu Lee, Nagesh B Lakshminarayana, Jaewoong Sim, Jieun Lim, and Tri Pho. 2012. Macsim: A cpu-gpu heterogeneous simulation framework user guide.Georgia Institute of Technology(2012), 1–57

    2012

  17. [17]

    Gunjae Koo, Hyeran Jeon, Zhenhong Liu, Nam Sung Kim, and Murali Annavaram

  18. [18]

    In2018 IEEE International 12 Parallel and Distributed Processing Symposium (IPDPS)

    Cta-aware prefetching and scheduling for gpu. In2018 IEEE International 12 Parallel and Distributed Processing Symposium (IPDPS). IEEE, 137–148

  19. [19]

    Gunjae Koo, Kiran Kumar Matam, Te I, HV Krishna Giri Narra, Jing Li, Hung-Wei Tseng, Steven Swanson, and Murali Annavaram. 2017. Summarizer: trading communication with computing near storage. InProceedings of the 50th Annual IEEE/ACM International Symposium on Microarchitecture. 219–231

    2017

  20. [20]

    Kwanhee Kyung, Yeon Ji Moon, Juhwan Cho, and Jung Ho Ahn. 2026. High- Bandwidth Flash for KV Caches: Endurance and Performance Implications.IEEE Computer Architecture Letters(2026)

    2026

  21. [21]

    Jaekyu Lee, Nagesh B Lakshminarayana, Hyesoon Kim, and Richard Vuduc

  22. [22]

    In 2010 43rd Annual IEEE/ACM International Symposium on Microarchitecture

    Many-thread aware prefetching mechanisms for GPGPU applications. In 2010 43rd Annual IEEE/ACM International Symposium on Microarchitecture. IEEE, 213–224

    2010

  23. [23]

    Aonian Li, Bangwei Gong, Bo Yang, Boji Shan, Chang Liu, Cheng Zhu, Chunhao Zhang, Congchao Guo, Da Chen, Dong Li, et al . 2025. Minimax-01: Scaling foundation models with lightning attention.arXiv preprint arXiv:2501.08313 (2025)

    Pith/arXiv arXiv 2025

  24. [24]

    Aixin Liu, Bei Feng, Bing Xue, Bingxuan Wang, Bochao Wu, Chengda Lu, Cheng- gang Zhao, Chengqi Deng, Chenyu Zhang, Chong Ruan, et al. 2024. Deepseek-v3 technical report.arXiv preprint arXiv:2412.19437(2024)

    Pith/arXiv arXiv 2024

  25. [25]

    Weile Luo, Ruibo Fan, Zeyu Li, Dayou Du, Hongyuan Liu, Qiang Wang, and Xiaowen Chu. 2025. Dissecting the nvidia hopper architecture through mi- crobenchmarking and multiple level analysis.arXiv preprint arXiv:2501.12084 (2025)

    Pith/arXiv arXiv 2025

  26. [26]

    Xiaoyu Ma and David Patterson. 2026. Challenges and Research Directions for Large Language Model Inference Hardware.arXiv preprint arXiv:2601.05047 (2026)

    arXiv 2026

  27. [27]

    Hwayong Nam, Seungmin Baek, Jumin Kim, Michael Jaemin Kim, and Jung Ho Ahn. 2026. RoMe: Row Granularity Access Memory System for Large Language Models. In2026 IEEE International Symposium on High Performance Computer Architecture (HPCA)

    2026

  28. [28]

    NVIDIA. 2023. NVIDIA H200 GPU Datasheet. https://resources.nvidia.com/en- us-gpu-resources/hpc-datasheet-sc23

    2023

  29. [29]

    NVIDIA. 2024. NVIDIA H100 Tensor Core GPU Architecture. https://resources. nvidia.com/en-us-hopper-architecture/nvidia-h100-tensor-c

    2024

  30. [30]

    NVIDIA. 2026. CUDA Toolkit Documentation 13.2. https://docs.nvidia.com/ cuda/

    2026

  31. [31]

    NVIDIA. 2026. cuTensorMapEncodeTiled — CUDA Driver API Refer- ence. https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR_ _MEMORY.html

    2026

  32. [32]

    NVIDIA. 2026. NVIDIA cuBLAS Documentation. https://docs.nvidia.com/cuda/ cublas/

    2026

  33. [33]

    NVIDIA. 2026. NVIDIA CUTLASS Documentation. https://docs.nvidia.com/ cutlass/latest/

    2026

  34. [34]

    NVIDIA. 2026. Parallel Thread Execution ISA Version 9.2. https://docs.nvidia. com/cuda/parallel-thread-execution/

    2026

  35. [35]

    Junho Park, Hyowon An, Haeseok Suh, Youngsu Yoon, Hyuni Lee, and Joungho Kim. 2026. HBM-HBF-Centric Memory Pooling Architecture With Custom Base Die for Terabyte-Scale LLM Inference.IEEE Computer Architecture Letters(2026)

    2026

  36. [36]

    Adam Paszke, Sam Gross, Soumith Chintala, Gregory Chanan, Edward Yang, Zachary DeVito, Zeming Lin, Alban Desmaison, Luca Antiga, and Adam Lerer

  37. [37]

    InNIPS-W

    Automatic differentiation in PyTorch. InNIPS-W

  38. [38]

    Zaid Qureshi, Vikram Sharma Mailthody, Isaac Gelado, Seungwon Min, Amna Masood, Jeongmin Park, Jinjun Xiong, Chris J Newburn, Dmitri Vainbrand, I-Hsin Chung, et al. 2023. GPU-initiated on-demand high-throughput storage access in the BaM system architecture. InProceedings of the 28th ACM International Conference on Architectural Support for Programming Lan...

    2023

  39. [39]

    Minsoo Rhu, Michael Sullivan, Jingwen Leng, and Mattan Erez. 2013. A locality- aware memory hierarchy for energy-efficient GPU architectures. InProceedings of the 46th Annual IEEE/ACM International Symposium on Microarchitecture. 86–98

    2013

  40. [40]

    SANDISK. 2025. SANDISK unveils the future of memory architecture for AI, introducing: High Bandwidth Flash. https://documents.sandisk.com/content/ dam/asset-library/en_us/assets/public/sandisk/collateral/company/Sandisk- HBF-Fact-Sheet.pdf

    2025

  41. [41]

    SemiAnalysis. 2026. Blackwell Micro-Architecture Benchmarking. https://github. com/SemiAnalysisAI/microbench-blackwell

    2026

  42. [42]

    Vivek Seshadri, Thomas Mullins, Amirali Boroumand, Onur Mutlu, Phillip B Gibbons, Michael A Kozuch, and Todd C Mowry. 2015. Gather-scatter DRAM: In-DRAM address translation to improve the spatial locality of non-unit strided accesses. InProceedings of the 48th International Symposium on Microarchitecture. 267–280

    2015

  43. [43]

    Changmin Shin, Jaeyong Song, Hongsun Jang, Dogeun Kim, Jun Sung, Taehee Kwon, Jae Hyung Ju, Frank Liu, Yeonkyu Choi, and Jinho Lee. 2025. Piccolo: Large-scale graph processing with fine-grained in-memory scatter-gather. In 2025 IEEE International Symposium on High Performance Computer Architecture (HPCA). IEEE, 641–656

    2025

  44. [44]

    Mohammad Shoeybi, Mostofa Patwary, Raul Puri, Patrick LeGresley, Jared Casper, and Bryan Catanzaro. 2019. Megatron-lm: Training multi-billion parameter language models using model parallelism.arXiv preprint arXiv:1909.08053(2019)

    Pith/arXiv arXiv 2019

  45. [45]

    Dowon Son, Yonggon Park, Hyunuk Cho, Hyungkyu Ham, Onur Mutlu, Sungjin Lee, Gwangsun Kim, and Jisung Park. 2026. Exploring High-Bandwidth Flash for Modern LLM Inference: Opportunities and Challenges.IEEE Computer Architec- ture Letters(2026)

    2026

  46. [46]

    Benjamin F Spector, Simran Arora, Aaryan Singhal, Daniel Y Fu, and Christopher Ré. 2024. Thunderkittens: Simple, fast, and adorable ai kernels.arXiv preprint arXiv:2410.20399(2024)

    Pith/arXiv arXiv 2024

  47. [47]

    I-Jui Sung, Geng Daniel Liu, and Wen-Mei W Hwu. 2012. DL: A data layout transformation system for heterogeneous computing. In2012 Innovative Parallel Computing (InPar). IEEE, 1–11

    2012

  48. [48]

    Texas Instruments. 2012. TMS320C6472/TMS320TCI648x DSP Enhanced DMA (EDMA3) Controller User’s Guide. https://www.ti.com/lit/ug/spru727e/spru727e. pdf

    2012

  49. [49]

    Philippe Tillet, Hsiang-Tsung Kung, and David Cox. 2019. Triton: an intermediate language and compiler for tiled neural network computations. InProceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Program- ming Languages. 10–19

    2019

  50. [50]

    Oreste Villa, Mark Stephenson, David Nellans, and Stephen W Keckler. 2019. Nvbit: A dynamic binary instrumentation framework for nvidia gpus. InProceed- ings of the 52nd Annual IEEE/ACM International Symposium on Microarchitecture. 372–383

    2019

  51. [51]

    An Yang, Anfeng Li, Baosong Yang, Beichen Zhang, Binyuan Hui, Bo Zheng, Bowen Yu, Chang Gao, Chengen Huang, Chenxu Lv, et al. 2025. Qwen3 technical report.arXiv preprint arXiv:2505.09388(2025)

    Pith/arXiv arXiv 2025

  52. [52]

    Jie Zhang and Myoungsoo Jung. 2020. ZnG: Architecting GPU multi-processors with new flash for scalable data analysis. In2020 ACM/IEEE 47th Annual Interna- tional Symposium on Computer Architecture (ISCA). IEEE, 1064–1075

    2020

  53. [53]

    Weiyu Zhou, Zheng Wang, Chao Chen, Yike Li, Yongkui Yang, Zhuoyu Wu, and Anupam Chattopadhyay. 2025. Tensor Manipulation Unit (TMU): Reconfig- urable, Near-Memory Tensor Manipulation for High-Throughput AI SoC.IEEE Transactions on Very Large Scale Integration (VLSI) Systems(2025). 13

    2025