GNStor: GPU-Native Remote All-Flash Array

• GNStor is a GPU-native high-performance remote all-flash array system that allows GPUs to directly access SSDs by bypassing the CPU in the I/O path.
• It employs GNoR for GPU-centric NVMe over RDMA and deEngine for decentralized SSD firmware functions to minimize overhead and improve scalability.
• The system achieves up to 3.2× higher throughput and reduced latency, making it ideal for data-intensive workloads such as LLM training and graph analytics.

Searching arXiv for GNStor and closely related GPU-storage work. GNStor is a GPU-native high-performance remote all-flash array (AFA) system that enables GPUs to directly access remote AFA without CPU intervention in the I/O path, thereby redesigning the storage path around the GPU as the active initiator rather than a passive compute endpoint behind CPU-managed storage software (Yi et al., 3 Jun 2026). The system is motivated by a mismatch in conventional GPU-AFA deployments: GPUs dominate computation in data-intensive workloads such as graph analytics, recommender systems, LLM training, and LLM inference, yet remote storage access and AFA-level services remain CPU-centric. GNStor addresses this disparity through two core mechanisms—GNoR, a GPU-centric NVMe over RDMA stack, and deEngine, a decentralized AFA engine embedded into SSD firmware—while retaining a CPU-resident GNStor daemon for infrequent control-plane tasks (Yi et al., 3 Jun 2026).

1. Conceptual basis and problem setting

GNStor is situated in the context of remote AFAs used to provide large capacity, efficient sharing, strong durability, and high performance for GPU applications. The motivating workloads include PB-scale LLM training corpora shared by many GPU clients, replicated checkpoints, graph analytics and GNN workloads with large shared datasets and latency-sensitive random accesses, and LLM inference with persistent shared KV cache (Yi et al., 3 Jun 2026). The central diagnosis is architectural: on the GPU node, the CPU orchestrates remote I/O, sends NVMe over RDMA requests, receives data through the NIC, and often forwards data between host memory and GPU memory; on the AFA node, a centralized CPU-resident AFA engine manages logical volumes, access control, metadata, redundancy, and consistency.

The paper identifies three major overhead sources in this conventional design: CPU-GPU interaction overhead, centralized orchestration overhead, and I/O traffic amplification. Even where GPUDirect improves data movement, orchestration remains CPU-bound. The reported impact is severe: CPU-centric design can cause up to 79.1% throughput degradation from constrained GPU-to-AFA access; simply pushing AFA-level coordination to clients can still cause 72.0% throughput degradation; and naively porting CPU-oriented NoR synchronization to GPUs can lead to 86.6% throughput degradation (Yi et al., 3 Jun 2026).

In this setting, “GPU-native” has a specific meaning. It does not denote the elimination of CPUs from the system as a whole. Rather, it denotes removal of CPU software from the I/O critical path while leaving infrequent or non-I/O-critical management operations on CPUs. This distinction is central to the design and addresses a common misunderstanding that GNStor is merely a GPUDirect-style data-path optimization; the paper’s claim is broader, namely an end-to-end redesign of control and metadata responsibilities around a GPU-initiated storage path (Yi et al., 3 Jun 2026).

2. System organization and end-to-end I/O path

GNStor consists of three major components: GNoR, deEngine, and the GNStor daemon. Their responsibilities are sharply partitioned between fast-path and control-plane work.

Component	Location	Role
GNoR	GPU side	GPU-centric NVMe over RDMA fast path
deEngine	SSD firmware	Decentralized AFA I/O-critical functions
GNStor daemon	AFA-side CPU	Infrequent management operations

The GNStor daemon handles volume creation, volume teardown, permission changes, client authentication or identity verification, and distributing volume metadata to SSDs. These actions are explicitly outside the steady-state I/O path. GNoR establishes GPU-usable NoR channels, stores I/O-critical queue structures in GPU memory, submits NoR command capsules directly from GPU, polls completions from GPU, manages pre-registered RDMA memory pools, and supports synchronous, asynchronous, and batched I/O. deEngine runs inside each SSD firmware and locally validates request legitimacy, checks whether the SSD is an intended request target, maintains the mapping from volume addresses to physical flash locations, persists metadata using SSD-internal durability mechanisms, and recovers metadata after failures (Yi et al., 3 Jun 2026).

The contrast between the conventional and GNStor I/O paths is fundamental. In the CPU-centric baseline path, a GPU requests I/O through the CPU; the CPU constructs and submits remote storage commands; data returns to host memory via the NIC and may then be copied to GPU memory; and the storage-side CPU-based AFA engine performs access control and metadata work before issuing SSD commands. In GNStor, a GPU kernel issues I/O directly through GNoR; GNoR creates and submits NoR capsules from GPU memory; NIC doorbells are rung directly by GPU through PCIe MMIO; the AFA NIC’s HCA parses NoR capsules and forwards them to SSDs via PCIe peer-to-peer communication; RDMA one-sided operations move data directly between GPU memory and the SSD-side path; and deEngine validates, executes, and updates metadata at the SSD (Yi et al., 3 Jun 2026).

The workflow is described as eight numbered steps: a client asks the daemon to create or access a volume; the daemon propagates access-control metadata to SSDs; the daemon returns volume metadata to the client; the client initializes GNoR channels to SSDs; the client uses a hash-based strategy to choose target SSDs for volume I/O; the GPU sends NoR capsules via GNoR; the AFA NIC HCA parses capsules and forwards NVMe commands to SSDs with PCIe P2P; and deEngine validates and executes requests and updates metadata. This decomposition is significant because it preserves AFA semantics while bypassing CPU software in steady-state data access (Yi et al., 3 Jun 2026).

3. GNoR: GPU-centric NVMe over RDMA

GNoR is the GPU-side remote storage stack that makes direct GPU-initiated remote SSD access feasible over NVMe over RDMA. Its relevance derives from the observation that prior GPU-centric storage systems primarily target local NVMe over PCIe access, whereas GNStor requires remote access with RDMA-specific queue management, memory-region registration, NoR capsule handling, and completion handling at GPU-scale concurrency. The paper explicitly argues that a CPU-style NoR implementation cannot be directly ported to GPU because CPU implementations assume far fewer threads and rely on heavier lock-based synchronization (Yi et al., 3 Jun 2026).

The main abstraction is the channel, which packages all I/O-critical state for a remote path into a structure containing an NVMe I/O queue pair, RDMA queue pairs, NIC doorbell register addresses, pre-registered memory buffers, and auxiliary state such as queue tail pointers. These channel structures reside in GPU memory. CPU responsibilities remain limited to RDMA connection establishment, keep-alive handling, and NVMe admin queues. Channel initialization proceeds through a mixed CPU/GPU sequence: the CPU establishes the NoR connection and creates the NVMe admin queue pair; for each channel, GNoR allocates channel data structures in GPU memory and starts a new NoR session with queues on GPU memory; the channel pointer is passed to the GPU; and the GPU completes setup by pre-posting RDMA receive requests and issuing the required Fabrics Connect command (Yi et al., 3 Jun 2026).

A GPU kernel then directly initiates NoR I/O: it encapsulates an I/O request as a NoR command capsule, submits it to the RDMA send queue, rings the NIC doorbell register via PCIe MMIO write, polls the RDMA completion queue, receives NVMe completion entries into buffers associated with RDMA Recvs, parses CQEs, reposts receives, and triggers callbacks for asynchronous I/O. This is not merely zero-copy transfer; it is GPU-side control over request issuance and completion processing (Yi et al., 3 Jun 2026).

A major obstacle in RDMA systems is memory region registration, which is expensive because it requires memory pinning and NIC-visible registration through system calls. GNoR therefore relies on pre-allocation and pre-registration of per-channel GPU memory pools. The paper rejects both a CPU-style buddy allocator and a simple fixed-size bitmap allocator, respectively because of complex lock-based synchronization and fragmentation-induced proliferation of RDMA operations. In their place, GNoR uses a GPU-friendly multi-level allocator with a small number of size levels such as 4 KB, 64 KB, and 1 MB, one bitmap per level, closest-match allocation, preferred contiguous blocks, selective splitting, opportunistic merging, and synchronization through atomic primitives such as CAS. If the pool is insufficient, it can expand dynamically, scaling by a factor of 2 by default (Yi et al., 3 Jun 2026).

The synchronization design is explicitly atomic-operation-based. GNoR “replaces heavyweight locks with lightweight atomic operations” and coordinates concurrent channel access using atomics. This design is mapped to the GPU SIMT model: a thread block cooperates on batched I/O; the warp is the 32-thread lockstep unit; and each thread may correspond to one request slot in a batch. In the batched workflow, a shared-memory bitmap tracks request status, threads allocate buffers and prepare requests in parallel, command capsules are appended to the RDMA send queue using atomics, a designated thread—usually thread 0—rings the doorbell, another designated thread polls the RDMA completion queue and dispatches CQEs, callbacks are executed, and incomplete slots remain pending into the next round. The stated rationale is reduced warp divergence, lower synchronization overhead, better exploitation of massive parallelism, and improved throughput (Yi et al., 3 Jun 2026).

GNStor exposes this functionality through libgnstor, which provides GPU-callable APIs for memory allocation and free, synchronous I/O, and asynchronous I/O. Batched I/O is presented as stage-wise primitives—submit, commit, poll, dispatch—rather than a monolithic call because a CUDA __device__ function can only be called by a single thread, not assigned to an entire warp as a first-class unit. This API structure reflects the paper’s broader design stance that the storage stack must be adapted to GPU execution semantics rather than merely recompiled for the device (Yi et al., 3 Jun 2026).

4. deEngine: decentralized AFA engine in SSD firmware

deEngine addresses the problem that a GPU-bypassed I/O path would otherwise bypass the storage-side centralized AFA engine as well. The key insight is that many AFA I/O-critical functions resemble existing SSD firmware functions: the AFA engine checks request legitimacy and permissions, while SSD firmware already validates request legality; the AFA engine maps volume addresses to SSD logical addresses, while the SSD FTL maps SSD logical addresses to physical pages. GNStor therefore decomposes AFA functionality and integrates the I/O-critical parts into each SSD firmware (Yi et al., 3 Jun 2026).

The division of labor is explicit. The GNStor daemon retains volume creation and teardown, permission granting and revocation, volume metadata dissemination, and write lease management. deEngine assumes responsibility for request target validation, access control enforcement for I/O, address mapping from volume space into flash space, and metadata update and persistence during I/O. This decentralization is not simply a sharding of a CPU engine; it is a relocation of functions to the storage device boundary so that correctness and AFA semantics remain available under direct GPU-to-SSD access (Yi et al., 3 Jun 2026).

Volume management proceeds through daemon-mediated metadata dissemination. When a client requests a new volume, it contacts the daemon via RPC; the daemon verifies identity; allocates a unique VID, a hash factor, and capacity and replication metadata; propagates this metadata to all SSDs using customized VOLUME ADD commands; causes each SSD to store the metadata in local DRAM and flash as a volume permission table; and returns the volume metadata to the client. The daemon uses VOLUME CHMOD for permission updates and VOLUME DELETE/CHMOD for release or removal. The paper states that at most one client may hold write permission to a volume at a time, controlled through a lease-based mechanism, with periodic renewal such as every 5 minutes. GNStor assumes trusted clients in private-cluster-style environments; access control is intended for correct sharing and consistency rather than defense against malicious clients (Yi et al., 3 Jun 2026).

Request routing uses a hash-based load-balance function over [VID, VBA] and a factor. For writes, multiple SSDs are selected according to the replication factor. Each deEngine recomputes the same hash and checks whether the SSD is a valid target and whether the client has permission according to the local volume permission table. The SSD then uses weight round robin scheduling, which the paper identifies as the default method in most commercial SSDs. An important invariant follows: the client specifies only [VID, VBA], while the final physical location is chosen by SSD firmware. This prevents different clients from claiming the same physical storage space and removes the need for a centralized AFA mapping-table lookup (Yi et al., 3 Jun 2026).

One of deEngine’s most distinctive mechanisms is mapping-table unification. A conventional GPU-AFA stack maintains an AFA-level mapping from [VID, VBA] to SSD number and SSD logical page address, alongside an SSD-internal FTL mapping from LPA to PPA. GNStor merges these inside SSD firmware by modifying the FTL mapping table from conventional LPA-to-PPA to [VID,VBA]-to-PPA. The paper notes that VID and VBA themselves need not be stored explicitly in the table and that cuckoo hash is used to find a mapping-table slot that stores the PPA for a given [VID,VBA]. This eliminates the need for GPU-side or AFA-side volume mapping tables (Yi et al., 3 Jun 2026).

Persistence and recovery leverage enterprise SSD mechanisms, particularly power loss protection (PLP). SSDs use power hold-up circuits and capacitors so that buffered metadata in DRAM can be flushed to flash upon power loss and FTL metadata can be recovered after reboot. Because GNStor integrates AFA-level mapping into SSD FTL, the paper states that this eliminates the need for additional logging or checkpointing in the AFA layer. On recovery, a GPU client re-establishes identity with the daemon, retrieves volume metadata, reconstructs GNoR channels, and resumes I/O; after an AFA reboot, each SSD restores local mapping tables and the volume permission table, then the daemon reconstructs global state from SSD-resident permission tables (Yi et al., 3 Jun 2026).

5. Protocol, implementation, and experimental characterization

The implementation modifies NVMe protocols based on NVMe Base Specification 2.3 and NVMe Command Set Specification 1.2. Customized management commands—VOLUME ADD, VOLUME DELETE, and VOLUME CHMOD—are implemented as NVMe admin commands. The paper also reuses the leftmost 32 bits of the start logical block address (SLBA) field in NVMe I/O commands for VID and client ID, split as 16 bits each. This supports up to 16,384 clients, each with 16,384 volumes, and each volume up to 16 TB, since that equals $2^{32} \times 4 \text{ KB}$ (Yi et al., 3 Jun 2026).

The evaluation platform uses one GPU node and one AFA node connected by Mellanox ConnectX-7 NICs at 200 Gbps, with 21.6 GB/s goodput in RDMA. The GPU node consists of AMD EPYC 9654, 96 cores at 2.4 GHz, 768 GB DDR5, and an NVIDIA A100 with 40 GB device memory. The AFA node uses Intel Xeon 5320, 26 cores at 2.2 GHz, and 256 GB DDR4. deEngine firmware is implemented in NVMeVirt, with each emulated SSD configured to match high-performance commercial devices at 6.8 GB/s peak read bandwidth and 4.8 GB/s peak write bandwidth. The deEngine hash and address-control logic was characterized from a Xilinx Kintex UltraScale+ KU15P FPGA at 300 MHz, adding 276 ns for hash computation used in address mapping and access control. The software stack includes Ubuntu 22.04, Linux 5.4.0, CUDA 12.9, DOCA 3.3.0, and SPDK 25.09. Implementation size is reported as 9 K LOC for GNoR, 3 K LOC for the GNStor daemon, and 2 K LOC for deEngine firmware modifications (Yi et al., 3 Jun 2026).

Three baselines are used: Basic, a conventional CPU-centric GPU-AFA design using cudaMemcpy and a centralized AFA engine; GD, which adds NVIDIA GPUDirect for direct NIC↔GPU transfer but retains CPU orchestration; and GNStor, the full system with GNoR and deEngine. The microbenchmarks use 4 KB and 64 KB request sizes, sequential and random read/write patterns, I/O depth 32, one client by default, one CPU thread and one GPU warp, one GNoR channel, and 4 SSDs with 2 replicas by default. The AFA engine in CPU-centric designs uses 8 CPU cores to align with SSD/CPU-core ratios in commercial AFAs (Yi et al., 3 Jun 2026).

Real applications include vector addition on two vectors of 1 billion elements, matrix multiplication on two $16384 \times 16384$ matrices, ImageNet-100 resize with bilinear interpolation, GAP graph workloads including BFS, Connected Components, and SSSP, and GPT-2 training with varying batch sizes and model sizes. The chosen metrics are throughput, latency, scalability with clients, scalability with SSD count, application execution time, and GPU memory footprint (Yi et al., 3 Jun 2026). This workload profile places GNStor in a broader trend of GPU-centric storage research for large-scale data-intensive learning and graph processing; related work on storage-intensive GPU pipelines includes out-of-core GNN systems such as AGNES, which likewise emphasizes storage-path bottlenecks but in the single-machine training setting (Jang et al., 4 Jan 2026).

6. Performance results, interpretation, and limitations

GNStor’s headline result is that it achieves 3.2× higher I/O throughput and reduces application execution time by 31.1% compared to state-of-the-art AFA systems (Yi et al., 3 Jun 2026). In microbenchmarks at I/O depth 32, Basic reaches only 0.5 GB/s read and 0.3 GB/s write in 4 KB tests; GD improves 4 KB throughput by 1.2× for reads and 1.3× for writes over Basic; and GNStor achieves the best throughput in all workloads, outperforming Basic by 3.2× and GD by 0.8× on average in 4 KB tests. In latency terms, GD reduces latency by 40.7% in 4 KB tests and 26.1% in 64 KB tests versus Basic, while GNStor further reduces read and write latency by 35.7% and 39.8%, respectively, through the fully CPU-bypassed shortest I/O path (Yi et al., 3 Jun 2026).

Scalability results are similarly emphasized. When clients scale from 1 to 32, GNStor reaches 11.8 GB/s for 4 KB read, 5.6 GB/s for 4 KB write, and 9.6 GB/s for 64 KB write, nearly reaching SSD throughput limits. In 64 KB tests, GNStor reaches 21.5 GB/s with only two clients, corresponding to 99.5% of network bandwidth and 85.7% of GPU PCIe throughput. As SSD count increases, Basic and GD improve only slightly because of CPU-centric synchronization bottlenecks, whereas GNStor scales well; with five SSDs, it reaches 13.6 GB/s in 4 KB sequential read and 6.6 GB/s in 4 KB sequential write, which are 3.6× and 4.8× higher than GD, respectively (Yi et al., 3 Jun 2026).

The ablation study separates the effects of deEngine and GNoR. Adding deEngine to GD as GD+deEngine improves write throughput relative to GD by 49.9% for 4 KB and 24.7% for 64 KB, with latency improvements of 32.4% for 4 KB and 11.1% for 64 KB. The explanation given is that deEngine benefits writes more because writes require metadata updates and consistency-related persistence, whereas reads in GD do not frequently update AFA-level mappings. Relative to GD+deEngine, full GNStor further improves throughput by 49.6% and reduces latency by 31.2% on average, isolating GNoR’s contribution in reducing CPU-GPU interaction and enabling fully GPU-centric remote access (Yi et al., 3 Jun 2026).

At the application level, GNStor reduces I/O time in vector addition by 85.5% versus Basic and 72.8% versus GD, yielding end-to-end speedups of 5.7× and 3.1×. In matrix multiplication, where compute dominates, GNStor still outperforms Basic by 24.7% and GD by 12.4%. For ImageNet-100 resize on 16 GB batches, throughput is 3.0× higher than Basic and 1.5× higher than GD, with I/O-time reductions of 73.5% and 58.4%. For BFS, GNStor achieves 7.3× I/O speedup over Basic and 4.0× over GD, with end-to-end execution time reduced by 77.6% and 62.5%; for CC and SSSP, execution time still improves by 29.2% and 17.3% on average over Basic and GD. In GPT-2 training, GNStor reduces total execution time for batch size 4 by 46.5% versus Basic and 26.4% versus GD, shortening Load by 84.0% and Checkpoint by 68.7% versus Basic; with larger models, the paper reports up to 1.3× speedup over GD for GPT-2 medium (Yi et al., 3 Jun 2026).

GPU memory footprint is characterized as modest because AFA mapping tables are not kept on the GPU. Per channel, queue structures consume about 50 KB for concurrency 128 and the memory pool consumes 8 MB. The paper states that 32 channels are enough to maximize throughput in the testbed, consuming about 258 MB, which is described as small relative to A100-class GPU memory (Yi et al., 3 Jun 2026).

Several limitations are explicitly identified. The main practical constraint is SSD firmware modification, which is difficult with off-the-shelf SSDs because firmware is manufacturer-controlled. GNStor also currently uses CUDA and DOCA, though the authors state that GNoR is not fundamentally tied to NVIDIA hardware and could be ported to GPUs with similar programming models and NICs capable of exposing doorbell registers to the GPU. The system assumes trusted clients and private-cluster deployment; stronger authentication would be needed in untrusted environments. The work also does not claim to solve all management or security issues of distributed storage, focusing instead on removal of CPU overhead from the I/O critical path while preserving needed AFA functionality (Yi et al., 3 Jun 2026).

A plausible implication is that GNStor marks a shift in remote GPU storage design from improving isolated links in the data path to redesigning the full end-to-end control, metadata, and execution substrate so that the GPU, NIC, and SSD firmware jointly implement the fast path. In that sense, its significance lies not only in throughput and latency gains, but in establishing a systems pattern: GPU-initiated NoR access, NIC NoR target offloading, and SSD-firmware-resident AFA services as a coherent architecture for remote flash arrays (Yi et al., 3 Jun 2026).