hipFile: Limit IO size to MAX_RW_COUNT on fastpath

.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 \

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.

examples/aiscp/aiscp.cpp

@@ -13,6 +13,7 @@
 #include <hipfile.h>
 #include <hip/hip_runtime_api.h>
 
+#include <algorithm>
 #include <cerrno>
 #include <cstdio>
 #include <cstdlib>
@@ -22,6 +23,13 @@
 #include <sys/types.h>
 #include <unistd.h>
 
+/// @brief AISCP_CHUNK_SIZE defines the maximum size of each aiscp IO operation.
+/// 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
+#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,25 @@ 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
+align_up(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[])
 {
@@ -80,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         nbytes;
+    ssize_t         nwrite{}, nread{}, nbytes{};
+    hoff_t          file_offset{};
 
     if (argc != 3) {
         fprintf(stderr, "Usage: %s SOURCE DEST\n", argv[0]);
@@ -98,6 +126,10 @@ main(int argc, char *argv[])
         }
         file_size  = static_cast<size_t>(statbuf.st_size);
         block_size = static_cast<size_t>(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,
@@ -106,42 +138,47 @@ main(int argc, char *argv[])
     }
 
     if (0 == file_size) {
+        exit_status = EXIT_SUCCESS;
         goto close_dst;
     }
 
     if (open_file(src_path, O_RDONLY, 0, &src_fd, &src_handle)) {
         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 = 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);
         goto close_src;
     }
 
-    nbytes = hipFileRead(src_handle, devbuf, buffer_size, 0, 0);
-    if (nbytes < 0 || file_size != static_cast<size_t>(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<size_t>(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, 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));
+            goto free_devbuf;
+        }
 
-    if (file_size < buffer_size) {
-        if (-1 == ftruncate(dst_fd, static_cast<off_t>(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, align_up(static_cast<size_t>(nread - nwrite), block_size),
+                             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));
+                goto free_devbuf;
+            }
+            nwrite += nbytes;
         }
+        file_offset += nread;
+    } while (nread > 0);
+
+    if (-1 == ftruncate(dst_fd, static_cast<off_t>(file_size))) {
+        fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno));
     }
 
     exit_status = EXIT_SUCCESS;

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

src/amd_detail/backend/fastpath.cpp

@@ -177,6 +177,12 @@ Fastpath::io(IoType type, shared_ptr<IFile> file, shared_ptr<IBuffer> 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/kfd 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<Hip>::get()->hipAmdFileRead(handle, devptr, size, file_offset);

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, IoSizeIsTruncatedToMaxRWCount)
+{
+    StrictMock<MHip> 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

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"