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 ```text 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. ```text 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 | Function | Purpose | Perf Impact | | ----------------------------------- | ----------------------------------------- | --------------- | | `cuFileDriverOpen()` | Initialize cuFile driver, negotiate API | One-time cost | | `cuFileDriverClose()` | Shutdown cuFile driver | Cleanup only | | `cuFileDriverGetProperties()` | Query driver version and GDS capabilities | Diagnostic tool | | `cuFileIsGpuDirectStorageCapable()` | Check if platform can support GDS | Diagnostic tool | | `cuFileIsGpuDirectStorageEnabled()` | Check if GDS is currently engaged | Diagnostic tool | ### Buffer Management | Function | Purpose | Perf Impact | | -------------------------- | ------------------------------------------ | ------------------------------- | | `cuFileBufRegister()` | Register GPU buffer for DMA | **Critical** — enables GDS path | | `cuFileBufDeregister()` | Unregister GPU buffer | Cleanup only | | `cuFileBufSetAllocFlags()` | Set allocation flags for registered buffer | Advanced tuning | ### File Handle Management | Function | Purpose | Perf Impact | | -------------------------- | ------------------------------------ | ------------------------------- | | `cuFileHandleRegister()` | Register file descriptor for GDS I/O | **Critical** — enables GDS path | | `cuFileHandleDeregister()` | Unregister file handle | Cleanup only | ### Synchronous I/O | Function | Purpose | Perf Impact | | --------------- | ------------------------------ | --------------------- | | `cuFileRead()` | Read from file into GPU buffer | Blocks calling thread | | `cuFileWrite()` | Write GPU buffer to file | Blocks calling thread | ### Asynchronous I/O (CUDA Streams) | Function | Purpose | Perf Impact | | ------------------------ | ----------------------------------------- | ------------------------------- | | `cuFileReadAsync()` | Read into GPU buffer on a CUDA stream | Overlaps I/O with compute | | `cuFileWriteAsync()` | Write GPU buffer to file on a CUDA stream | Overlaps I/O with compute | | `cuFileStreamRegister()` | Associate a CUDA stream with cuFile | Required before first async I/O | ### Batch I/O | Function | Purpose | Perf Impact | | -------------------------- | ------------------------------------ | ------------------------------------- | | `cuFileBatchIOSetUp()` | Create batch I/O handle | Setup cost, amortized across requests | | `cuFileBatchIOSubmit()` | Submit multiple I/O requests at once | **Critical** — 2-10× throughput gain | | `cuFileBatchIOGetStatus()` | Check completion status of batch | Polling overhead | | `cuFileBatchIOCancel()` | Cancel pending batch requests | Error recovery | | `cuFileBatchIODestroy()` | Tear down batch I/O handle | Cleanup only | ### Error & Diagnostic Utilities | Function | Purpose | | ------------------------ | ----------------------------------------- | | `cuFileGetErrorString()` | Convert error code to human-readable text | | `cuFileOpError()` | Query last error for a specific operation | ## 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." ```bash # 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 | Workload | Recommended IO Size | Batch? | Reasoning | | ----------------------------- | ------------------- | ---------------- | ---------------------------------------------- | | Small random reads (4KB-64KB) | 4KB-64KB | **Yes** | Batch amortizes per-IO cost 2–5× | | Medium sequential reads | 256KB-1MB | Optional | Large enough that per-IO cost is low | | Large sequential reads | 1MB-16MB | No | Single IO saturates PCIe bandwidth | | Checkpoint / model save | 16MB-256MB | No | Max throughput via single large write | | Data pipeline (train+ingest) | 1MB-8MB | Yes (double buf) | Batch N reads while kernel processes batch N-1 | ## 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 ```json { "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 } } } } ``` | Key | Default | Description | Tuning Guidance |