feat: opt-in host-to-device copy for aisio-cpu/posix

README.md

@@ -59,6 +59,7 @@ filperf <device-uri>[,<device-uri>,...] [options]
 | `--max-file-size <n>` | _(required)_ | Max file size in bytes; required for the `aisio-cpu`/`aisio-gpu`/`aisio-p2p` backends to size the upcie heap |
 | `--warmup <n>` | `0` | Un-timed batches to run before starting the measurement window |
 | `--buffered` | off | Disable `O_DIRECT` when using `posix` backend |
+| `--copy-to-gpu` | off | Copy each file from host to device memory after reading (`aisio-cpu` and `posix` only) |
 | `--async` | off | Use async API when using `cufile` backend |
 | `--summary` | off | Print I/O and dataset statistics after completion |
 | `--help` | | Print usage |

include/fil_iter.h

@@ -23,6 +23,7 @@ struct fil_dev {
 	uint32_t io_errors; ///< Reads that failed in the current batch (CPU/P2P path)
 	const char *data_dir;
 	void **buffers;
+	void **gpu_buffers; ///< Per-buffer GPU destinations for the aisio-cpu copy_to_gpu path
 	uint64_t buf;
 	uint32_t n_buffers;
 	uint32_t nsid;

include/libfil.h

@@ -43,6 +43,7 @@ struct fil_opts {
 	bool buffered;       ///< Whether to use O_DIRECT with POSIX
 	bool async;          ///< Whether to use async API with cuFile
 	bool register_bufs;  ///< Register cuFile device buffers with cuFileBufRegister
+	bool copy_to_gpu;    ///< aisio-cpu/posix: copy files host->device after read
 };
 
 /**

meson.build

@@ -1,5 +1,5 @@
 project('fil', 'c', 'cuda',
-  version: '0.3.2',
+  version: '0.3.3',
   meson_version: '>=0.55.0',
   default_options: [
     'c_std=gnu11',

python/filmodule.c

@@ -17,23 +17,32 @@ init(PyObject *self, PyObject *args, PyObject *keywds)
 	struct fil_stats *stats;
 	struct fil_iter *iter;
 	char *dev_uri;
+	int buffered = 0, use_async = 0, register_bufs = 0, copy_to_gpu = 0;
 	int err;
 
 	if (FilIter) {
 		PyErr_SetString(FilError, "A FIL iterator is already initialized");
 		return NULL;
 	}
 
-	static char *kwlist[] = {"dev_uri",    "data_dir",    "mnt",           "backend",
-				 "iosize",     "gpu_nqueues", "max_file_size", "queue_depth",
-				 "batch_size", NULL};
+	// 'async' is a reserved word in Python, so it is exposed under the keyword 'async_'.
+	static char *kwlist[] = {"dev_uri",     "data_dir",    "mnt",           "backend",
+				 "iosize",      "gpu_nqueues", "max_file_size", "queue_depth",
+				 "batch_size",  "buffered",    "async_",        "register_bufs",
+				 "copy_to_gpu", NULL};
 
-	if (!PyArg_ParseTupleAndKeywords(args, keywds, "s|$sssiikii", kwlist, &dev_uri,
+	// Bools are parsed via 'p' into ints, then folded back into opts.
+	if (!PyArg_ParseTupleAndKeywords(args, keywds, "s|$sssiikiipppp", kwlist, &dev_uri,
 					 &opts.data_dir, &opts.mnt, &opts.backend, &opts.iosize,
 					 &opts.gpu_nqueues, &opts.max_file_size, &opts.queue_depth,
-					 &opts.batch_size)) {
+					 &opts.batch_size, &buffered, &use_async, &register_bufs,
+					 &copy_to_gpu)) {
 		return NULL;
 	}
+	opts.buffered = buffered;
+	opts.async = use_async;
+	opts.register_bufs = register_bufs;
+	opts.copy_to_gpu = copy_to_gpu;
 
 	err = fil_init(&iter, &dev_uri, 1, &opts);
 	if (err) {

python/pyproject.toml

@@ -4,4 +4,4 @@ build-backend = "setuptools.build_meta"
 
 [project]
 name = "fil"
-version = "0.3.2"
+version = "0.3.3"

src/cli.c

@@ -46,6 +46,8 @@ print_help(const char *name)
 	fprintf(stderr, "\t --async \t | \t Use the async API when using cuFile\n");
 	fprintf(stderr, "\t --register_bufs | \t Register device buffers with cuFileBufRegister "
 			"(cuFile backend)\n");
+	fprintf(stderr, "\t --copy-to-gpu \t | \t Copy each file from host to device memory after "
+			"reading. Only relevant for backends: 'aisio-cpu' and 'posix'\n");
 	fprintf(stderr, "\t --summary \t | \t Print IO and dataset stats\n");
 	fprintf(stderr, "\t --help \t | \t Print this message\n");
 }
@@ -114,6 +116,8 @@ parse_args(int argc, char *argv[], struct fil_cli_args *args, struct fil_opts *o
 			opts->async = true;
 		} else if (strcmp(argv[i], "--register_bufs") == 0) {
 			opts->register_bufs = true;
+		} else if (strcmp(argv[i], "--copy-to-gpu") == 0) {
+			opts->copy_to_gpu = true;
 		} else if (strcmp(argv[i], "--summary") == 0) {
 			args->summary = true;
 		} else if (strcmp(argv[i], "--help") == 0) {

src/io.c

@@ -198,6 +198,24 @@ _submit_device(struct fil_iter *iter, struct fil_dev *device, uint32_t dev_id)
 	return err;
 }
 
+static int
+_copy_buffers_to_gpu(struct fil_iter *iter, struct fil_dev *device, uint32_t dev_id)
+{
+	uint32_t slot;
+	int err;
+
+	for (uint32_t i = 0; i < device->n_buffers; i++) {
+		slot = i + dev_id * device->n_buffers;
+		err = cudaMemcpy(device->gpu_buffers[i], device->buffers[i],
+				 iter->output->buf_len[slot], cudaMemcpyHostToDevice);
+		if (err) {
+			fprintf(stderr, "Could not copy data to GPU memory, err: %d\n", err);
+			return err;
+		}
+	}
+	return 0;
+}
+
 int
 fil_cpu_submit(struct fil_iter *iter)
 {
@@ -211,6 +229,9 @@ fil_cpu_submit(struct fil_iter *iter)
 		clock_gettime(CLOCK_MONOTONIC_RAW, &start);
 
 		err = _submit_device(iter, device, i);
+		if (!err && iter->opts->copy_to_gpu) {
+			err = _copy_buffers_to_gpu(iter, device, i);
+		}
 
 		clock_gettime(CLOCK_MONOTONIC_RAW, &end);
 		iter->stats->io_time += ELAPSED(start, end);
@@ -446,11 +467,15 @@ fil_file_submit(struct fil_iter *iter)
 				bytes_read += err;
 			} while ((uint64_t)bytes_read != file->size);
 
-			err = cudaMemcpy(buffer, bounce, file->size, cudaMemcpyHostToDevice);
-			if (err) {
-				fprintf(stderr, "Could not copy data to GPU memory, err: %ld\n",
-					err);
-				return err;
+			if (iter->opts->copy_to_gpu) {
+				err = cudaMemcpy(buffer, bounce, file->size,
+						 cudaMemcpyHostToDevice);
+				if (err) {
+					fprintf(stderr,
+						"Could not copy data to GPU memory, err: %ld\n",
+						err);
+					return err;
+				}
 			}
 		}
 		clock_gettime(CLOCK_MONOTONIC_RAW, &end);

src/iter.c

@@ -426,6 +426,15 @@ _alloc(struct fil_iter *iter, uint32_t n_buffers)
 			fprintf(stderr, "Could not allocate array of buffers: %d\n", err);
 			return err;
 		}
+		if (iter->type == FIL_CPU && iter->opts->copy_to_gpu) {
+			device->gpu_buffers = malloc(sizeof(void *) * device->n_buffers);
+			if (!device->gpu_buffers) {
+				err = errno;
+				fprintf(stderr, "Could not allocate array of GPU buffers: %d\n",
+					err);
+				return err;
+			}
+		}
 		for (uint32_t j = 0; j < device->n_buffers; j++) {
 			switch (iter->type) {
 			case FIL_GPU:
@@ -436,6 +445,16 @@ _alloc(struct fil_iter *iter, uint32_t n_buffers)
 			case FIL_CPU:
 				device->buffers[j] =
 					xnvme_buf_alloc(device->dev, iter->buffer_size);
+				if (device->buffers[j] && iter->opts->copy_to_gpu) {
+					err = cudaMalloc(&device->gpu_buffers[j],
+							 iter->buffer_size);
+					if (err) {
+						fprintf(stderr,
+							"Could not allocate gpu_buffers[%d]: %d\n",
+							i, err);
+						return err;
+					}
+				}
 				break;
 			case FIL_FILE:
 				err = cudaMalloc(&device->buffers[j], iter->buffer_size);
@@ -462,7 +481,13 @@ _alloc(struct fil_iter *iter, uint32_t n_buffers)
 				fprintf(stderr, "Could not allocate buffers[%d]: %d\n", i, err);
 				return err;
 			}
-			iter->output->buffers[j + i * device->n_buffers] = device->buffers[j];
+			if (iter->type == FIL_CPU && iter->opts->copy_to_gpu) {
+				iter->output->buffers[j + i * device->n_buffers] =
+					device->gpu_buffers[j];
+			} else {
+				iter->output->buffers[j + i * device->n_buffers] =
+					device->buffers[j];
+			}
 		}
 
 		if (iter->type == FIL_FILE) {
@@ -636,7 +661,11 @@ fil_term(struct fil_iter *iter)
 		case FIL_CPU:
 			for (uint32_t j = 0; j < device->n_buffers; j++) {
 				xnvme_buf_free(device->dev, device->buffers[j]);
+				if (device->gpu_buffers) {
+					cudaFree(device->gpu_buffers[j]);
+				}
 			}
+			free(device->gpu_buffers);
 			xnvme_queue_term(device->queue);
 			break;
 		case FIL_FILE:
@@ -727,6 +756,14 @@ fil_init(struct fil_iter **iter, char **dev_uris, uint32_t n_devs, struct fil_op
 		return EINVAL;
 	}
 
+	if (opts->copy_to_gpu && strcmp(opts->backend, "aisio-cpu") != 0 &&
+	    strcmp(opts->backend, "posix") != 0) {
+		fprintf(stderr, "opts->copy_to_gpu == true is only compatible with the aisio-cpu "
+				"and posix "
+				"backends\n");
+		return EINVAL;
+	}
+
 	if (!opts->max_file_size &&
 	    (strcmp(opts->backend, "aisio-cpu") == 0 || strcmp(opts->backend, "aisio-gpu") == 0 ||
 	     strcmp(opts->backend, "aisio-p2p") == 0)) {
@@ -881,7 +918,8 @@ fil_opts_default()
 				.batch_size = 1,
 				.buffered = false,
 				.async = false,
-				.register_bufs = false};
+				.register_bufs = false,
+				.copy_to_gpu = false};
 
 	return opts;
 }