Pith. sign in

REVIEW 3 major objections 5 minor 75 references

AGILE: Lightweight and Efficient Asynchronous GPU-SSD Integration

T0 review · 3 major / 5 minor · reviewed 2026-08-16 · deepseek-v4-flash

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

Pith's one-line read A lightweight GPU-side service kernel lets GPU threads issue NVMe requests asynchronously without deadlock, overlapping I/O with computation and achieving up to 1.88x end-to-end speedup over synchronous GPU-centric I/O.

desk verdict AGILE is a genuine first in GPU-centric async I/O with a sensible design and supportive experiments; the deadlock-free claim rests on an unmeasured service-kernel liveness assumption that the authors should scope or verify. read the letter →

arxiv 2504.19365 v3 pith:HXYGBX7X submitted 2025-04-27 cs.DC

Zhuoping Yang , Jinming Zhuang , Xingzhen Chen , Alex K. Jones , Peipei Zhou This is my paper · ORCID
classification cs.DC
keywords GPUNVMeasynchronousI/OsoftwarecachedeadlockavoidanceHBMstoragesystems
verification ladder T0 review T1 audit T2 compute T3 formal

The pith

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

The reading

AGILE is a GPU-centric I/O library that lets GPU threads issue NVMe storage requests asynchronously, rather than making each thread stall until the SSD replies. The paper's central claim is that a lightweight service kernel running on the GPU can own completion-queue polling and release submission-queue entry locks, eliminating the deadlock that otherwise blocks asynchronous GPU-side I/O. On a microbenchmark spanning computation-to-communication ratios, asynchronous AGILE reaches up to 1.88x lower end-to-end time than a synchronous counterpart, and on DLRM inference it reports up to 1.75x speedup over the state-of-the-art synchronous GPU-centric baseline while cutting per-thread register usage by up to 1.32x. The library also provides a flexible HBM software cache with user-customizable policies and a coherence mechanism for user-specified buffers.

What carries the argument

The load-bearing mechanism is a lightweight kernel daemon on the GPU that owns completion-queue polling: a warp checks 32 physically contiguous CQEs per round, tracks phase bits, advances the CQ head and doorbell, and uses the command identifier to release the correct submission-queue entry locks when completions arrive. This breaks the deadlock cycle in which a submission queue fills before any thread can consume a completion. Supporting machinery includes Algorithm 2's serialized SQ tail/doorbell update with per-SQE lock states (EMPTY, UPDATED, ISSUED), a two-level warp-plus-cache coalescing step for duplicate requests, a software-managed HBM cache with user-pluggable replacement policies and a four-state line protocol (INVALID, BUSY, READY, MODIFIED), and a Share Table that extends MOESI-style coherence to user-specified buffers via pointer sharing and reference counts.

What would settle it

Run the prefetch workload with a single NVMe queue pair of depth 64 while a user kernel occupies nearly all streaming multiprocessors, so the AGILE service's polling warps compete for scheduling; if the submission queue fills and end-to-end time approaches the synchronous baseline, the dependence on service progress is demonstrated. If throughput stays flat, the deadlock-avoidance design would hold even when the service is starved.

Watch

Extended reading notes

Core claim

AGILE establishes that a GPU-centric asynchronous I/O model is practical: a lightweight service kernel resident on the GPU continuously polls all registered NVMe completion queues in a warp-centric round-robin fashion, matches each completion to its submission entry via the NVMe Command Identifier, and releases the entry locks once the SSD finishes. User threads never retain a lock after enqueueing a command; they hand the lock to the service and receive a barrier, so even a completely full submission queue cannot deadlock the issuers. On top of this, AGILE serializes submission-queue doorbell updates by scanning SQE states (EMPTY, UPDATED, ISSUED), coalesces duplicate warp-level requests, and layers an HBM software cache with pluggable policies plus a Share Table that applies a MOESI-inspired coherence protocol to user buffers. The evaluation claims up to 1.88x speedup over a synchronous model, up to 1.75x over the synchronous GPU-centric baseline on DLRM, up to 3.12x and 2.85x reductions in software-cache and NVMe-I/O API overhead on graph workloads, and up to 1.32x lower per-thread register usage.

Load-bearing premise

The design rests on the AGILE service kernel being scheduled on the GPU often enough to drain NVMe completion queues faster than user threads can fill the submission queues; if user kernels crowd out or delay the service, queues fill, threads stall, and the asynchronous advantage collapses.

Editorial extensions

If this is right

  • GPU threads can issue NVMe reads and writes without host-CPU involvement, so the CPU no longer needs to synchronize with every in-flight GPU request.
  • Asynchronous prefetching hides SSD latency behind compute only when the software cache is large enough to avoid evicting not-yet-needed data; the paper's cache-size sweep shows async mode overtakes sync mode once the cache reaches roughly 64 MB in its DLRM workload.
  • Users can supply custom cache replacement and sharing policies through the CRTP-based interface, letting the same library adapt to different access patterns without kernel recompilation.
  • The service kernel offloads polling logic from application kernels, which is the stated reason for up to 1.32x lower per-thread register usage and retained warp scheduling flexibility.
  • Aggregate bandwidth scales with the number of SSDs: around 3.7, 7.4, and 11.1 GB/s for 4KB random reads with one, two, and three SSDs respectively.

Reading between the lines

Editorial extensions of the paper, not claims the author makes directly.

  • I infer the lock-handoff pattern could generalize to any producer-consumer queue on the GPU, such as network queues or CXL-attached memory controllers, whenever a dedicated progress engine is needed to prevent lock-holding deadlocks.
  • Because the paper notes programmers currently find overlap opportunities by hand, a natural extension is a compiler pass that automatically inserts epoch-ahead prefetches and uses AGILE's barrier returns for synchronization.
  • The small-cache degradation suggests a tunable cache-size-versus-prefetch-distance policy is worth building: the API could expose the cache capacity or prefetch distance as an autotuned knob to keep async mode above its operating threshold.
  • The Share Table's pointer-based MOESI adaptation appears to be a ready building block for multi-GPU coherent caching over NVLink, which the paper lists as future work rather than a demonstrated result.
Share X Bluesky LinkedIn Reddit HN

Signed reviews

No signed human review yet.

Editorial analysis

A structured set of objections, weighed in public.

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

Referee Report

3 major / 5 minor

Summary. The paper presents AGILE, a GPU-centric I/O library that lets GPU threads issue NVMe requests asynchronously. The central mechanism is a lightweight persistent GPU service kernel that polls NVMe completion queues, releases submission-queue locks and barriers, and is started before and stopped after user kernels (Section 3.2, Listing 1). The design also includes a software-managed HBM cache with pluggable policies, warp-level request coalescing, and a Share Table for coherent user buffers. The evaluation compares AGILE with a synchronous I/O model and with BaM on microbenchmarks, DLRM inference, BFS/SpMV, and register counts, claiming up to 1.88x speedup over synchronous I/O, up to 1.75x over BaM on DLRM, 3.12x lower software-cache API overhead, 2.85x lower I/O API overhead, and 1.32x lower per-thread register usage.

Significance. If the results hold, AGILE is a useful contribution: it addresses a real problem (GPU threads stalling on synchronous NVMe accesses), and the service-kernel/lock-handoff design is a plausible alternative to CPU-mediated asynchronous I/O. The paper is careful to present the ideal speedup as a reference bound rather than a fitted model, and the authors open-source the implementation. The main claims, however, rest on empirical speedups and on an unverified liveness property of the service kernel, so significance is conditional on closing those gaps.

major comments (3)
  1. [Section 3.2, Listing 1, Section 4.6] The deadlock-elimination and overlap claims rely on the AGILE service kernel making progress while the user kernel occupies the GPU, but the paper does not establish this. CUDA does not guarantee forward progress for co-resident kernels, and the manuscript never measures the service kernel's achieved CQ polling rate, scheduling latency, or occupancy under full user-kernel load; Section 4.6 only reports 37 registers per thread, which is not a liveness guarantee. Figures 9 and 10 show that the asynchronous advantage collapses to synchronous or worse behavior with one queue pair or a small software cache, confirming that progress is conditional. To support the abstract's claim that AGILE 'eliminates deadlock risks,' the authors should either provide a forward-progress argument (for example, guaranteed service occupancy or preemption-resilient polling) or measure service progress under the evaluated user kernels with varying occupancy.
  2. [Sections 4.1-4.5] All speedup figures are normalized ratios with no repetition counts, no error bars, and no absolute execution times. The headline claims (1.88x over synchronous I/O, 1.75x over BaM) are therefore not statistically grounded; a single measurement could support any of the reported ratios. Please add the number of runs, variance or error bars, and absolute execution times (or a table) for at least the headline comparisons.
  3. [Sections 3.4 and 3.5] The abstract claims AGILE 'eliminates deadlock risks,' but the software-cache path can still block indefinitely: case (d) in Section 3.4 states that a BUSY cache line cannot be evicted until processing finishes and that the user-specified cache policy may decide to wait, and the debug lock-chain in Section 3.5 only detects circular dependencies after the fact. No mechanism is described that prevents a cache-policy-induced deadlock. The paper should either restrict the deadlock-elimination claim to NVMe-queue deadlocks or provide a deadlock-avoidance protocol for the software cache.
minor comments (5)
  1. [Abstract and Section 4.5] The abstract and conclusion report a 3.12x cache-overhead reduction, while Section 4.5 reports 3.17x for SpMV on Kronecker graphs; harmonize these numbers.
  2. [Figure 4] The x-axis labels and legend in Figure 4 are garbled and hard to read; the figure should be regenerated with clear axis labels.
  3. [Section 3.1 and Listing 1] There is a typo in Section 3.1 ('SDDs' should be 'SSDs'), and Listing 1 uses 'initNvme' while the surrounding text uses 'NVMe'; please make the capitalization consistent.
  4. [Section 4.4] The sentence 'The results also indicate that the AGILE async benefits more when the batch size is smaller and near 16' should be rephrased; the intended claim about the location and reason for the peak at batch size 16 needs a more precise statement.
  5. [Section 4.6] The text says the AGILE service kernel 'can assist multiple CUDA kernels simultaneously,' but no experiment demonstrates concurrent user kernels; either add such a test or soften the claim.

Circularity Check

0 steps flagged · score 0.0 of 10

No significant circularity: AGILE's performance claims are empirical comparisons and an analytic reference bound, with no fitted parameter or self-citation chain doing load-bearing work.

full rationale

AGILE's central claims (async overlap, deadlock avoidance, cache flexibility, speedups over BaM) are supported by system design and measured experiments, not by equations that reduce to their inputs. The ideal-speedup formula (Eq. 1) is an analytic bound used as a reference: 'Ideally, when computation and communication perfectly overlap with each other, the speedup can be defined by Equation 1'; the experimental speedups are measured against a synchronous baseline and against BaM, so the 1.88x, 1.75x, 3.12x, and 2.85x numbers are empirical outcomes, not constructed identities. No parameter is fitted to a subset of data and then reported as a prediction. The only author self-citations (e.g., refs. 16, 23, 24, 64, 67-70) appear in background material on FPGAs and future-work discussion and are not invoked to justify the asynchronous I/O design or the deadlock-avoidance mechanism. The skeptical concern that the AGILE service kernel needs a forward-progress guarantee to realize the claimed overlap is a legitimate correctness/robustness risk, but the paper does not derive the speedup from that guarantee; it measures performance under the tested configurations. Therefore no circular derivation is present.

Assumptions & free parameters 0 free parameters · 5 assumptions · 2 invented entities

AGILE introduces software components (service, share table) rather than free parameters. The main axiomatic input is the correctness of the NVMe protocol and GPU scheduling model, both standard domain assumptions. The design does not fit any parameters to data.

assumptions (5)
  • domain assumption NVMe queue-pair and doorbell semantics are as described in Section 2.1.
    The design relies on standard NVMe protocol behavior (SQ/CQ, tail/head pointers, doorbells) being exactly as described; any deviation would break the issuing and completion logic.
  • domain assumption GPU thread-block scheduling and warp-level preemption behavior behave as described in Section 2.2.
    The motivation for asynchronous I/O assumes that stalled warps cause underutilization that can be improved by overlapping I/O.
  • domain assumption The AGILE service kernel can run concurrently on the GPU and poll CQs without causing resource exhaustion.
    Deadlock avoidance depends on the service making progress while user threads run. The paper shows 37 registers per thread and low warp usage, but does not formally prove schedulability.
  • domain assumption The MOESI-inspired coherency protocol applies to user-specified buffers.
    The Share Table adopts a modified MOESI model for software-managed buffers; correctness is argued informally in Section 3.4.1.
  • domain assumption NVMe SSDs expose only a limited number of I/O queue pairs (for example 128), so completion processing must be efficient.
    The warp-centric CQ polling design is motivated by this hardware property; on systems with many queue pairs the benefit may shrink.
invented entities (2)
  • AGILE service (GPU-side completion daemon) independent evidence
    purpose: Polls NVMe completion queues and releases SQ locks, preventing deadlock in asynchronous issue.
    The service is a software component shipped in the open-source repository; its effect is measurable through the reported speedups.
  • Share Table independent evidence
    purpose: Tracks user-specified buffer ownership to extend cache coherency to arbitrary buffers.
    Implemented in software and exercised via async_issue; the open-source code provides a falsifiable handle.

how reviews work

0 comments
Cite this review

Pith. "Pith review of AGILE: Lightweight and Efficient Asynchronous GPU-SSD Integration." pith.science (2026). https://pith.science/paper/HXYGBX7X

@misc{pith2026250419365,
  author       = {Pith},
  title        = {Pith review of: AGILE: Lightweight and Efficient Asynchronous GPU-SSD Integration},
  year         = {2026},
  howpublished = {\url{https://pith.science/paper/HXYGBX7X}},
  note         = {Machine review of arXiv:2504.19365}
}
abstract

GPUs are critical for compute-intensive applications, yet emerging workloads such as recommender systems, graph analytics, and data analytics often exceed GPU memory capacity. Existing solutions allow GPUs to use CPU DRAM or SSDs as external memory, and the GPU-centric approach enables GPU threads to directly issue NVMe requests, further avoiding CPU intervention. However, current GPU-centric approaches adopt synchronous I/O, forcing threads to stall during long communication delays. We propose AGILE, a lightweight asynchronous GPU-centric I/O library that eliminates deadlock risks and integrates a flexible HBM-based software cache. AGILE overlaps computation and I/O, improving performance by up to 1.88$\times$ across workloads with diverse computation-to-communication ratios. Compared to BaM on DLRM, AGILE achieves up to 1.75$\times$ speedup through efficient design and overlapping; on graph applications, AGILE reduces software cache overhead by up to 3.12$\times$ and NVMe I/O overhead by up to 2.85$\times$; AGILE also lowers per-thread register usage by up to 1.32$\times$.

Figures

Figures reproduced from arXiv: 2504.19365 by the authors.

Figure 1
Figure 1. A deadlock example caused by sharing NVMe [PITH_FULL_IMAGE:figures/full_fig_p003_1.png] view at source ↗
Figure 2
Figure 2. Overview of system architecture adopting AGILE. [PITH_FULL_IMAGE:figures/full_fig_p004_2.png] view at source ↗
Figure 3
Figure 3. Avoiding NVMe Queue Deadlocks in AGILE [PITH_FULL_IMAGE:figures/full_fig_p005_3.png] view at source ↗
Figures from the paper (8 more)
Figure 4
Figure 4. Figure 4: Speedup comparison of asynchronous I/O over syn [PITH_FULL_IMAGE:figures/full_fig_p008_4.png]
Figure 5
Figure 5. Figure 5: AGILE 4KB random read on multiple SSDs Bandwidth (GB/s) 8 6 4 2 0 #Request per SSD 1 8 64 512 4096 32768 262144 1 SSD 2 SSDs 3 SSDs [PITH_FULL_IMAGE:figures/full_fig_p008_5.png]
Figure 6
Figure 6. Figure 6: AGILE 4KB random write on multiple SSDs 4.3 AGILE 4KB random read and write on multiple SSDs We evaluate the scalability of AGILE using 4 KB random read and write using 1, 2, and 3 SSDs, as shown in [PITH_FULL_IMAGE:figures/full_fig_p008_6.png]
Figure 8
Figure 8. Figure 8: Speedup comparison of AGILE (async and sync [PITH_FULL_IMAGE:figures/full_fig_p009_8.png]
Figure 9
Figure 9. Figure 9: Speedup comparison of AGILE (async and sync [PITH_FULL_IMAGE:figures/full_fig_p009_9.png]
Figure 7
Figure 7. Figure 7: Speedup comparison of AGILE (async and sync [PITH_FULL_IMAGE:figures/full_fig_p009_7.png]
Figure 11
Figure 11. Figure 11: Execution time breakdown of BaM and AGILE [PITH_FULL_IMAGE:figures/full_fig_p010_11.png]
Figure 12
Figure 12. Figure 12: Per-thread register usage comparison between [PITH_FULL_IMAGE:figures/full_fig_p011_12.png]

Discussion (0). Continue with ORCID to comment.

Reference graph

Works this paper leans on

75 extracted references · 62 canonical work pages

  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)

    arXiv 2023

  2. [2]

    Tyler Allen and Rong Ge. 2021. In-depth analyses of unified virtual memory sys- tem for GPU accelerated computing. InProceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis . 1–15

    work page 2021

  3. [3]

    Jonghyun Bae, Jongsung Lee, Yunho Jin, Sam Son, Shine Kim, Hakbeom Jang, Tae Jun Ham, and Jae W Lee. 2021. {FlashNeuron}:{SSD-Enabled}{ Large- Batch} training of very deep neural networks. In 19th USENIX Conference on File and Storage Technologies (FAST 21). 387–401

    work page 2021

  4. [4]

    Shai Bergman, Tanya Brokhman, Tzachi Cohen, and Mark Silberstein. 2019. SPIN: Seamless operating system integration of peer-to-peer DMA between SSDs and GPUs. ACM Transactions on Computer Systems (TOCS) 36, 2 (2019), 1–26

    work page 2019

  5. [5]

    Xiao Bi, Deli Chen, Guanting Chen, Shanhuang Chen, Damai Dai, Chengqi Deng, Honghui Ding, Kai Dong, Qiushi Du, Zhe Fu, et al. 2024. Deepseek llm: Scaling open-source language models with longtermism. arXiv preprint arXiv:2401.02954 (2024)

    arXiv 2024

  6. [6]

    Aydin Buluç and Kamesh Madduri. 2011. Parallel breadth-first search on dis- tributed memory systems. In Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis . 1–12

    work page 2011

  7. [7]

    Chia-Hao Chang, Jihoon Han, Anand Sivasubramaniam, Vikram Sharma Mailthody, Zaid Qureshi, and Wen-Mei Hwu. 2024. Gmt: Gpu orchestrated memory tiering for the big data era. In Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 3 . 464–478

    work page 2024

  8. [8]

    Chang Chen, Xiuhong Li, Qianchao Zhu, Jiangfei Duan, Peng Sun, Xingcheng Zhang, and Chao Yang. 2024. Centauri: Enabling efficient scheduling for communication-computation overlap in large model training via communication partitioning. In Proceedings of the 29th ACM International Conference on Archi- tectural Support for Programming Languages and Operating...

    work page 2024

Show all 75 references
  1. [9]

    Avery Ching, Sergey Edunov, Maja Kabiljo, Dionysios Logothetis, and Sambavi Muthukrishnan. 2015. One trillion edges: Graph processing at facebook-scale. Proceedings of the VLDB Endowment 8, 12 (2015), 1804–1815

    2015

  2. [10]

    Fernando J Corbato. 1968. A paging experiment with the multics system . Mas- sachusetts Institute of Technology

    1968

  3. [11]

    Neal C Crago, Sana Damani, Karthikeyan Sankaralingam, and Stephen W Keckler

  4. [12]

    Criteo AI Lab. 2025. Download Criteo 1TB Click Logs dataset - Criteo AI Lab. https://ailab.criteo.com/download-criteo-1tb-click-logs-dataset/

    2025

  5. [13]

    John D Davis and Eric S Chung. 2012. SpMV: A memory-bound application on the GPU stuck between a rock and a hard place. Microsoft Research Silicon Valley, Technical Report14 September 2012 (2012)

    2012

  6. [14]

    Debendra Das Sharma and Ishwar Agarwal. 2023. CXL_3.0_white- paper_FINAL. https://computeexpresslink.org/wp-content/uploads/2023/12/ CXL_3.0_white-paper_FINAL.pdf

    2023

  7. [15]

    Dell. 2021. Dell Enterprise Agnostic NVMe Drive Technical Specifications. https://dl.dell.com/manuals/all-products/esuprt_data_center_infra_int/esuprt_ data_center_infra_storage_adapters/dell-poweredge-exp-fsh-nvme-pcie- ssd_Users-Guide7_en-us.pdf

    2021

  8. [16]

    Peiyan Dong, Jinming Zhuang, Zhuoping Yang, Shixin Ji, Yanyu Li, Dongkuan Xu, Heng Huang, Jingtong Hu, Alex K Jones, Yiyu Shi, et al. 2024. EQ-ViT: Algorithm- hardware co-design for end-to-end acceleration of real-time vision transformer inference on Versal ACAP architecture. ...

    2024

  9. [17]

    Gil Einziger, Roy Friedman, and Ben Manes. 2017. Tinylfu: A highly efficient cache admission policy. ACM Transactions on Storage (ToS) 13, 4 (2017), 1–31

    2017

  10. [18]

    Eran Gal and Sivan Toledo. 2005. Algorithms and data structures for flash memories. ACM Computing Surveys (CSUR) 37, 2 (2005), 138–163

    2005

  11. [19]

    Amir Gholami, Zhewei Yao, Sehoon Kim, Coleman Hooper, Michael W Mahoney, and Kurt Keutzer. 2024. AI and memory wall. IEEE Micro (2024)

    2024

  12. [20]

    Pieter Hijma, Stijn Heldens, Alessio Sclocco, Ben Van Werkhoven, and Henri E Bal. 2023. Optimization techniques for GPU programming. Comput. Surveys 55, 11 (2023), 1–81

    2023

  13. [21]

    Jeongmin Hong, Sungjun Cho, Geonwoo Park, Wonhyuk Yang, Young-Ho Gong, and Gwangsun Kim. 2024. Bandwidth-effective dram cache for gpu s with storage-class memory. In2024 IEEE International Symposium on High-Performance Computer Architecture (HPCA). IEEE, 139–155

    2024

  14. [22]

    Guyue Huang, Yang Bai, Liu Liu, Yuke Wang, Bei Yu, Yufei Ding, and Yuan Xie

  15. [23]

    Shixin Ji, Xingzhen Chen, Jinming Zhuang, Wei Zhang, Zhuoping Yang, Sarah Schultz, Yukai Song, Jingtong Hu, Alex Jones, Zheng Dong, and Peipei Zhou

  16. [24]

    Shixin Ji, Zhuoping Yang, Xingzhen Chen, Wei Zhang, Jinming Zhuang, Alex K Jones, Zheng Dong, and Peipei Zhou. 2025. CLARE: Deterministic Cycle-Level Accelerator on REconfigurable platforms in DNN-Enabled Real-Time Safety- Critical Systems. In The 46th IEEE Real-Time Systems S...

    2025

  17. [25]

    Diya Joseph, Juan Luis Aragón, Joan-Manuel Parcerisa, and Antonio Gonzalez

  18. [26]

    Mark J Kilgard and Jeff Bolz. 2012. GPU-accelerated path rendering. ACM Transactions on Graphics (TOG) 31, 6 (2012), 1–10

    2012

  19. [27]

    Gyusun Lee, Seokha Shin, Wonsuk Song, Tae Jun Ham, Jae W Lee, and Jinkyu Jeong. 2019. Asynchronous {I/O} stack: A low-latency kernel {I/O} stack for{Ultra-Low} latency{SSDs}. In 2019 USENIX Annual Technical Conference (USENIX ATC 19). 603–616

    2019

  20. [28]

    Ruihao Li, Sanjana Yadav, Qinzhe Wu, Krishna Kavi, Gayatri Mehta, Neeraja J Yadwadkar, and Lizy K John. 2023. Performance Implications of Async Mem- cpy and UVM: A Tale of Two Data Transfer Modes. In 2023 IEEE International Symposium on Workload Characterization (IISWC). IEEE, 115–127

    2023

  21. [29]

    arXiv preprint arXiv:2404.06156 (2024)

    Wasp: Warp scheduling to mimic prefetching in graphics workloads. arXiv preprint arXiv:2404.06156 (2024)

    2024 arXiv

  22. [30]

    Matthieu Tardy and Carter Edwards. 2020. Controlling Data Move- ment to Boost Performance on the NVIDIA Ampere Architecture. https://developer.nvidia.com/blog/controlling-data-movement-to-boost- performance-on-ampere-architecture/

    2020

  23. [31]

    Avinash Maurya, Jie Ye, M Mustafa Rafique, Franck Cappello, and Bogdan Nicolae

  24. [32]

    Microsoft. 2025. DeepNVMe. https://www.deepspeed.ai/tutorials/deepnvme/

    2025

  25. [33]

    Haikun Liu, Yujie Chen, Xiaofei Liao, Hai Jin, Bingsheng He, Long Zheng, and Ren- tong Guo. 2017. Hardware/software cooperative caching for hybrid DRAM/NVM memory architectures. In Proceedings of the International Conference on Super- computing. 1–10

    2017

  26. [34]

    Maxim Naumov, Dheevatsa Mudigere, Hao-Jun Michael Shi, Jianyu Huang, Narayanan Sundaraman, Jongsoo Park, Xiaodong Wang, Udit Gupta, Carole- Jean Wu, Alisson G Azzolini, et al. 2019. Deep learning recommendation model for personalization and recommendation systems.arXiv preprin...

    2019 arXiv

  27. [35]

    Giovanni Neglia, Damiano Carra, and Pietro Michiardi. 2018. Cache policies for linear utility maximization. IEEE/ACM Transactions on Networking 26, 1 (2018), 302–313

    2018

  28. [36]

    In Proceedings of the 14th Workshop on AI and Scientific Computing at Scale using Flexible Computing Infrastructures

    Breaking the memory wall: A study of i/o patterns and gpu memory utiliza- tion for hybrid cpu-gpu offloaded optimizers. In Proceedings of the 14th Workshop on AI and Scientific Computing at Scale using Flexible Computing Infrastructures . 9–16

  29. [37]

    Nvidia. 2025. cuBLAS | NVIDIA Developer. https://developer.nvidia.com/cublas

    2025

  30. [38]

    Sparsh Mittal and Shraiysh Vaishay. 2019. A survey of techniques for optimizing deep learning on GPUs. Journal of Systems Architecture 99 (2019), 101635

    2019

  31. [39]

    Nvidia. 2025. cuda::memcpy_async — libcudacxx 3.1 documentation. https://nvidia.github.io/cccl/libcudacxx/extended_api/asynchronous_ operations/memcpy_async.html?utm_source=ainews&utm_medium=email& utm_campaign=ainews-a-quiet-weekend

    2025

  32. [40]

    Nvidia. 2019. GPUDirect Storage: A Direct Path Between Storage and GPU Memory | NVIDIA Technical Blog. https://developer.nvidia.com/blog/gpudirect- storage/

    2019

  33. [41]

    Nvidia. 2025. 1. Introduction — PTX ISA 8.8 documentation. https://docs.nvidia. com/cuda/parallel-thread-execution/

    2025

  34. [42]

    Nvidia. 2025. NVIDIA/gdrcopy: A fast GPU memory copy library based on NVIDIA GPUDirect RDMA technology. https://github.com/NVIDIA/gdrcopy

    2025

  35. [43]

    Nvidia. 2025. CUDA C++ Programming Guide. https://docs.nvidia.com/cuda/ pdf/CUDA_C_Programming_Guide.pdf

    2025

  36. [44]

    Nvidia. 2013. Unified Memory in CUDA 6 | NVIDIA Technical Blog. https: //developer.nvidia.com/blog/unified-memory-in-cuda-6/

    2013

  37. [45]

    Nvidia. 2018. Using CUDA Warp-Level Primitives | NVIDIA Technical Blog. https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/

    2018

  38. [46]

    Nvidia. 2025. NVIDIA Ampere GPU Architecture Tuning Guide. https://docs. nvidia.com/cuda/ampere-tuning-guide/index.html

    2025

  39. [47]

    Stéfani Pires, Adriana Ribeiro, and Leobino N Sampaio. 2024. On learning suitable caching policies for in-network caching. IEEE Transactions on Machine Learning in Communications and Networking (2024)

    2024

  40. [48]

    Nvidia. 2025. RTX 5000 Ada Generation Graphics Card | NVIDIA. https://www. nvidia.com/en-us/design-visualization/rtx-5000/

    2025

  41. [49]

    Mohaimenul Azam Khan Raiaan, Md Saddam Hossain Mukta, Kaniz Fatema, Nur Mohammad Fahad, Sadman Sakib, Most Marufatul Jannat Mim, Jubaer Ah- mad, Mohammed Eunus Ali, and Sami Azam. 2024. A review on large Language Models: Architectures, applications, taxonomies, open issues and...

    2024

  42. [50]

    Samyam Rajbhandari, Olatunji Ruwase, Jeff Rasley, Shaden Smith, and Yuxiong He. 2021. Zero-infinity: Breaking the gpu memory wall for extreme scale deep learning. In Proceedings of the international conference for high performance com- puting, networking, storage and analysis . 1–14

    2021

  43. [51]

    NVM Express. 2025. NVM Express. https://nvmexpress.org/

    2025

  44. [52]

    2021.{Zero-offload}: Democratizing{billion-scale} model training

    Jie Ren, Samyam Rajbhandari, Reza Yazdani Aminabadi, Olatunji Ruwase, Shuangyan Yang, Minjia Zhang, Dong Li, and Yuxiong He. 2021.{Zero-offload}: Democratizing{billion-scale} model training. In 2021 USENIX Annual Technical Conference (USENIX ATC 21). 551–564

    2021

  45. [53]

    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. In Proceedings of t...

    2023

  46. [54]

    Samsung. 2025. Samsung 990 PRO PCIe 4.0 SSD | Samsung Semiconductor Global. https://semiconductor.samsung.com/consumer-storage/internal-ssd/990-pro/

    2025

  47. [55]

    Scott Beamer. 2024. sbeamer/gapbs: GAP Benchmark Suite. https://github.com/ sbeamer/gapbs

    2024

  48. [56]

    Shaina Raza and Chen Ding. 2019. Progress in context-aware recommender systems—An overview. Computer Science Review 31 (2019), 84–97

    2019

  49. [57]

    Tom’s Hardware. 2021. Samsung 980 Pro M.2 NVMe SSD Review: Redefining Gen4 Performance | Tom’s Hardware. https://www.tomshardware.com/reviews/ samsung-980-pro-m-2-nvme-ssd-review

    2021

  50. [58]

    Xiaowei Ren and Mieszko Lis. 2021. Chopin: Scalable graphics rendering in multi-gpu systems via parallel image composition. In 2021 IEEE International Symposium on High-Performance Computer Architecture (HPCA) . IEEE, 709–722

    2021

  51. [59]

    Yangzihao Wang, Andrew Davidson, Yuechao Pan, Yuduo Wu, Andy Riffel, and John D Owens. 2016. Gunrock: A high-performance graph processing library on the GPU. In Proceedings of the 21st ACM SIGPLAN symposium on principles and practice of parallel programming . 1–12

    2016

  52. [60]

    Yang Wang, Jiwu Shu, Guangyan Zhang, Wei Xue, and Weimin Zheng. 2010. Sopa: Selecting the optimal caching policy adaptively. ACM Transactions on Storage (TOS) 6, 2 (2010), 1–18

    2010

  53. [61]

    Paul Sweazey and Alan Jay Smith. 1986. A class of compatible cache consistency protocols and their support by the IEEE futurebus. ACM SIGARCH Computer Architecture News 14, 2 (1986), 414–423

    1986

  54. [62]

    Kun Wu, Jeongmin Brian Park, Xiaofan Zhang, Mert Hidayetoğlu, Vikram Sharma Mailthody, Sitao Huang, Steven Sam Lumetta, and Wen-mei Hwu. 2024. SSDTrain: An Activation Offloading Framework to SSDs for Faster Large Language Model Training. arXiv preprint arXiv:2408.10013 (2024)

    2024 arXiv

  55. [63]

    Pengyu Wang, Jing Wang, Chao Li, Jianzong Wang, Haojin Zhu, and Minyi Guo. 2021. Grus: Toward unified-memory-efficient high-performance graph processing on gpu. ACM Transactions on Architecture and Code Optimization (TACO) 18, 2 (2021), 1–25

    2021

  56. [64]

    Jones, and Peipei Zhou

    Zhuoping Yang, Jinming Zhuang, Jiaqi Yin, Cunxi Yu, Alex K. Jones, and Peipei Zhou. 2023. AIM: Accelerating Arbitrary-precision Integer Multiplication on Heterogeneous Reconfigurable Computing Platform Versal ACAP. In ICCAD

    2023

  57. [65]

    Haoyang Zhang, Yirui Zhou, Yuqi Xue, Yiqi Liu, and Jian Huang. 2023. G10: Enabling an efficient unified gpu memory and storage architecture with smart tensor migrations. In Proceedings of the 56th Annual IEEE/ACM International Symposium on Microarchitecture. 395–410

    2023

  58. [66]

    Zeke Wang, Hongjing Huang, Jie Zhang, Fei Wu, and Gustavo Alonso. 2022. {FpgaNIC}: An{FPGA-based} versatile 100gb{SmartNIC} for{GPUs}. In 2022 USENIX Annual Technical Conference (USENIX ATC 22) . 967–986

    2022

  59. [67]

    Jinming Zhuang, Jason Lau, Hanchen Ye, Zhuoping Yang, Yubo Du, Jack Lo, Kristof Denolf, Stephen Neuendorffer, Alex Jones, Jingtong Hu, Deming Chen, Jason Cong, and Peipei Zhou. 2023. CHARM: Composing Heterogeneous AcceleR- ators for Matrix Multiply on Versal ACAP Architecture....

    2023

  60. [68]

    Shao-Peng Yang, Minjae Kim, Sanghyun Nam, Juhyung Park, Jin-Yong Choi, Eyee Hyun Nam, Eunji Lee, Sungjin Lee, and Bryan S Kim. 2023. Overcoming the memory wall with{CXL-Enabled}{ SSDs}. In 2023 USENIX Annual Technical Conference (USENIX ATC 23). 601–617

    2023

  61. [69]

    Jones, Jingtong Hu, Yiyu Shi, and Peipei Zhou

    Jinming Zhuang, Zhuoping Yang, Shixin Ji, Heng Huang, Alex K. Jones, Jingtong Hu, Yiyu Shi, and Peipei Zhou. 2024. SSR: Spatial Sequential Hybrid Architecture for Latency Throughput Tradeoff in Transformer Acceleration. In Proceedings of the 2024 ACM/SIGDA International Sympos...

    2024

  62. [70]

    Jinming Zhuang, Zhuoping Yang, and Peipei Zhou. 2023. High Performance, Low Power Matrix Multiply Design on ACAP: from Architecture, Design Challenges and DSE Perspectives. In 2023 60th ACM/IEEE Design Automation Conference (DAC). 1–6. doi:10.1109/DAC56929.2023.10247981

    2023

  63. [71]

    Yu Zhang, Yuxuan Liang, Jin Zhao, Fubing Mao, Lin Gu, Xiaofei Liao, Hai Jin, Haikun Liu, Song Guo, Yangqing Zeng, et al. 2022. Egraph: efficient concurrent GPU-based dynamic graph processing. IEEE Transactions on Knowledge and Data Engineering 35, 6 (2022), 5823–5836

    2022

  64. [73]

    Jinming Zhuang, Shaojie Xiang, Hongzheng Chen, Niansong Zhang, Zhuoping Yang, Tony Mao, Zhiru Zhang, and Peipei Zhou. 2025. ARIES: An Agile MLIR- Based Compilation Flow for Reconfigurable Devices with AI Engines. In Proceed- ings of the 2025 ACM/SIGDA International Symposium o...

    2025

  65. [2023]

    Proceedings of Machine Learning and Systems 5 (2023), 680–694

    Alcop: Automatic load-compute pipelining in deep learning compiler for ai-gpus. Proceedings of Machine Learning and Systems 5 (2023), 680–694

    2023

  66. [2024]

    In 2024 IEEE International Symposium on High- Performance Computer Architecture (HPCA)

    Wasp: Exploiting gpu pipeline parallelism with hardware-accelerated automatic warp specialization. In 2024 IEEE International Symposium on High- Performance Computer Architecture (HPCA). IEEE, 1–16

    2024

  67. [2025]

    In Proceedings of the 2025 ACM Great Lakes Symposium on VLSI (GLSVLSI ’25)

    ART: Customizing Accelerators for DNN-Enabled Real-Time Safety-Critical Systems. In Proceedings of the 2025 ACM Great Lakes Symposium on VLSI (GLSVLSI ’25)

    2025

Pith tools

Reviewed August 16, 2026 · model on record in the stance chip above.