From c6ec6f292616b306ea2ddd02a971482ab96c0780 Mon Sep 17 00:00:00 2001 From: Karl Bonde Torp Date: Tue, 30 Jun 2026 16:39:52 +0200 Subject: [PATCH 1/2] refactor(io): rename gds backend to cufile The backend is built on the cuFile userspace library, which ships with the CUDA toolkit. True GPUDirect Storage additionally needs the kernel-side nvidia-fs driver; without it cuFile falls back to a host-staged path. Name the backend 'cufile' so we don't promise GDS when only the library is present. Renames the backend string, the fil_cufile_io struct, fil_cufile_async_submit, and related symbols/messages, plus CLI help and README. Signed-off-by: Karl Bonde Torp --- README.md | 14 +++++++------ include/fil_io.h | 4 ++-- include/fil_iter.h | 2 +- include/libfil.h | 2 +- src/cli.c | 6 +++--- src/io.c | 47 ++++++++++++++++++++++--------------------- src/iter.c | 50 +++++++++++++++++++++++----------------------- 7 files changed, 64 insertions(+), 61 deletions(-) diff --git a/README.md b/README.md index bb3d7ca..b807d61 100644 --- a/README.md +++ b/README.md @@ -49,8 +49,8 @@ filperf [,,...] [options] | Option | Default | Description | |---|---|---| | `--data-dir ` | _(none)_ | Root directory name containing class subdirectories | -| `--backend ` | `aisio-cpu` | I/O backend: `aisio-cpu`, `aisio-gpu`, `aisio-p2p`, `posix`, `gds` | -| `--mnt ` | `/mnt` | Mountpoint of the drive (for `posix` and `gds` backends) | +| `--backend ` | `aisio-cpu` | I/O backend: `aisio-cpu`, `aisio-gpu`, `aisio-p2p`, `posix`, `cufile` | +| `--mnt ` | `/mnt` | Mountpoint of the drive (for `posix` and `cufile` backends) | | `--batch-size ` | `1` | Number of files per batch | | `--batches ` | `1` | Number of batches to read | | `--iosize ` | `4096` | Number of bytes per I/O (`aisio-cpu` and `aisio-gpu` only) | @@ -59,7 +59,7 @@ filperf [,,...] [options] | `--max-file-size ` | _(required)_ | Max file size in bytes; required for the `aisio-cpu`/`aisio-gpu`/`aisio-p2p` backends to size the upcie heap | | `--warmup ` | `0` | Un-timed batches to run before starting the measurement window | | `--buffered` | off | Disable `O_DIRECT` when using `posix` backend | -| `--async` | off | Use async API when using `gds` backend | +| `--async` | off | Use async API when using `cufile` backend | | `--summary` | off | Print I/O and dataset statistics after completion | | `--help` | | Print usage | @@ -88,8 +88,10 @@ Dataset stats: Average size of files in the dataset (KiB): 109.862692 ``` -`posix` and `gds` both require a block device path and the device mounted at `--mnt`. -`gds` additionally requires CUDA with cuFile for direct NVMe-to-GPU transfers. +`posix` and `cufile` both require a block device path and the device mounted at `--mnt`. +`cufile` additionally requires CUDA with the cuFile library. Direct NVMe-to-GPU +(GPUDirect Storage) transfers need the kernel-side GDS driver (`nvidia-fs`); without +it, cuFile falls back to a compatibility path that stages through host memory. For these file-level backends, IOPS equals File/s since each file is a single I/O. ``` @@ -114,7 +116,7 @@ Dataset stats: ``` ``` -$ filperf /dev/nvme0n1 --batches 10 --batch-size 1024 --backend gds --mnt /mnt/datasets --data-dir imagenetish --summary +$ filperf /dev/nvme0n1 --batches 10 --batch-size 1024 --backend cufile --mnt /mnt/datasets --data-dir imagenetish --summary Time, Batches, IOPS, MiB/s 1.255441, 4, 3262.598873, 348.563276 2.413140, 8, 3538.085634, 381.345474 diff --git a/include/fil_io.h b/include/fil_io.h index cd8e479..80b0783 100644 --- a/include/fil_io.h +++ b/include/fil_io.h @@ -28,7 +28,7 @@ struct fil_file_io { void *buffer; }; -struct fil_gds_io { +struct fil_cufile_io { CUfileDescr_t *descr; CUfileHandle_t *handle; size_t *expected; @@ -59,6 +59,6 @@ int fil_file_submit(struct fil_iter *iter); int -fil_gds_async_submit(struct fil_iter *iter); +fil_cufile_async_submit(struct fil_iter *iter); #endif diff --git a/include/fil_iter.h b/include/fil_iter.h index 33aae64..abfefc3 100644 --- a/include/fil_iter.h +++ b/include/fil_iter.h @@ -34,7 +34,7 @@ struct fil_iter { struct fil_stats *stats; struct fil_opts *opts; struct fil_output *output; - struct fil_gds_io *gds_io; + struct fil_cufile_io *cufile_io; int (*io_fn)(struct fil_iter *iter); struct fil_time *time; uint64_t buffer_size; diff --git a/include/libfil.h b/include/libfil.h index 112cdd6..9f5266a 100644 --- a/include/libfil.h +++ b/include/libfil.h @@ -41,7 +41,7 @@ struct fil_opts { ///< aisio-p2p backends to size the upcie heap correctly. uint32_t batch_size; ///< The number of files per batch bool buffered; ///< Whether to use O_DIRECT with POSIX - bool async; ///< Whether to use async API with GDS + bool async; ///< Whether to use async API with cuFile }; /** diff --git a/src/cli.c b/src/cli.c index 2d006ca..62ecb98 100644 --- a/src/cli.c +++ b/src/cli.c @@ -26,9 +26,9 @@ print_help(const char *name) "\t \t \t | \t The data-dir should be a name of a directory, not a path\n"); fprintf(stderr, "\t \t \t | \t The name of the data-dir should be unique\n"); fprintf(stderr, "\t --backend \t | \t The backend to use for reading files (aisio-cpu " - "[default], aisio-gpu, aisio-p2p, posix, gds)\n"); + "[default], aisio-gpu, aisio-p2p, posix, cufile)\n"); fprintf(stderr, "\t --mnt \t \t | \t The mountpoint of the drive (default = /mnt). Only " - "relevant for backends: 'posix' and 'gds'\n"); + "relevant for backends: 'posix' and 'cufile'\n"); fprintf(stderr, "\t --iosize \t | \t The number of bytes per I/O (default = 4096). Only " "relevant for backends: 'aisio-cpu' and 'aisio-gpu'\n"); fprintf(stderr, "\t --gpu-nqueues \t | \t The number of GPU queues to create (default = " @@ -43,7 +43,7 @@ print_help(const char *name) fprintf(stderr, "\t --warmup \t | \t The number of un-timed warmup batches to read before " "the timed run (default = 0)\n"); fprintf(stderr, "\t --buffered \t | \t Don't open with O_DIRECT when using POSIX\n"); - fprintf(stderr, "\t --async \t | \t Use the async API when using GDS\n"); + fprintf(stderr, "\t --async \t | \t Use the async API when using cuFile\n"); fprintf(stderr, "\t --summary \t | \t Print IO and dataset stats\n"); fprintf(stderr, "\t --help \t | \t Print this message\n"); } diff --git a/src/io.c b/src/io.c index 9d4dac2..2db4af5 100644 --- a/src/io.c +++ b/src/io.c @@ -370,9 +370,9 @@ fil_file_submit(struct fil_iter *iter) char *prefix, *path; int fd, flags; ssize_t err, bytes_read; - bool is_gds = strcmp(iter->opts->backend, "gds") == 0; + bool is_cufile = strcmp(iter->opts->backend, "cufile") == 0; flags = O_RDONLY; - if (is_gds || !iter->opts->buffered) { + if (is_cufile || !iter->opts->buffered) { flags |= O_DIRECT; } @@ -397,7 +397,7 @@ fil_file_submit(struct fil_iter *iter) strcat(path, file->name); nbytes = file->size; - if (!is_gds && !iter->opts->buffered) { + if (!is_cufile && !iter->opts->buffered) { // POSIX O_DIRECT requires aligned nbytes nbytes = (1 + ((file->size - 1) / xal_blksize)) * xal_blksize; } @@ -409,7 +409,7 @@ fil_file_submit(struct fil_iter *iter) return err; } - if (is_gds) { + if (is_cufile) { memset(&descr, 0, sizeof(CUfileDescr_t)); descr.handle.fd = fd; descr.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; @@ -426,7 +426,7 @@ fil_file_submit(struct fil_iter *iter) clock_gettime(CLOCK_MONOTONIC_RAW, &start); bytes_read = 0; - if (is_gds) { + if (is_cufile) { bytes_read = cuFileRead(fh, buffer, nbytes, 0, 0); if (bytes_read < 0) { @@ -476,11 +476,11 @@ fil_file_submit(struct fil_iter *iter) } int -fil_gds_async_submit(struct fil_iter *iter) +fil_cufile_async_submit(struct fil_iter *iter) { struct xal_inode *dir, *file; struct timespec start, end; - struct fil_gds_io *gds_io; + struct fil_cufile_io *cufile_io; CUfileError_t status; uint32_t buf_id, dev_id; void *buffer; @@ -488,7 +488,7 @@ fil_gds_async_submit(struct fil_iter *iter) int fd, flags, err = 0; off_t offset = 0; - gds_io = iter->gds_io; + cufile_io = iter->cufile_io; clock_gettime(CLOCK_MONOTONIC_RAW, &start); for (uint32_t i = 0; i < iter->opts->batch_size; i++) { dev_id = i % iter->n_devs; @@ -516,20 +516,21 @@ fil_gds_async_submit(struct fil_iter *iter) return err; } - memset(&gds_io->descr[i], 0, sizeof(CUfileDescr_t)); - gds_io->descr[i].handle.fd = fd; - gds_io->descr[i].type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; - status = cuFileHandleRegister(&gds_io->handle[i], &gds_io->descr[i]); + memset(&cufile_io->descr[i], 0, sizeof(CUfileDescr_t)); + cufile_io->descr[i].handle.fd = fd; + cufile_io->descr[i].type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; + status = cuFileHandleRegister(&cufile_io->handle[i], &cufile_io->descr[i]); if (status.err != CU_FILE_SUCCESS) { fprintf(stderr, "Could not register file, err: %d\n", status.err); close(fd); return status.err; } - gds_io->expected[i] = file->size; - gds_io->actual[i] = 0; + cufile_io->expected[i] = file->size; + cufile_io->actual[i] = 0; - status = cuFileReadAsync(gds_io->handle[i], buffer, &gds_io->expected[i], &offset, - &offset, &gds_io->actual[i], gds_io->streams[i]); + status = cuFileReadAsync(cufile_io->handle[i], buffer, &cufile_io->expected[i], + &offset, &offset, &cufile_io->actual[i], + cufile_io->streams[i]); if (status.err != CU_FILE_SUCCESS) { fprintf(stderr, "cuFileReadAsync failed, err: %d\n", status.err); err = status.err; @@ -542,19 +543,19 @@ fil_gds_async_submit(struct fil_iter *iter) clock_gettime(CLOCK_MONOTONIC_RAW, &start); for (uint32_t i = 0; i < iter->opts->batch_size; i++) { - err = cudaStreamSynchronize(gds_io->streams[i]); + err = cudaStreamSynchronize(cufile_io->streams[i]); if (err) { fprintf(stderr, "Could not synchronize CUDA Stream, err: %d\n", err); goto teardown; } - if (gds_io->actual[i] < 0) { - err = gds_io->actual[i]; + if (cufile_io->actual[i] < 0) { + err = cufile_io->actual[i]; fprintf(stderr, "Reading failed, err: %d\n", err); goto teardown; } - if ((size_t)gds_io->actual[i] != gds_io->expected[i]) { + if ((size_t)cufile_io->actual[i] != cufile_io->expected[i]) { fprintf(stderr, "Could not read entire file, expected: %lu, actual: %lu\n", - gds_io->expected[i], gds_io->actual[i]); + cufile_io->expected[i], cufile_io->actual[i]); err = EIO; goto teardown; } @@ -564,8 +565,8 @@ fil_gds_async_submit(struct fil_iter *iter) teardown: for (uint32_t i = 0; i < iter->opts->batch_size; i++) { - fd = gds_io->descr[i].handle.fd; - cuFileHandleDeregister(gds_io->handle[i]); + fd = cufile_io->descr[i].handle.fd; + cuFileHandleDeregister(cufile_io->handle[i]); close(fd); } return err; diff --git a/src/iter.c b/src/iter.c index 2cd816d..3a5dd91 100644 --- a/src/iter.c +++ b/src/iter.c @@ -91,7 +91,7 @@ _xnvme_setup(struct fil_iter *iter, struct fil_dev *device, const char *uri) } else if (strcmp(backend, "posix") == 0) { opts.be = "linux"; iter->type = FIL_FILE; - } else if (strcmp(backend, "gds") == 0) { + } else if (strcmp(backend, "cufile") == 0) { opts.be = "linux"; iter->type = FIL_FILE; status = cuFileDriverOpen(); @@ -539,50 +539,50 @@ _alloc(struct fil_iter *iter, uint32_t n_buffers) } if (iter->opts->async) { - iter->gds_io = malloc(sizeof(struct fil_gds_io)); - if (!iter->gds_io) { + iter->cufile_io = malloc(sizeof(struct fil_cufile_io)); + if (!iter->cufile_io) { err = errno; - fprintf(stderr, "Could not allocate GDS IO struct: %d\n", err); + fprintf(stderr, "Could not allocate cuFile IO struct: %d\n", err); return err; } - iter->gds_io->descr = malloc(sizeof(CUfileDescr_t) * iter->opts->batch_size); - if (!iter->gds_io->descr) { + iter->cufile_io->descr = malloc(sizeof(CUfileDescr_t) * iter->opts->batch_size); + if (!iter->cufile_io->descr) { err = errno; fprintf(stderr, "Could not allocate cuFile descriptors: %d\n", err); return err; } - iter->gds_io->handle = malloc(sizeof(CUfileHandle_t) * iter->opts->batch_size); - if (!iter->gds_io->handle) { + iter->cufile_io->handle = malloc(sizeof(CUfileHandle_t) * iter->opts->batch_size); + if (!iter->cufile_io->handle) { err = errno; fprintf(stderr, "Could not allocate cuFile handles: %d\n", err); return err; } - iter->gds_io->expected = malloc(sizeof(size_t) * iter->opts->batch_size); - if (!iter->gds_io->expected) { + iter->cufile_io->expected = malloc(sizeof(size_t) * iter->opts->batch_size); + if (!iter->cufile_io->expected) { err = errno; fprintf(stderr, "Could not allocate array of expected values: %d\n", err); return err; } - iter->gds_io->actual = malloc(sizeof(ssize_t) * iter->opts->batch_size); - if (!iter->gds_io->actual) { + iter->cufile_io->actual = malloc(sizeof(ssize_t) * iter->opts->batch_size); + if (!iter->cufile_io->actual) { err = errno; fprintf(stderr, "Could not allocate array of actual values: %d\n", err); return err; } - iter->gds_io->streams = malloc(sizeof(cudaStream_t) * iter->opts->batch_size); - if (!iter->gds_io->streams) { + iter->cufile_io->streams = malloc(sizeof(cudaStream_t) * iter->opts->batch_size); + if (!iter->cufile_io->streams) { err = errno; fprintf(stderr, "Could not allocate array of CUDA Streams: %d\n", err); return err; } for (uint32_t i = 0; i < iter->opts->batch_size; i++) { - err = cudaStreamCreateWithFlags(&iter->gds_io->streams[i], + err = cudaStreamCreateWithFlags(&iter->cufile_io->streams[i], cudaStreamNonBlocking); if (err) { fprintf(stderr, "Could not setup CUDA Stream, err: %d\n", err); @@ -656,15 +656,15 @@ fil_term(struct fil_iter *iter) free(device->buffers); free(device); } - if (iter->gds_io) { - free(iter->gds_io->descr); - free(iter->gds_io->handle); - free(iter->gds_io->expected); - free(iter->gds_io->actual); + if (iter->cufile_io) { + free(iter->cufile_io->descr); + free(iter->cufile_io->handle); + free(iter->cufile_io->expected); + free(iter->cufile_io->actual); for (uint32_t i = 0; i < iter->opts->batch_size; i++) { - cudaStreamDestroy(iter->gds_io->streams[i]); + cudaStreamDestroy(iter->cufile_io->streams[i]); } - free(iter->gds_io->streams); + free(iter->cufile_io->streams); } if (iter->data) { free(iter->data->entries); @@ -713,8 +713,8 @@ fil_init(struct fil_iter **iter, char **dev_uris, uint32_t n_devs, struct fil_op return EINVAL; } - if (opts->async && strcmp(opts->backend, "gds") != 0) { - fprintf(stderr, "opts->async == true is only compatible with GDS backend"); + if (opts->async && strcmp(opts->backend, "cufile") != 0) { + fprintf(stderr, "opts->async == true is only compatible with cuFile backend"); return EINVAL; } @@ -808,7 +808,7 @@ fil_init(struct fil_iter **iter, char **dev_uris, uint32_t n_devs, struct fil_op break; case FIL_FILE: if (_iter->opts->async) { - _iter->io_fn = fil_gds_async_submit; + _iter->io_fn = fil_cufile_async_submit; } else { _iter->io_fn = fil_file_submit; } From 010b116bd61ca9a6fe6ef2a7f2d8a8d333919843 Mon Sep 17 00:00:00 2001 From: Karl Bonde Torp Date: Tue, 30 Jun 2026 16:45:07 +0200 Subject: [PATCH 2/2] ver: bump to v0.3.1 Signed-off-by: Karl Bonde Torp --- meson.build | 2 +- python/pyproject.toml | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/meson.build b/meson.build index 554c62a..026cf5b 100644 --- a/meson.build +++ b/meson.build @@ -1,5 +1,5 @@ project('fil', 'c', 'cuda', - version: '0.3.0', + version: '0.3.1', meson_version: '>=0.55.0', default_options: [ 'c_std=gnu11', diff --git a/python/pyproject.toml b/python/pyproject.toml index b44c1ae..485e86c 100644 --- a/python/pyproject.toml +++ b/python/pyproject.toml @@ -4,4 +4,4 @@ build-backend = "setuptools.build_meta" [project] name = "fil" -version = "0.3.0" +version = "0.3.1"