{"id":"03c17fab-e037-496b-b1ac-877aaae8e114","arxiv_id":"2504.19365","paper_version":3,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":7.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":0,"one_line_summary":"AGILE is a GPU-centric asynchronous I/O library that uses a background GPU service to avoid deadlock while issuing NVMe requests, with a customizable HBM software cache and up to 1.75x speedup over BaM on recommendation models.","lead":"AGILE is a new library that lets a GPU read and write SSDs directly from its threads without waiting for each disk request to finish. It hides slow SSD delays behind computation, giving up to 1.88x speedups on data-heavy GPU jobs compared with older synchronous methods.","discovery_kind":"new_method","skeptic_critique":{"model":"deepseek-v4-flash","headline":"AGILE's deadlock-free asynchronous model depends on the service kernel being scheduled and making progress, which the paper never verifies under full user-kernel load.","rationale":"The reader's weakest_assumption identifies precisely the service kernel's ability to coexist and make progress as the load-bearing assumption. My stress test confirms this is the single most critical unverified precondition: the entire asynchronous model and the deadlock-elimination claim are void if the service starves. The paper provides strong empirical evidence that the system works in the evaluated configurations, but no general forward-progress guarantee, no overhead characterization of the service's polling loop under full user-kernel load, and no analysis of the scheduling interaction between service and user kernels. The reader's conditional verdict is appropriate: the paper should be accepted only if the authors address this by providing either a scheduling guarantee (e.g., dedicated SM or priority stream evidence) or a quantitative evaluation of service polling latency under varying occupancy. I therefore do not propose changing the reader's verdict, but I underscore the importance of this specific check.","tokens_in":21528,"tokens_out":5888,"duration_ms":64767,"concrete_test":"Instrument the AGILE service kernel to timestamp each CQ doorbell update and record the number of CQEs processed per fixed interval, then run a microbenchmark at a balanced computation-to-communication ratio (e.g., CTC≈1) while sweeping user-kernel occupancy from 25% to 100% of each SM's warp slots (via __launch_bounds__ and varying block count). If the service's CQE processing rate drops below the SSD completion rate and the async speedup over synchronous I/O collapses toward 1.0 as occupancy approaches 100%, the forward-progress concern is confirmed.","verdict_should_be":"UNCHANGED","load_bearing_attack":"The central async/deadlock-free claim rests on the AGILE service kernel (Section 3.2) running concurrently with user kernels and polling CQs fast enough to free SQ entries and complete transactions. Listing 1 shows the service is a separate persistent kernel launched before the user kernel (line 40) and the user kernel runs afterward on the same GPU (line 42). CUDA provides no forward-progress guarantee for co-resident kernels; if the user kernel occupies all SM issue slots, the service warp can be starved, CQs go unpolled, SQ entries are never released, and the asynchronous pipeline degrades toward synchronous behavior. The paper reports the service uses 37 registers/thread (Section 4.6) but never measures its achieved polling rate, scheduling latency, or impact on the user kernel's occupancy. The service's liveness is assumed, not established. The authors' own experiments show the async advantage is conditional: with one queue pair (Figure 9) or small software cache (Figure 10), async collapses to sync or worse. Those results are explained as resource contention, but they underscore that the claimed 'elimination' of deadlock and robust overlap require a progress guarantee the paper does not provide. Without evidence that the service always makes progress, the headline 1.88x overlap and 1.75x DLRM speedups cannot be claimed as generally guaranteed.","agreement_with_reader":"agree"},"referee_report":{"model":"deepseek-v4-flash","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.","tokens_in":21777,"tokens_out":4432,"duration_ms":37224,"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":[{"comment":"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.","section":"Section 3.2, Listing 1, Section 4.6"},{"comment":"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.","section":"Sections 4.1-4.5"},{"comment":"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.","section":"Sections 3.4 and 3.5"}],"minor_comments":[{"comment":"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.","section":"Abstract and Section 4.5"},{"comment":"The x-axis labels and legend in Figure 4 are garbled and hard to read; the figure should be regenerated with clear axis labels.","section":"Figure 4"},{"comment":"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.","section":"Section 3.1 and Listing 1"},{"comment":"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.","section":"Section 4.4"},{"comment":"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.","section":"Section 4.6"}],"recommendation":"major_revision","confidential_remarks":"The manuscript header indicates the paper has been accepted at SC'25, which I treated as neutral contextual information and not part of my technical assessment. My main technical concern is that the liveness of the service kernel is asserted rather than demonstrated; this is fixable in revision and does not require a new architecture. The lack of variance reporting in the evaluation should also be addressed. I would not reject the paper, but I would ask for these load-bearing points to be resolved."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"Colleague,\n\nShort version: AGILE is the first GPU-centric asynchronous NVMe I/O library I know of that lets GPU threads issue commands without holding locks, using a small co-resident service kernel to reap completions. The design is sensible, the implementation is open-sourced, and the evaluation against BaM on DLRM and graph workloads supports the headline speedups. The paper is worth a serious referee and likely deserves acceptance after minor revision; the one thing that nags me is liveness.\n\nWhat's genuinely new: BaM is synchronous, DeepNVMe and GPUDirect Storage are CPU-initiated. AGILE's deadlock avoidance by handing lock ownership to a dedicated service kernel is a real step forward. The flexible software cache with pluggable policy and the share-table MOESI-inspired coherency are useful additions. The empirical work is solid: DLRM speedups up to 1.75x over BaM, graph API overhead reductions, and register savings. The ideal-speedup formula is a reference bound, not a fitted model, so no circularity. Citation pattern is fine; self-citations are confined to FPGA background and are not load-bearing.\n\nSoft spots in proportion: the evaluation lacks error bars or repetition counts, and reports only normalized speedups, not absolute times. Small inconsistency: abstract says 3.12x cache-overhead reduction, the body text gives 3.17x for SpMV K-graph. Minor but should be fixed. More substantive: the stress-test concern is real. The deadlock-free claim depends on the AGILE service kernel making progress while the user kernel occupies the GPU. CUDA gives no forward-progress guarantee for co-resident kernels, and the paper never measures the service's achieved polling rate, scheduling latency, or worst-case behavior under full SM load. The authors' own experiments show async advantage collapses with one queue pair or small cache, which they attribute to SQE and cache-line contention. That is consistent with the service periodically being unable to recycle entries fast enough. It doesn't invalidate the design, but \"eliminates deadlock risks\" should be scoped to conditions where the service is schedulable; a brief discussion or a simple measurement of service polling delay would close the gap.\n\nNet: this is an honest, well-engineered systems paper. It deserves peer review, and I'd accept it with the requested revisions on reproducibility and liveness. I'd cite it if I worked on GPU-storage systems, and I'd bring it to a reading group focused on GPU memory expansion.","headline":"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.","tokens_in":22310,"tokens_out":2321,"would_cite":true,"duration_ms":23172,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"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.","keywords":["GPU","NVMe","asynchronous I/O","software cache","deadlock avoidance","HBM","storage systems"],"falsifier":"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.","tokens_in":21337,"feed_emoji":"⚡","tokens_out":7630,"duration_ms":70768,"temperature":0.7,"pith_summary":"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.","feed_headline":"GPU threads issue NVMe I/O asynchronously, up to 1.88x faster","feed_subtitle":"A lightweight GPU-side service kernel polls completion queues, freeing threads from stalls and unlocking overlap.","key_machinery":"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.","core_discovery":"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.","pith_inferences":["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."],"forward_implications":["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."],"supporting_citations":[{"why":"The synchronous GPU-centric storage access baseline that AGILE extends and compares against in DLRM and graph experiments.","marker":"[48]"},{"why":"An asynchronous I/O stack for SSDs that motivates the latency-hiding model AGILE brings into the GPU.","marker":"[27]"},{"why":"The GPU instruction-level asynchronous copy primitive whose memory-path limits AGILE's async_issue extends.","marker":"[36]"},{"why":"The asynchronous copy API that AGILE's async_issue(src,dst) mimics while allowing arbitrary GPU buffers.","marker":"[39]"},{"why":"Provides the GPU memory pinning and physical-address mapping needed for direct SSD-to-HBM DMA.","marker":"[42]"},{"why":"Defines the MOESI coherence states that the Share Table adapts for user-specified buffers.","marker":"[56]"}],"fun_headline_variants":["AGILE: async GPU I/O, up to 1.88x faster","Lightweight async GPU-SSD with deadlock-free I/O","AGILE achieves 1.88x via async NVMe polling","GPU-centric async I/O library: AGILE speeds up to 1.88x"],"cache_read_input_tokens":3200,"weakest_assumption_plain":"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.","fun_headline_variants_meta":{"raw":{"variants":["AGILE: async GPU I/O, up to 1.88x faster","Lightweight async GPU-SSD with deadlock-free I/O","AGILE achieves 1.88x via async NVMe polling","GPU-centric async I/O library: AGILE speeds up to 1.88x"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000134,"raw_usage":{"total_tokens":1165,"prompt_tokens":993,"completion_tokens":172,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":609,"completion_tokens_details":{"reasoning_tokens":91}},"tokens_in":609,"tokens_out":172,"duration_ms":2184,"temperature":1.0,"reasoning_tokens":91,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-16T05:53:39.313046+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"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.","supporting_citations":[{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"The synchronous GPU-centric storage access baseline that AGILE extends and compares against in DLRM and graph experiments."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"An asynchronous I/O stack for SSDs that motivates the latency-hiding model AGILE brings into the GPU."},{"cited_title":"In Proceedings of the 14th Workshop on AI and Scientific Computing at Scale using Flexible Computing Infrastructures","cited_arxiv_id":null,"evidence_quote":"The GPU instruction-level asynchronous copy primitive whose memory-path limits AGILE's async_issue extends."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"The asynchronous copy API that AGILE's async_issue(src,dst) mimics while allowing arbitrary GPU buffers."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Provides the GPU memory pinning and physical-address mapping needed for direct SSD-to-HBM DMA."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Defines the MOESI coherence states that the Share Table adapts for user-specified buffers."}],"review_version":1}