Architecture & Measurement-Design Analysis

Quantifying Performance Gains of GPUDirect Storage

Source: Inupakutika, D.; Davis, B.; Yang, Q.; Kim, D.; Akopian, D. 2022 IEEE International Conference on Networking, Architecture and Storage (NAS), pp. 1-8. DOI: 10.1109/NAS55553.2022.9925516 Authors: Samsung (San Diego, CA) + University of Texas at San Antonio. Reader: Direct PDF read (gemini-reader quota exhausted; codex-reader unavailable). 9-page paper read whole via the pages parameter. Analyst: Vishwakarma Date: 2026-06-15

Table of Contents

1. Evaluation Harness Architecture (the "instrument")
2. System-Under-Test Architecture (GDS vs the bounce-buffer baseline)
3. Design-Space Diagram (axes swept vs axes held fixed)
4. Algorithm / Control Flow Diagrams (the two data paths)
5. Quantitative Results — Empirical Findings by Regime
6. Configuration-Regime Trade-off Tables
7. Bottlenecks & Insights Surfaced by the Measurements
8. Limitations of the Methodology
9. Note on NCCL Tuning
10. Analogy

1. Evaluation Harness Architecture (the "instrument")

The harness is a two-path A/B comparator: every workload is run twice over the same storage device — once on the GDS path (cuda_io=cufile, O_DIRECT, direct=1) and once on the traditional bounce-buffer path (cuda_io=posix, cudaMemcpy) — and the difference in throughput, IOPS, latency, CPU utilization, and GPU utilization is reported. The comparison is deliberately a paired measurement: the SSD, the GPU, the filesystem, and the workload are held identical across the two arms, so the only independent variable is which data path the bytes travel. This is the cleanest possible isolation of GDS's effect.

Crucially, the harness is run over two storage backends so that the GDS benefit can be measured both where the device is close (local DAS) and where it is far (a disaggregated Weka cluster across a 100 GbE fabric). The paper calls these Phase 1 (local, EXT4) and Phase 2 (remote, WekaFS).

+-------------------------------------------------------------------+
|                  Measurement Harness                              |
|                                                                   |
|  +---------------------+   +-----------------------------------+  |
|  | Workload Driver     |-->| Workload Loader                   |  |
|  | (FIO 3.26 +         |   |  Synthetic: 4K..256K random read  |  |
|  |  GDSIO,             |   |   (1 TB read IOs, 100% random)    |  |
|  |  DALI+TensorFlow)   |   |  Real:      ResNet-50 v1.5 train  |  |
|  | metrics: IOPS,      |   |   (ImageNet 2012, BS 4..256)      |  |
|  |  GiB/s, usec, CPU%  |   +-----------------+-----------------+  |
|  +---------------------+                     |                    |
|              |                               v                    |
|              v          +---------------------------------------+ |
|  +--------------------+ | Path-Switch Layer (the A/B knob)      | |
|  | Precondition stage | |                                       | |
|  |  - Seq precond     | |   GDS arm:  cuda_io = cufile          | |
|  |    (128K writes)   | |     O_DIRECT, direct=1, libcufile     | |
|  |  - Rand precond    | |                                       | |
|  |    (4K writes)     | |   Baseline: cuda_io = posix           | |
|  |  -> SSD steady     | |     pread + cudaMemcpyHostToDevice    | |
|  |     state          | +-------------------+-------------------+ |
|  +--------------------+                     |                     |
|                                             v                     |
|  +-----------------------------------------------------------+    |
|  | I/O Engine + GDS Stack                                     |    |
|  |   libcufile / cuFile API (engine=libcufile)                |    |
|  |   nvidia-fs.ko kernel driver  (GDS device path)            |    |
|  |   CUDA 11.4 toolkit                                         |    |
|  +-----------------------------------------------------------+    |
|              |                                                    |
|              v                                                    |
|  +-----------------------------------------------------------+    |
|  | Storage Backend (one of two)                               |    |
|  |   Phase 1 LOCAL : Samsung PM9A3 NVMe, EXT4, PCIe Gen4 x4    |    |
|  |   Phase 2 REMOTE: WekaFS over 100 GbE, 6 targets x 4 drives |    |
|  +-----------------------------------------------------------+    |
|              |                                                    |
|              v                                                    |
|  +-----------------------------------------------------------+    |
|  | Result Aggregator                                          |    |
|  |   (path x size/batch) -> Fig. 4 (local), Fig. 5 (remote)   |    |
|  |   4 panels each: CPU cores, IOPS, latency, throughput      |    |
|  +-----------------------------------------------------------+    |
+-------------------------------------------------------------------+
^ Fig 1: Measurement harness — a paired A/B comparator. The only
  swept primary variable is the data path (cufile vs posix); the SSD,
  GPU, filesystem and workload are held fixed within each cell.

The harness has two unusual properties worth flagging. First, it preconditions the SSD twice (a sequential 128K-granularity fill, then a 4K random-write fill) before every measurement so that the drive is in steady state rather than fresh-out-of-box (FOB) — the authors report ~3.365 GB/s sequential-write and ~199K random-write IOPS at steady state, matching Samsung's published PM9A3 spec. Without this, GDS-vs-baseline numbers would be polluted by the SSD's write-cliff behavior. Second, it deliberately uses two completely different benchmark front-ends for the two workload classes: GDSIO (NVIDIA's in-house tool built on libcufile) for synthetic random reads, and DALI + TensorFlow for the real ResNet-50 pipeline. This means "synthetic vs real" is also implicitly "GDSIO vs DALI"; the authors accept that conflation in order to use each tool's production-grade integration.

Methodology specifics extracted verbatim:

The metric set is broader than a pure-bandwidth storage benchmark: by reporting USR and SYS CPU cores separately the harness can attribute savings to the elimination of the bounce-buffer copy (a SYS/kernel cost) rather than to user-space changes. That USR/SYS split is the single most diagnostic feature of the instrument.

2. System-Under-Test Architecture (GDS vs the bounce-buffer baseline)

The specimen is NVIDIA GPUDirect Storage (GDS) — a direct DMA path between an NVMe device and GPU memory that bypasses a staging copy through host (CPU) DRAM. The paper's whole contribution is to quantify GDS against the path it replaces, so the two paths must be drawn side by side.

2.1 The two data paths

  (a) TRADITIONAL PATH (cuda_io = posix)   "bounce buffer"

   +--------+   (1) DMA    +-----------+  (2) cudaMemcpy   +--------+
   |  NVMe  | ===========> |  CPU /    | ================> |  GPU   |
   |  SSD   |  pread()     |  RAM      |  HostToDevice     | memory |
   +--------+              | (bounce   |                   +--------+
        ^                  |  buffer)  |
        |                  +-----------+
   PCIe switch                  ^
   routes bytes UP to CPU,      |  CPU cycles burned staging the copy;
   then back DOWN to GPU        |  kernel page-cache + read-ahead engaged

  (b) GDS PATH (cuda_io = cufile, O_DIRECT, direct=1)

   +--------+   (1) direct DMA over PCIe / fabric    +--------+
   |  NVMe  | ====================================> |  GPU   |
   |  SSD   |   cuFileRead(fh, d_buf, size, 0, 0)    | memory |
   +--------+   NO bounce buffer, NO CPU copy        +--------+
        ^
        |
   PCIe switch routes bytes straight ACROSS to the GPU BAR;
   all PCIe devices share one physical address space (BAR addr)
^ Fig 2: The two data paths. (a) burns CPU cycles and DRAM bandwidth on
  a staging copy; (b) DMAs storage->GPU directly, freeing the CPU.

The mechanism that makes (b) possible: in GPUDirect RDMA / GDS, all PCIe devices see the same physical address space and can read/write each other's Base Address Register (BAR) regions. So the SSD's DMA engine can be programmed to target GPU memory directly. The traditional path cannot do this because POSIX pread lands data in host memory, after which cudaMemcpy(HostToDevice) is mandatory to reach the GPU.

2.2 The GDS software stack

+------------------------------------------------+
| Applications on CPU                            |  user app / framework
+------------------------------------------------+
| cuFile API  (libcufile.so)                     |  GDS user API
+------------------------------------------------+
| CUDA                                           |  runtime
+------------------------------------------------+   <== user / kernel line
| nvidia-fs.ko kernel driver                     |  GDS device path
+------------------------------------------------+
| Virtual File System (VFS)                      |  Linux I/O stack
+------------------------------------------------+
| Filesystem driver | Block I/O driver |         |  third-party / OS kernel
| Storage driver                       |         |
+------------------------------------------------+
|        Storage (SSDs) DMA engine               |  hardware DMA
+------------------------------------------------+

  PARALLEL PATH for proprietary distributed FS (e.g. WekaFS):
  cuFile API  -->  Distributed FS (user-space)  -->  Storage DMA
  (no nvidia-fs / VFS patching required -> "Linux compatibility")
^ Fig 3: GDS software stack (paper Fig. 2). nvidia-fs sits in the kernel
  I/O stack for local FS; proprietary distributed FS take a user-space
  shortcut that needs no kernel patching.

This stack has a consequential fork. For a local filesystem (EXT4), GDS routes through the nvidia-fs.ko kernel driver, which hooks into VFS to expose GPU memory buffers to the block/disk I/O layers. For a proprietary distributed filesystem (WekaFS), the vendor implements its own user-space GDS path, so no kernel patching is needed — the paper calls this WekaFS's "Linux compatibility." A third mode exists when no GDS-supported filesystem is present: cuFile falls back to a compatibility mode that internally uses CPU bounce buffers — i.e. it quietly degrades to path (a) while keeping the cuFile API surface.

2.3 The two physical testbeds

  PHASE 1 — LOCAL DAS (Table I)
  +-----------------------------------------------------------+
  | Client: Dell DSS8440                                       |
  |   GPU : 8x NVIDIA Tesla V100S, 32 GB each                  |
  |   CPU : 2x Intel Xeon Platinum 8268, 2.9 GHz,              |
  |         48 (24x2) total cores, 768 GB system memory        |
  |   Drive: Samsung PM9A3 (3.49 TiB) x4 per host,            |
  |          PCIe Gen4 x4, U.2                                 |
  |   OS  : Ubuntu 20.04 focal                                |
  |   FS  : EXT4 (512-byte logical blocks)                     |
  +-----------------------------------------------------------+

  PHASE 2 — REMOTE DISAGGREGATED (Table II, Fig. 3)
  +-----------------------------------------------------------+
  |  Client 1 (Dell DSS8440, 8x V100S)                        |
  |        |                                                   |
  |   100 GbE x1  (NIC: Connect-5 dual-port)                  |
  |        |                                                   |
  |   [ Data network 100 GbE ]                                |
  |        |                                                   |
  |   100 GbE x1                                              |
  |        |                                                   |
  |  Weka Target Cluster: 6 nodes (Dell R730xd)               |
  |    each: 2x Xeon E5-2640 v3 @2.6 GHz (15 cores),          |
  |          128 GB RAM, 4x PM9A3 drives, WekaFS 3.12.0       |
  |    total capacity 37.72 TiB (21.72 TiB allocated)         |
  |    DPDK data plane (kernel-bypass networking)             |
  +-----------------------------------------------------------+
^ Fig 4: The two testbeds. Phase 1 measures GDS where the SSD is local
  over PCIe Gen4; Phase 2 measures it across a 100 GbE fabric to a 6-node
  Weka cluster using kernel-bypass DPDK networking.

The remote testbed is the more demanding case because the bytes must cross a 100 GbE link and a distributed filesystem before reaching GPU memory — yet it is also where GDS has the most CPU work to eliminate, because Weka's DPDK-based data plane already consumes CPU cores for kernel-bypass networking. GDS lets those cores be reclaimed.

3. Design-Space Diagram (axes swept vs axes held fixed)

The experiment sweeps two primary axes per workload class and runs each sweep over two storage backends. The synthetic-read sweep varies the I/O transfer (block) size; the ResNet-50 sweep varies the mini-batch size. Both sweeps cross the data-path axis (GDS vs baseline), which is the variable of interest.

                   DESIGN SPACE (per workload class)
  +---------------------------------------------------------------+
  |                                                               |
  |  Axis 1: DATA PATH (2 levels) -- the variable of interest     |
  |    [ GDS   : cuda_io = cufile, O_DIRECT, direct=1 ]           |
  |    [ POSIX : cuda_io = posix, cudaMemcpy bounce buffer ]      |
  |                                                               |
  |  Axis 2a: SYNTHETIC -> I/O TRANSFER SIZE (7 levels)           |
  |    [4K] [8K] [16K] [32K] [64K] [128K] [256K]                  |
  |    workload: 100% random read, 1 TB read IOs                  |
  |                                                               |
  |  Axis 2b: REAL DL -> MINI-BATCH SIZE (7 levels)               |
  |    [4] [8] [16] [32] [64] [128] [256]                         |
  |    workload: ResNet-50 v1.5, ImageNet 2012, DALI pipeline     |
  |                                                               |
  |  Axis 3: STORAGE BACKEND (2 levels)                           |
  |    [ LOCAL  DAS : PM9A3 / EXT4 / PCIe Gen4 ]  -> Fig. 4       |
  |    [ REMOTE Weka: PM9A3 / WekaFS / 100 GbE ] -> Fig. 5        |
  |                                                               |
  |  Axis 4 (remote only): CLIENT MOUNT CPU CORES (3 levels)      |
  |    [2 cores] [4 cores] [6 cores]  -> WekaFS throughput study  |
  |                                                               |
  |  Held FIXED:                                                  |
  |    - GPU model     : NVIDIA Tesla V100S, 32 GB                |
  |    - SSD model     : Samsung PM9A3, PCIe Gen4 x4              |
  |    - CUDA toolkit  : 11.4                                     |
  |    - Workload type : 100% RANDOM read (synthetic);           |
  |                      ResNet-50 v1.5 (real) -- no other models |
  |    - Read vs write : READ only (no write-path GDS results)    |
  |    - Queue depth   : not swept in reported results            |
  |    - GPU count     : 1 (synthetic), 8 (ResNet local),        |
  |                      3 (ResNet remote)                        |
  |    - SSD state     : steady state (preconditioned twice)      |
  +---------------------------------------------------------------+
^ Fig 5: Design space. The data-path axis is crossed against transfer
  size (synthetic) or batch size (real DL), over two storage backends.
  Everything else -- GPU, SSD, CUDA, workload type -- is pinned.

Three absences scope the study. First, only the read path is characterized; although the abstract mentions reads and writes, all reported figures (Fig. 4, Fig. 5) are read workloads, and the preconditioning writes serve only to put the SSD in steady state. Second, only one real model (ResNet-50) is benchmarked — the future work explicitly names "different MLPerf benchmark models" as not-yet- done. Third, queue depth is not swept in the reported panels even though GDSIO exposes it; the only multi-threading axis that is varied is the remote-storage client mount core count (2/4/6), which controls how many CPU cores feed the Weka data plane.

4. Algorithm / Control Flow Diagrams (the two data paths)

The two I/O paths differ in exactly which system calls issue the transfer. The paper gives both call sequences explicitly; the control flow below makes the divergence point visible.

4.1 The read-issue control flow

  START (application wants file bytes in GPU memory)
       |
       v
  (1) Open the file
       baseline: fd  = open("data_file", O_RDONLY, ...)
       GDS     : fh  = open("data_file", O_RDONLY | O_DIRECT, ...)
       |                                  ^^^^^^^^^
       |                                  O_DIRECT bypasses page cache
       v
  (2) Allocate the destination buffer
       baseline: buf   = malloc(size)        // HOST memory
       GDS     : cudaMalloc(&d_buf, size)    // DEVICE memory
       |
       v
  (3) Is a GDS-supported filesystem present?
       |
       +--- NO ----> COMPATIBILITY MODE: cuFile internally uses
       |             CPU bounce buffers (silently behaves like path a)
       |
       +--- YES ---> continue to direct path
       |
       v
  (4) Issue the read
       baseline: pread(fd, buf, size, 0)            // SSD -> HOST
                 cudaMemcpy(d_buf, buf, size,        // HOST -> DEVICE
                            cudaMemcpyHostToDevice)   // the bounce copy
       GDS     : cuFileRead(fh, d_buf, size, 0, 0)   // SSD -> DEVICE
                 // direct DMA, no host copy, no CPU staging
       |
       v
  (5) Bytes now resident in GPU memory -> compute proceeds
       |
       v
  END
^ Fig 6: Read-issue control flow. The fork is at step (4): the baseline
  needs TWO calls (pread + cudaMemcpy); GDS needs ONE (cuFileRead). Step
  (3) is the silent compatibility-mode fallback that erases GDS's gain.

The decisive design point is at step (4): the baseline path is a two-stage transfer (pread then cudaMemcpy), and the second stage — the HostToDevice copy — is precisely the bounce-buffer cost GDS deletes. The single cuFileRead call collapses both stages into one DMA. The silent hazard is step (3): if no GDS-supported filesystem is configured, cuFile drops into compatibility mode and reinstates the bounce buffer without changing the API, so an unaware caller pays the baseline cost while believing it is on the fast path.

4.2 The real-workload data flow (ResNet-50 + DALI)

  +-----------+  ══ raw image files ══>  +------------------------+
  |  NVMe SSD |                          |  Data path (per fig 2) |
  | (PM9A3)   |  ◄══ (no return) ════════|  GDS: direct to GPU    |
  +-----------+                          |  POSIX: via CPU bounce |
                                         +-----------+------------+
                                                     ║
                                          ══ decoded minibatch ══>
                                                     ║
                                                     v
  +----------------------------------------------------------------+
  | GPU (V100S) -- DALI pipeline                                   |
  |   decompress -> random crop -> resize -> flip  (PREPROCESS)    |
  |        |                                                        |
  |        v                                                        |
  |   ResNet-50 forward + backward (COMPUTE on current minibatch)  |
  |        |    (overlaps with prefetch of NEXT minibatch)         |
  |        v                                                        |
  |   gradient -> weight update                                    |
  +----------------------------------------------------------------+
^ Fig 7: Real-workload data flow. With GDS the storage->GPU edge is a
  direct DMA, and DALI does decompress/crop/resize ON the GPU, so the
  CPU is freed from both the staging copy and the preprocessing.

The pipeline matters to the result: a deep-learning training step prefetches and decodes the next minibatch while the GPU computes on the current one. Because DALI performs decompression, random cropping, resizing, and flipping on the GPU, the GDS path lets the entire ingest+preprocess chain avoid the CPU — which is why CPU-utilization savings are large even when training-throughput gains are modest (the GPU is the bottleneck once data arrives, so freeing the CPU does not always speed the step, but it always frees cores).

5. Quantitative Results — Empirical Findings by Regime

5.1 Local DAS — synthetic random read (Fig. 4 a-d)

Headline for local single-SSD storage, stated in the conclusion: GDS delivers a 14% throughput increase and 12% latency decrease, with a 12% CPU-utilization reduction concentrated at medium-to-high transfer sizes. The small-read regime (4K-16K) actually shows reduced GDS read rates because querying the CUDA runtime for the GPU memory buffer location (pointer metadata) adds per-op overhead that dominates when each op is tiny; the OS read-ahead/prefetch of the non-O_DIRECT baseline can out-run GDS at the smallest sizes.

5.2 Local DAS — ResNet-50 training (Fig. 4 e-h)

The ResNet result is the clearest demonstration of GDS's primary value: it is a CPU-offload technology more than a speed technology. At batch 64+ GDS reclaims roughly 15 CPU cores (about 42% of the ingest+preprocess CPU work) while only improving training throughput by 12-15%. The GPU is already the bottleneck during compute, so freeing the CPU does not multiply throughput — but it does free ~15 cores for other work, which is the real prize in a shared HPC node.

5.3 Remote Weka — synthetic random read (Fig. 5 a-d)

The remote result is much stronger than the local result because the baseline path on a Weka client is doubly CPU-bound: WekaFS already runs a DPDK kernel-bypass data plane that consumes CPU cores, and the bounce buffer adds a second CPU burden on top. GDS removes the bounce copy, so the freed cores can feed the DPDK data plane — and the small-block throughput jumps 2.8x while overall latency improves 45%. The remarkable 256 us floor latency at 4K is attributed to WekaFS bypassing the kernel entirely (DPDK + user-space FS in a Linux container).

A key scaling observation: WekaFS nearly saturates the 100 GbE link at ~11 GiB/s with as few as 2, 4, or 6 client cores, and IOPS jumps from ~400K (2 cores) to ~740K (4 cores) but does not improve at 6 cores — i.e. 4 cores is the sweet spot, and adding more cores or more GPUs does not help random read (it is link-bound, not core-bound, past 4 cores).

5.4 Remote Weka — ResNet-50 training (Fig. 5 e-h)

The remote ResNet run uses 3 GPUs (not all 8), because the authors found that adding more GPUs did not improve random-read throughput on the Weka backend — the storage/network is the bottleneck, so 3 GPUs already achieve the filesystem's maximum. Training-throughput gains mirror the local case (2-8% small batch, 12-15% large batch), but CPU savings are larger (39-45% at batch 32+) for the same reason as §5.3: the Weka client was more CPU-bound to begin with.

5.5 Headline numbers, consolidated

  +-----------------------------------------------------------+
  | LOCAL single NVMe (PM9A3, EXT4):                          |
  |   throughput  : +14%                                       |
  |   latency     : -12%                                       |
  |   CPU util    : -12% (-18% to -21% at 32K..1M)            |
  |                                                            |
  | REMOTE WekaFS (6 nodes, 24 PM9A3, 100 GbE):              |
  |   IOPS (4K)   : 740K to a single client (4 cores)         |
  |   throughput  : 11 GiB/s (link-saturating)                |
  |   small-block : 2.8x throughput advantage (4K)            |
  |   latency     : -45% overall; -88%/-49% at 32K/256K       |
  |   CPU util    : -18% to -21% (2-3x lower) at 32K+         |
  |                                                            |
  | ResNet-50 (both backends):                                |
  |   throughput  : +2-8% (small batch), +12-15% (batch 32+)  |
  |   CPU cores   : -5..10 (batch 4-32), -15 (batch 64+)      |
  |   GPU util    : comparable (no change)                    |
  +-----------------------------------------------------------+
^ Fig 8: Consolidated headline numbers. GDS's gain is small-to-moderate
  on local DAS and large on remote Weka, and is dominated by CPU/latency
  savings rather than raw throughput -- larger for bigger transfers.

6. Configuration-Regime Trade-off Tables

6.1 Data path (GDS vs traditional bounce buffer)

For data-intensive GPU workloads, prefer GDS once transfers reach ~32K and above — below that the CUDA pointer-lookup overhead and the loss of OS read-ahead can erase or invert the benefit, while above it the elimination of the bounce copy wins decisively on both latency and CPU.

6.2 Storage backend (local DAS vs remote Weka)

The GDS benefit is larger on the remote Weka backend because the baseline there was more CPU-bound (DPDK data plane + bounce buffer), so removing the bounce copy frees more cores; and because WekaFS's user-space, kernel-bypass design pairs naturally with GDS's direct DMA.

6.3 Transfer/batch size regime

GDS pays off most at large transfer / large batch sizes, where the per-op fixed costs (pointer lookup, access overhead) are amortized over more bytes, and where the bounce copy would have consumed the most CPU and DRAM bandwidth. At the smallest sizes the latency is dominated by access overhead rather than transfer, so GDS's transfer-path advantage is muted.

7. Bottlenecks & Insights Surfaced by the Measurements

7.1 The bounce buffer is a CPU-and-DRAM tax, not a throughput tax

The most important insight is that GDS's headline benefit is CPU and latency, not raw bandwidth. On local DAS, throughput rose only 14% but CPU fell 12% (and up to 42% on ResNet at batch 64+). The bounce buffer's real cost is the kernel cycles and DRAM bandwidth spent staging the copy — freeing those frees ~15 cores at large batch, which matters far more in a shared node than a 14% throughput bump.

7.2 The small-I/O inversion (the pointer-lookup tax)

At 4K-16K reads on local DAS, GDS is slower than the baseline, because each cuFileRead must query the CUDA runtime for the GPU buffer's memory location (pointer metadata) and because O_DIRECT forfeits the OS read-ahead/page-cache that helps small sequential-ish reads. This is a genuine crossover: the GDS fast path has a fixed per-op cost that only amortizes once the transfer is large enough. Below the crossover, the "slow" buffered path wins.

7.3 Link saturation caps the remote benefit past 4 cores / 3 GPUs

On Weka, throughput plateaus at ~11 GiB/s — the 100 GbE link limit — with as few as 4 client cores, and IOPS does not improve at 6 cores or with more than 3 GPUs. The bottleneck shifts from CPU (which GDS relieves) to the network fabric. This means GDS's value on remote storage is bounded by link bandwidth: it gets you to the link ceiling with fewer cores, but it cannot exceed the ceiling.

7.4 WekaFS's kernel-bypass design is a GDS multiplier

The 256 us floor latency at 4K and the 2-3x CPU savings on Weka are attributed to WekaFS running its data plane in a user-space Linux container with DPDK kernel-bypass networking. Because Weka already avoids kernel-specific dependencies and time-sharing, GDS's user-space direct path composes cleanly with it (no nvidia-fs/VFS patching). The lesson: GDS's benefit is amplified when the surrounding storage stack is also kernel-bypass, because the two technologies remove overlapping CPU/kernel overheads.

7.5 Latency composition flips with transfer size

The measurements reveal that at small block sizes latency is dominated by data-access overhead (metadata, queue, runtime lookup), while beyond 32K it is dominated by data-transfer time (moving the bytes). GDS attacks the transfer component, so its latency advantage grows with size — at 256K remote, latency is 49% lower; at 32K, 88% lower. Below 32K the access component, which GDS does not reduce, dominates.

7.6 The silent compatibility-mode hazard

cuFile's compatibility mode reinstates CPU bounce buffers without changing the API when no GDS-supported filesystem is present. A deployment that believes it is on the GDS path but lacks nvidia-fs or a GDS-aware FS will silently pay the full baseline cost. This is an operational bottleneck of measurement, not of hardware: confirming direct=1 actually took effect (the paper's "verify if set" note) is required to trust any GDS result.

8. Limitations of the Methodology

The most consequential limitation is the read-only, single-model, single-GPU-class scope: the paper establishes a clean baseline for GDS random-read on PM9A3 + V100S, but it does not yet map how the benefit shifts across PCIe generations, GPU classes, write workloads, or other distributed filesystems — all of which the authors flag as future work.

9. Note on NCCL Tuning

GDS and a collective-communication library tune different segments of the same byte journey, and the paper's central finding — that the win is a CPU/latency win that grows with transfer size — rhymes with the small-vs-large message regimes that govern protocol selection in collective libraries. GDS removes the storage-to-GPU staging copy; a collective library removes (or amortizes) the GPU-to-GPU staging copy, and both see their fixed per-operation overhead dominate at small sizes and vanish at large ones. The paper's crossover — buffered/read-ahead wins below ~16K, direct DMA wins above ~32K — is the same shape as a low-latency-vs-bandwidth protocol crossover, and the same diagnostic (split CPU USR/SYS cores to attribute a copy's cost) would let a tuner decide when bypassing a host-staging step is worth its fixed setup cost.

10. Analogy

GDS is a freight rail spur that runs straight from the dock to the factory floor, replacing a depot where every container had to be unloaded onto trucks and re-loaded. In the old arrangement (the bounce-buffer baseline), goods arriving at the port (the NVMe SSD) are trucked to a central warehouse (CPU/host DRAM), unloaded, then re-loaded onto a second truck bound for the factory (the GPU) — two handlings, two crews, and a warehouse full of idling forklifts (CPU cores spent on cudaMemcpy). GDS lays a dedicated rail spur so a container rolls from dock to factory in one move, and the warehouse crew (the CPU) is freed for other work. The catch the paper measures precisely: for a single small parcel the rail spur is slower, because hooking up the locomotive and clearing the track (querying the CUDA runtime for the buffer address, forfeiting the OS read-ahead) costs more than the trucks would have — so the spur only pays off once you are shipping full containers (32K+ transfers, large batches). And on the long-haul remote route (the Weka cluster across 100 GbE), the spur helps most of all, because that route's trucks were already stuck in traffic (the DPDK data plane burning CPU cores) — but no matter how good the spur, you can never move freight faster than the single track allows (the 11 GiB/s link ceiling that 4 cores already saturate). The warehouse, not the spur, was the hidden cost; GDS bills it back to the workload.