From 709f6f8020f59c95cb50519a2df204088cd9768e Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 15 Jan 2026 22:43:32 +0000 Subject: [PATCH 01/13] aiscp: Handle empty source files --- examples/aiscp/aiscp.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 6132be06..02db22b1 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -106,6 +106,7 @@ main(int argc, char *argv[]) } if (0 == file_size) { + exit_status = EXIT_SUCCESS; goto close_dst; } From f93121a558e328a8ac8f9452da63d76a88f3b2ac Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Tue, 20 Jan 2026 21:38:06 +0000 Subject: [PATCH 02/13] hipFile: AIS path: Truncate IO sizes to MAX_RW_COUNT --- src/amd_detail/backend/fastpath.cpp | 6 ++++++ test/amd_detail/fastpath.cpp | 23 +++++++++++++++++++++++ 2 files changed, 29 insertions(+) diff --git a/src/amd_detail/backend/fastpath.cpp b/src/amd_detail/backend/fastpath.cpp index 257bd29e..fd1891f0 100644 --- a/src/amd_detail/backend/fastpath.cpp +++ b/src/amd_detail/backend/fastpath.cpp @@ -177,6 +177,12 @@ Fastpath::io(IoType type, shared_ptr file, shared_ptr buffer, si throw std::invalid_argument("IO could overflow buffer"); } + // Currently, when IO sizes > MAX_RW_COUNT are submitted to amdgpu/kfd an + // Illegal Seek error is returned. To avoid this, hipFile limits IO size to + // MAX_RW_COUNT. When amdgpu/kdf properly handles IO sizes > MAX_RW_COUNT + // this can be removed. + size = std::min(size, MAX_RW_COUNT); + switch (type) { case IoType::Read: nbytes = Context::get()->hipAmdFileRead(handle, devptr, size, file_offset); diff --git a/test/amd_detail/fastpath.cpp b/test/amd_detail/fastpath.cpp index 05efe19b..e365e665 100644 --- a/test/amd_detail/fastpath.cpp +++ b/test/amd_detail/fastpath.cpp @@ -446,6 +446,29 @@ TEST_P(FastpathIoParam, IoDoesNotMaskSystemError) std::system_error); } +// Ensure IO size is truncated to MAX_RW_COUNT +TEST_P(FastpathIoParam, IoSizeIsTrucatedToMaxRWCount) +{ + StrictMock mhip; + + const size_t buffer_size{SIZE_MAX}; + const size_t io_size{SIZE_MAX}; + + expect_io(DEFAULT_UNBUFFERED_FD, DEFAULT_BUFFER_ADDR, buffer_size); + switch (GetParam()) { + case IoType::Read: + EXPECT_CALL(mhip, hipAmdFileRead(_, _, MAX_RW_COUNT, _)).WillOnce(Return(MAX_RW_COUNT)); + break; + case IoType::Write: + EXPECT_CALL(mhip, hipAmdFileWrite(_, _, MAX_RW_COUNT, _)).WillOnce(Return(MAX_RW_COUNT)); + break; + default: + FAIL() << "Invalid IoType"; + } + + ASSERT_EQ(Fastpath().io(GetParam(), mfile, mbuffer, io_size, 0, 0), MAX_RW_COUNT); +} + INSTANTIATE_TEST_SUITE_P(FastpathTest, FastpathIoParam, Values(IoType::Read, IoType::Write)); HIPFILE_WARN_NO_GLOBAL_CTOR_ON From c90e41e92b4e02ad42fbb59000f291a0042e527f Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 15 Jan 2026 22:35:44 +0000 Subject: [PATCH 03/13] aiscp: Handle files larger than the maximum allowed IO size on linux --- examples/aiscp/aiscp.cpp | 68 ++++++++++++++++++++++++++-------------- 1 file changed, 45 insertions(+), 23 deletions(-) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 02db22b1..9342ad49 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -13,6 +13,7 @@ #include #include +#include #include #include #include @@ -22,6 +23,13 @@ #include #include +/// @brief AISCP_CHUNK_SIZE defines the maximum size of each aiscp IO operation. +/// Currently AISCP_CHUNKS_SIZE is defined to match the maximum IO size on Linux +/// (MAX_RW_COUNT). +#ifndef AISCP_CHUNK_SIZE +#define AISCP_CHUNK_SIZE 0x7ffff000LU +#endif + /// @brief Open and register a file /// @param path [in] Path to the file /// @param flags [in] flags: Flags to pass to open (2) @@ -70,6 +78,16 @@ close_file(const char *path, int fd, hipFileHandle_t handle) return 0; } +/// @brief Round value to the next multiple of align. Align _must_ be a power of 2. +/// @param value The value to round up. +/// @param align Value will be rounded up to a multiple of align +/// @return Value rounded up to a multiple of align. +static inline size_t +alignUp(size_t value, size_t align) +{ + return (value + align - 1) & ~(align - 1); +} + int main(int argc, char *argv[]) { @@ -80,7 +98,7 @@ main(int argc, char *argv[]) hipError_t hip_err; int exit_status = EXIT_FAILURE; size_t buffer_size, file_size, block_size; - ssize_t nbytes; + ssize_t ncopy{}, nwrite{}, nread{}, nbytes{}; if (argc != 3) { fprintf(stderr, "Usage: %s SOURCE DEST\n", argv[0]); @@ -114,35 +132,39 @@ main(int argc, char *argv[]) goto close_dst; } - buffer_size = file_size; - // If needed, round buffer_size up to the next multiple of block_size - if (buffer_size & (block_size - 1)) { - buffer_size = (buffer_size + block_size) & ~(block_size - 1); - } - hip_err = hipMalloc(&devbuf, buffer_size); + buffer_size = alignUp(std::min(file_size, AISCP_CHUNK_SIZE), block_size); + hip_err = hipMalloc(&devbuf, buffer_size); if (hipSuccess != hip_err) { fprintf(stderr, "Could not allocate device buffer (%d)", hip_err); goto close_src; } - nbytes = hipFileRead(src_handle, devbuf, buffer_size, 0, 0); - if (nbytes < 0 || file_size != static_cast(nbytes)) { - fprintf(stderr, "Could not read from %s (%zd) (%s)\n", src_path, nbytes, - IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno)); - goto free_devbuf; - } - - nbytes = hipFileWrite(dst_handle, devbuf, buffer_size, 0, 0); - if (nbytes < 0 || buffer_size != static_cast(nbytes)) { - fprintf(stderr, "Could not write to %s (%zd) (%s)\n", src_path, nbytes, - IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno)); - goto free_devbuf; - } + // Copy the file chunk-by-chunk until we hit EOF + do { + nread = hipFileRead(src_handle, devbuf, buffer_size, ncopy, 0); + if (nread < 0) { + fprintf(stderr, "Could not read from %s (%zd) (%s)\n", src_path, nread, + IS_HIPFILE_ERR(nread) ? HIPFILE_ERRSTR(nread) : strerror(errno)); + goto free_devbuf; + } - if (file_size < buffer_size) { - if (-1 == ftruncate(dst_fd, static_cast(file_size))) { - fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno)); + nwrite = 0; + while (nwrite < nread) { + nbytes = + hipFileWrite(dst_handle, devbuf, alignUp(static_cast(nread - nwrite), block_size), + ncopy + nwrite, static_cast(nwrite)); + if (nbytes < 0) { + fprintf(stderr, "Could not write to %s (%zd) (%s)\n", dst_path, nbytes, + IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno)); + goto free_devbuf; + } + nwrite += nbytes; } + ncopy += nread; + } while (nread > 0); + + if (-1 == ftruncate(dst_fd, static_cast(file_size))) { + fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno)); } exit_status = EXIT_SUCCESS; From 1da4d1e7f42d110c0df8504dcefae869dfd4ed40 Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Tue, 20 Jan 2026 22:42:53 +0000 Subject: [PATCH 04/13] review: Fix build error (useless cast) --- examples/aiscp/aiscp.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 9342ad49..9ca4cca6 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -152,7 +152,7 @@ main(int argc, char *argv[]) while (nwrite < nread) { nbytes = hipFileWrite(dst_handle, devbuf, alignUp(static_cast(nread - nwrite), block_size), - ncopy + nwrite, static_cast(nwrite)); + ncopy + nwrite, nwrite); if (nbytes < 0) { fprintf(stderr, "Could not write to %s (%zd) (%s)\n", dst_path, nbytes, IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno)); From 475e00015a3bf1175a700e61c473098f4b2ffcee Mon Sep 17 00:00:00 2001 From: kurtmcmillan Date: Wed, 21 Jan 2026 09:14:43 -0700 Subject: [PATCH 05/13] copilot: fix typo: kdf -> kfd Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- src/amd_detail/backend/fastpath.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/amd_detail/backend/fastpath.cpp b/src/amd_detail/backend/fastpath.cpp index fd1891f0..1238b778 100644 --- a/src/amd_detail/backend/fastpath.cpp +++ b/src/amd_detail/backend/fastpath.cpp @@ -179,7 +179,7 @@ Fastpath::io(IoType type, shared_ptr file, shared_ptr buffer, si // Currently, when IO sizes > MAX_RW_COUNT are submitted to amdgpu/kfd an // Illegal Seek error is returned. To avoid this, hipFile limits IO size to - // MAX_RW_COUNT. When amdgpu/kdf properly handles IO sizes > MAX_RW_COUNT + // MAX_RW_COUNT. When amdgpu/kfd properly handles IO sizes > MAX_RW_COUNT // this can be removed. size = std::min(size, MAX_RW_COUNT); From 7ad5fa7a92a1ae88e0f12daafea2a65bd17f4da2 Mon Sep 17 00:00:00 2001 From: kurtmcmillan Date: Wed, 21 Jan 2026 09:15:17 -0700 Subject: [PATCH 06/13] copilot: fix typo: AISCP_CHUNKS_SIZE -> AISCP_CHUNK_SIZE Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- examples/aiscp/aiscp.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 9ca4cca6..bba1d8b5 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -24,7 +24,7 @@ #include /// @brief AISCP_CHUNK_SIZE defines the maximum size of each aiscp IO operation. -/// Currently AISCP_CHUNKS_SIZE is defined to match the maximum IO size on Linux +/// Currently AISCP_CHUNK_SIZE is defined to match the maximum IO size on Linux /// (MAX_RW_COUNT). #ifndef AISCP_CHUNK_SIZE #define AISCP_CHUNK_SIZE 0x7ffff000LU From 4c94332b353dcf53feb6a5c1d286008ba6d70c16 Mon Sep 17 00:00:00 2001 From: kurtmcmillan Date: Wed, 21 Jan 2026 09:18:17 -0700 Subject: [PATCH 07/13] copilot: fix typo: IoSizeIsTrucatedToMaxRWCount -> IoSizeIsTruncatedToMaxRWCount Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- test/amd_detail/fastpath.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/amd_detail/fastpath.cpp b/test/amd_detail/fastpath.cpp index e365e665..3e4564b0 100644 --- a/test/amd_detail/fastpath.cpp +++ b/test/amd_detail/fastpath.cpp @@ -447,7 +447,7 @@ TEST_P(FastpathIoParam, IoDoesNotMaskSystemError) } // Ensure IO size is truncated to MAX_RW_COUNT -TEST_P(FastpathIoParam, IoSizeIsTrucatedToMaxRWCount) +TEST_P(FastpathIoParam, IoSizeIsTruncatedToMaxRWCount) { StrictMock mhip; From 65bc4f7e097be2b2feab74438bd6e093c1e8818d Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 17:47:50 +0000 Subject: [PATCH 08/13] review: Document hipFileRead()/hipFileWrite() transfer limit --- CHANGELOG.md | 1 + include/hipfile.h | 6 ++++++ 2 files changed, 7 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 7f717aa0..4eb5cb34 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,7 @@ * Renamed `hipFileOpStatusError()` to `hipFileGetOpErrorString()` * The `hipfile-doc` CMake target no longer exists (just use `doc`) * The `doc` CMake target only exists if the `AIS_BUILD_DOCS` option is enabled +* `hipFileRead()`/`hipFileWrite()` will transfer at most 0x7ffff000 (2,147,479,552) bytes returning the number of bytes actually transferred ### Removed * The rocFile library has been completely removed and the code is now a part of hipFile. diff --git a/include/hipfile.h b/include/hipfile.h index 07a2e405..783622e7 100644 --- a/include/hipfile.h +++ b/include/hipfile.h @@ -528,6 +528,9 @@ hipFileError_t hipFileBufDeregister(const void *buffer_base); * @brief Synchronously read data from a file into a GPU buffer * @ingroup file * + * hipFileRead() will transfer at most 0x7ffff000 (2,147,479,552) bytes, + * returning the number of bytes actually transferred. + * * @param [in] fh Opaque file handle for GPU IO operations * @param [in] buffer_base Base address of the GPU buffer * @param [in] size Number of bytes that should be read @@ -546,6 +549,9 @@ ssize_t hipFileRead(hipFileHandle_t fh, void *buffer_base, size_t size, hoff_t f * @brief Synchronously write data from a GPU buffer to a file * @ingroup file * + * hipFileWrite() will transfer at most 0x7ffff000 (2,147,479,552) bytes, + * returning the number of bytes actually transferred. + * * @param [in] fh Opaque file handle for GPU IO operations * @param [in] buffer_base Base address of the GPU buffer * @param [in] size Number of bytes that should be written From 2fed075c6b3609ab3792a3e1c5722c3b5d02e8fd Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 17:48:36 +0000 Subject: [PATCH 09/13] review: Ensure block size is a power of two --- examples/aiscp/aiscp.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index bba1d8b5..679d4b38 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -88,6 +88,15 @@ alignUp(size_t value, size_t align) return (value + align - 1) & ~(align - 1); } +/// @brief Determines if value is a power of two +/// @param value The value to inspect +/// @return True if value is a power of two, false otherwise +static inline bool +is_power_of_two(size_t value) +{ + return (value > 0) && ((value & (value - 1)) == 0); +} + int main(int argc, char *argv[]) { @@ -116,6 +125,10 @@ main(int argc, char *argv[]) } file_size = static_cast(statbuf.st_size); block_size = static_cast(statbuf.st_blksize); + if (!is_power_of_two(block_size)) { + fprintf(stderr, "Blocksize is not a power of two (%zu)", block_size); + goto program_exit; + } } if (open_file(dst_path, O_WRONLY | O_CREAT, S_IWUSR | S_IRUSR | S_IRGRP | S_IROTH, &dst_fd, From 12e868ba835c327debf47b50528262cc963e1e19 Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 17:49:29 +0000 Subject: [PATCH 10/13] review: Rename alignUp -> align_up --- examples/aiscp/aiscp.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 679d4b38..b900c520 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -83,7 +83,7 @@ close_file(const char *path, int fd, hipFileHandle_t handle) /// @param align Value will be rounded up to a multiple of align /// @return Value rounded up to a multiple of align. static inline size_t -alignUp(size_t value, size_t align) +align_up(size_t value, size_t align) { return (value + align - 1) & ~(align - 1); } @@ -145,7 +145,7 @@ main(int argc, char *argv[]) goto close_dst; } - buffer_size = alignUp(std::min(file_size, AISCP_CHUNK_SIZE), block_size); + buffer_size = align_up(std::min(file_size, AISCP_CHUNK_SIZE), block_size); hip_err = hipMalloc(&devbuf, buffer_size); if (hipSuccess != hip_err) { fprintf(stderr, "Could not allocate device buffer (%d)", hip_err); @@ -164,7 +164,7 @@ main(int argc, char *argv[]) nwrite = 0; while (nwrite < nread) { nbytes = - hipFileWrite(dst_handle, devbuf, alignUp(static_cast(nread - nwrite), block_size), + hipFileWrite(dst_handle, devbuf, align_up(static_cast(nread - nwrite), block_size), ncopy + nwrite, nwrite); if (nbytes < 0) { fprintf(stderr, "Could not write to %s (%zd) (%s)\n", dst_path, nbytes, From 312c75333be11a8eec8a1e2a06f696fde4876c80 Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 18:15:26 +0000 Subject: [PATCH 11/13] review: Change "ssize_t ncopy" -> "hoff_t file_offset" --- examples/aiscp/aiscp.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index b900c520..65341a73 100644 --- a/examples/aiscp/aiscp.cpp +++ b/examples/aiscp/aiscp.cpp @@ -107,7 +107,8 @@ main(int argc, char *argv[]) hipError_t hip_err; int exit_status = EXIT_FAILURE; size_t buffer_size, file_size, block_size; - ssize_t ncopy{}, nwrite{}, nread{}, nbytes{}; + ssize_t nwrite{}, nread{}, nbytes{}; + hoff_t file_offset{}; if (argc != 3) { fprintf(stderr, "Usage: %s SOURCE DEST\n", argv[0]); @@ -154,7 +155,7 @@ main(int argc, char *argv[]) // Copy the file chunk-by-chunk until we hit EOF do { - nread = hipFileRead(src_handle, devbuf, buffer_size, ncopy, 0); + nread = hipFileRead(src_handle, devbuf, buffer_size, file_offset, 0); if (nread < 0) { fprintf(stderr, "Could not read from %s (%zd) (%s)\n", src_path, nread, IS_HIPFILE_ERR(nread) ? HIPFILE_ERRSTR(nread) : strerror(errno)); @@ -165,7 +166,7 @@ main(int argc, char *argv[]) while (nwrite < nread) { nbytes = hipFileWrite(dst_handle, devbuf, align_up(static_cast(nread - nwrite), block_size), - ncopy + nwrite, nwrite); + file_offset + nwrite, nwrite); if (nbytes < 0) { fprintf(stderr, "Could not write to %s (%zd) (%s)\n", dst_path, nbytes, IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno)); @@ -173,7 +174,7 @@ main(int argc, char *argv[]) } nwrite += nbytes; } - ncopy += nread; + file_offset += nread; } while (nread > 0); if (-1 == ftruncate(dst_fd, static_cast(file_size))) { From aa28daaf4015e6e20e0d4fd16979cd7a6119da54 Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 18:19:09 +0000 Subject: [PATCH 12/13] ci: Add AIS aiscp test --- .github/workflows/test-ais-system.yml | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/.github/workflows/test-ais-system.yml b/.github/workflows/test-ais-system.yml index 45396d5f..184f119d 100644 --- a/.github/workflows/test-ais-system.yml +++ b/.github/workflows/test-ais-system.yml @@ -166,6 +166,17 @@ jobs: /ais/hipFile/util/ci-aiscp-test.sh \ /ais/hipFile/build/examples/aiscp/aiscp ' + - name: hipFile fastpath/AIS IO test using aiscp + run: | + docker exec \ + -t \ + -w "/mnt/ais-fs/${AIS_CONTAINER_NAME}" \ + "${AIS_CONTAINER_NAME}" \ + /bin/bash -c ' + HIPFILE_ALLOW_COMPAT_MODE=false + /ais/hipFile/util/ci-aiscp-test.sh \ + /ais/hipFile/build/examples/aiscp/aiscp + ' - name: Configure fio run: | docker exec \ From bfc75ad46f7ea39b5d7b06db1b570593c4ea497d Mon Sep 17 00:00:00 2001 From: Kurt McMillan Date: Thu, 22 Jan 2026 18:19:33 +0000 Subject: [PATCH 13/13] ci: Reduce file size of aiscp tests from 1G to 1M Reduce the rate at which we burn NVMe drives on CI systems. --- util/ci-aiscp-test.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/util/ci-aiscp-test.sh b/util/ci-aiscp-test.sh index c9ad7c20..4bcd6471 100755 --- a/util/ci-aiscp-test.sh +++ b/util/ci-aiscp-test.sh @@ -27,7 +27,7 @@ DST="$TMPDIR/rand.dat.cp" echo "Source: $SRC" echo "Dest: $DST" -head --bytes 1G /dev/urandom > "$SRC" +head --bytes 1M /dev/urandom > "$SRC" "$AISCP" "$SRC" "$DST"