Cufile Skill

cuFile / GPUDirect Storage 高性能编程

Core Philosophy

GDS is a DMA shortcut. GPUDirect Storage enables the NVMe controller to DMA data directly into GPU memory via PCIe peer-to-peer (P2P) transfers. The CPU does not touch the data — no bounce buffer, no cudaMemcpy, no kernel-mode copies.

Compatibility mode is THE #1 pitfall. cuFile silently falls back to a CPU bounce buffer (two PCIe hops: NVMe → CPU → GPU) when GDS prerequisites aren't met. Your code still works, but performance is 2–3× worse. Always verify with cuFileDriverGetProperties() and gdscheck before measuring performance.

Registration has cost. Reuse has value. cuFileBufRegister and cuFileHandleRegister pin pages and set up DMA mappings — these are not free. Register once at init, reuse for the lifetime of the workload. Never register/deregister inside a hot loop.

The API is simple — the performance is in the details. cuFile has only ~15 core functions. The hard part is alignment, IO sizing, stream ordering, batch aggregation, and knowing whether GDS is actually engaged.

GPUDirect Storage: 30-Second Architecture

Traditional I/O (no GDS):
  NVMe SSD ──PCIe──► CPU DRAM (bounce buffer) ──PCIe──► GPU DRAM
  ─ Two PCIe hops, CPU memory bandwidth consumed, CPU involved

GPUDirect Storage (GDS enabled):
  NVMe SSD ──PCIe P2P──► GPU DRAM (BAR-mapped)
  ─ Single PCIe hop, zero CPU memory, CPU only submits commands

Compatibility mode (GDS unavailable, fallback):
  NVMe SSD ──PCIe──► CPU DRAM ──PCIe──► GPU DRAM
  ─ Looks the same to your code, but 2–3× worse throughput

How it works under the hood:

1. cuFileBufRegister(devPtr, size) → Pins GPU virtual address range, creates a BAR mapping visible to the PCIe root complex
2. cuFileHandleRegister(fd, path) → Verifies the file is on a GDS-compatible filesystem, creates a GDS file descriptor
3. cuFileRead(file_handle, devPtr, size, offset, 0) → Submits an NVMe read command with the GPU physical address as the data pointer. The NVMe controller DMAs directly to GPU memory.
4. NVMe completion interrupt → cuFile library processes completion → GPU buffer has the data

Key requirement: The NVMe SSD and GPU must be on the same PCIe root complex with PCIe Access Control Services (ACS) disabled for P2P. gdscheck validates this topology.

cuFile API Lifecycle: The 5-Step Pattern

Every cuFile application follows this exact sequence. Get the order wrong = runtime errors.

Step 1: cuFileDriverOpen()
        └── Initialize GDS driver, negotiate API version
Step 2: Allocate & Register GPU Buffers
        ├── cudaMalloc(&devPtr, size)
        └── cuFileBufRegister(devPtr, size, 0)
Step 3: Open & Register File
        ├── fd = open(path, O_DIRECT | O_RDONLY)
        └── cuFileHandleRegister(&file_handle, &descr)
Step 4: Perform I/O
        ├── cuFileRead(file_handle, devPtr, size, offset, 0)
        ├── cuFileWrite(file_handle, devPtr, size, offset, 0)
        ├── cuFileReadAsync(file_handle, devPtr, size, offset, stream)
        └── cuFileWriteAsync(file_handle, devPtr, size, offset, stream)
Step 5: Cleanup (reverse order of setup)
        ├── cuFileHandleDeregister(file_handle)
        ├── close(fd)
        ├── cuFileBufDeregister(devPtr)
        ├── cudaFree(devPtr)
        └── cuFileDriverClose()

Critical ordering rule: Cleanup in strict reverse order. Deregister the handle BEFORE closing the file descriptor. Deregister the buffer BEFORE freeing GPU memory. Otherwise you get CU_FILE_INVALID_HANDLE or worse — silent memory corruption if the GPU frees a page the DMA engine still references.

API Quick Reference

Driver & Lifecycle

Buffer Management

File Handle Management

Synchronous I/O

Asynchronous I/O (CUDA Streams)

Batch I/O

Error & Diagnostic Utilities

Performance Tuning: The Complete Checklist

Before You Start: Verify GDS is Working

This is the single most important step. Many "perf tuning" sessions end with "GDS was never engaged."

# 1. Check GDS toolchain
gdscheck -p              # Platform check: GPU topology, PCIe P2P, ACS state
gdscheck -f /mnt/nvme    # Filesystem check: mount options, O_DIRECT support
gdsio -f /mnt/nvme/test  # Functional test: write + read + verify

# 2. Verify in code
CUfileDrvProps_t props;
cuFileDriverGetProperties(&props);
printf("GDS capable: %s\n", props.is_gds_capable ? "YES" : "NO");
printf("GDS enabled:  %s\n", props.is_gds_enabled ? "YES" : "NO");
printf("Version:      %d.%d\n", props.major, props.minor);

Performance Tuning Checklist (ordered by impact)

• Buffer Alignment — GPU buffers must be 4KB-aligned for GDS. Use cuMemAlloc (not cudaMalloc) for guaranteed alignment. Misaligned buffers silently fall back to compat mode.
• IO Size — Minimum 4KB for GDS path. For high throughput, use IO sizes ≥ 1MB. Small IOs (≤ 64KB) have high per-operation overhead. Batch them.
• File Open Flags — Files MUST be opened with O_DIRECT for the GDS path. Without it, you get compat mode. Also use O_RDONLY or O_WRONLY — O_RDWR has additional locking overhead.
• Buffer Registration Strategy — Register buffers once at initialization, reuse for thousands of IOs. Registration pins GPU pages — pinning is expensive. Unregistration unpins. A hot loop that registers/unregisters is wasting 90%+ of its time in pinning.
• Batch I/O — For workloads with many small IOs, use cuFileBatchIOSetUp + cuFileBatchIOSubmit. Reduces per-IO submission overhead by 2–10×. See references/batch-io.md.
• Async I/O + CUDA Stream Overlap — Use cuFileReadAsync on one stream while your kernel runs on another. This hides I/O latency behind compute. See references/async-io.md.
• CUDA Stream Registration — Call cuFileStreamRegister() once per stream before first async I/O. Don't re-register the same stream.
• Thread to File Mapping — For multi-file workloads, dedicate one CPU thread per file. cuFile handles are not thread-safe for concurrent operations on the same handle.
• NUMA Affinity — CPU threads submitting I/O should be on the NUMA node closest to the PCIe root complex serving the NVMe device. Use numactl --cpunodebind=N or sched_setaffinity().
• cufile.json Tuning — The global configuration file at /etc/cufile.json (or $CUFILE_CONFIG) controls GDS behavior. Key parameters:
  • max_direct_io_size: Maximum IO size for direct GDS path (default: 1GB)
  • max_device_cache_size: Maximum device-internal cache for GDS (default: 128MB)
  • max_device_pinned_mem_size: Maximum device pinned memory (default: unlimited)
  • enable_compat_mode: Allow fallback to compat mode (default: true — turn off to get hard errors instead of silent fallback)
• Filesystem Choice — GDS requires: ext4, xfs, or GPFS with specific mount options. NFS/CIFS/FUSE are NOT supported. Verify with gdscheck -f <mountpoint>.
• PCIe Topology — GPU and NVMe must be on the same PCIe root complex with ACS disabled for P2P. Multi-socket systems often have separate root complexes — if GPU is on socket 0 and NVMe on socket 1, GDS may not work.

IO Size Guidelines

Configuration: cufile.json Reference

The cuFile configuration file controls global GDS behavior. Default location: /etc/cufile.json. Override with environment variable CUFILE_CONFIG.

Complete Key Reference

{
  "properties": {
    "max_device_cache_size": 134217728,
    "max_device_pinned_mem_size": 0,
    "max_direct_io_size": 1073741824,
    "max_device_allocated_mem_size": 0,
    "enable_compat_mode": true,
    "profiling": {
      "enable": false,
      "trace": false
    },
    "fs": {
      "ext4": {
        "max_direct_io_chunk_size": 0
      },
      "xfs": {
        "max_direct_io_chunk_size": 0
      }
    }
  }
}

Per-Filesystem Tuning

{
  "properties": {
    "fs": {
      "ext4": {
        "max_direct_io_chunk_size": 8388608
      },
      "xfs": {
        "max_direct_io_chunk_size": 16777216
      }
    }
  }
}

GDS Debug Configuration

During development, set these to catch silent fallbacks:

{
  "properties": {
    "enable_compat_mode": false,
    "profiling": {
      "enable": true,
      "trace": false
    }
  }
}

With enable_compat_mode: false, operations that can't use GDS will fail with an explicit error instead of silently falling back. This is essential for verifying that your alignment, IO size, and file configuration are correct.

Common Performance Traps

Error Handling

Error Model

cuFile returns CUfileError_t from most functions. Check .err field — 0 means success. Any non-zero value is an error.

CUfileError_t status;
status = cuFileRead(fh, devPtr, size, offset, 0);
if (status.err != 0) {
    fprintf(stderr, "cuFileRead failed: %s (err=%d)\n",
            cuFileGetErrorString(status), status.err);
}

Common Error Codes

Error Handling Pattern

CUfileError_t cuFileCheck(CUfileError_t status, const char *op) {
    if (status.err != 0) {
        fprintf(stderr, "ERROR: %s failed: %s (err=%d)\n",
                op, cuFileGetErrorString(status), status.err);
        exit(EXIT_FAILURE);
    }
    return status;
}

// Usage:
cuFileCheck(cuFileRead(fh, devPtr, size, offset, 0), "cuFileRead");

GDS Status Diagnostic Pattern

Always include this diagnostic block in your initialization:

void check_gds_status() {
    CUfileDrvProps_t props;
    CUfileError_t st = cuFileDriverGetProperties(&props);
    if (st.err != 0) {
        fprintf(stderr, "WARNING: cuFileDriverGetProperties failed: %s\n",
                cuFileGetErrorString(st));
        return;
    }
    printf("cuFile driver: v%d.%d\n", props.major, props.minor);
    printf("GDS capable:  %s\n", props.is_gds_capable ? "YES" : "NO");
    printf("GDS enabled:   %s\n", props.is_gds_enabled ? "YES" : "NO");
    printf("Max IO size:   %lu bytes (%.1f GB)\n",
           props.max_device_direct_io_size,
           props.max_device_direct_io_size / 1e9);
    if (!props.is_gds_capable) {
        printf("\nACTION REQUIRED: GDS not capable on this system.\n");
        printf("  Run: gdscheck -p\n");
    }
    if (props.is_gds_capable && !props.is_gds_enabled) {
        printf("\nACTION REQUIRED: GDS capable but not enabled.\n");
        printf("  Check /etc/cufile.json: enable_compat_mode\n");
    }
}

Integration Patterns

Pattern 1: Double-Buffered Prefetch Pipeline

The most common high-performance pattern. While the GPU kernel processes batch N, cuFile reads batch N+1 into the other buffer.

Time ─────────────────────────────────────────────────►

Stream 1 (I/O):   [Read B0]            [Read B1]            [Read B2]
Stream 2 (Kernel):          [Process B0]         [Process B1]         [Process B2]

Result: I/O latency is hidden behind compute. Throughput = max(IO_bw, compute_bw),
        not IO_bw + compute_bw.

See examples/05-end-to-end-pipeline.cu for complete code.

Pattern 2: Checkpoint/Restart

Model weights live in GPU memory. cuFile writes them directly to NVMe without a CPU copy.

// Save checkpoint: GPU memory → NVMe (no CPU involved)
cuFileWrite(model_fh, model_weights_gpu, model_size, checkpoint_offset, 0);

// Restore checkpoint: NVMe → GPU memory (no CPU involved)
cuFileRead(model_fh, model_weights_gpu, model_size, checkpoint_offset, 0);

Performance: With GDS, 40GB model checkpoint writes at 20+ GB/s on a PCIe Gen4 x16 NVMe. Traditional CPU-bounce path caps at ~12 GB/s (limited by CPU memory bandwidth).

Pattern 3: Multi-GPU Multi-NVMe Striping

For extreme throughput, stripe data across multiple NVMe devices, each serving a dedicated GPU.

GPU 0 ◄──PCIe P2P──► NVMe 0  (stream 0 data)
GPU 1 ◄──PCIe P2P──► NVMe 1  (stream 1 data)
GPU 2 ◄──PCIe P2P──► NVMe 2  (stream 2 data)
GPU 3 ◄──PCIe P2P──► NVMe 3  (stream 3 data)

Key constraint: Each GPU-NVMe pair must be on the SAME PCIe root complex. gdscheck -p reveals the topology.

Pattern 4: Streaming Ingest with In-Place Processing

Data lands in GPU memory via cuFile, then a CUDA kernel processes it in-place — zero copies, zero CPU involvement.

// Step 1: Read data directly into GPU memory
cuFileReadAsync(fh, gpu_buffer, chunk_size, offset, stream1);

// Step 2: No cudaMemcpy needed — data is already on GPU
// Step 3: Launch kernel on stream2 (waits on stream1 completion)
cudaStreamWaitEvent(stream2, done_event);
process_kernel<<<grid, block, 0, stream2>>>(gpu_buffer, chunk_size);

// Step 4: Write results out (if needed)
cuFileWriteAsync(out_fh, gpu_buffer, result_size, out_offset, stream2);

Relationship with NVMe Programming Skill

cuFile is a higher-level abstraction over NVMe. You do NOT construct NVMe commands, manage PRP lists, or ring doorbell registers — cuFile handles all of that internally.

When to use cuFile (this skill) vs NVMe (nvme-programming skill):

What cuFile handles for you (you don't need to know these to use cuFile):

• PRP/SGL construction for GPU physical addresses
• NVMe command submission and completion polling
• Doorbell register management
• Queue pair allocation and management
• PCIe P2P BAR mapping setup

What you still need to understand (covered in this skill):

• O_DIRECT and filesystem requirements
• Buffer alignment and registration
• IO size vs throughput tradeoffs
• Batch submission strategies
• CUDA stream synchronization
• GDS vs compat mode detection

For deep NVMe understanding (when you need to debug at the protocol level), see the nvme-programming skill.

Reference Files

Detailed reference material is organized by topic — load the relevant file when going deep:

Code Examples

Ready-to-compile CUDA C examples using the cuFile API:

Shared helper code in examples/common/:

• cufile_utils.h / cufile_utils.cu — check_gds_available(), check_alignment(), print_error(), measure_bandwidth()

Search Patterns

When working with the NVIDIA cuFile documentation or source:

# Find API function documentation
grep -r "cuFileRead\|cuFileWrite\|cuFileBufRegister" /path/to/cufile-docs/

# Find error codes
grep -r "CU_FILE_\|CUfileError_t" /path/to/cufile-docs/

# Find configuration reference
grep -r "cufile.json\|max_direct_io_size\|max_device_cache_size" /path/to/cufile-docs/

# Find GDS version history
grep -r "GDS\|GPUDirect Storage\|gdscheck\|gdsio" /path/to/cufile-docs/

# Find examples in NVIDIA's repo
grep -r "cuFileRead\|cuFileHandleRegister" /usr/local/cuda/gds/samples/

Quick Reference

One-page cheat sheet covering core types, API layering, error codes, GDS requirements, cufile.json keys, IO size selection, and performance diagnostics. See references/quick-reference.md.