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 \ 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/examples/aiscp/aiscp.cpp b/examples/aiscp/aiscp.cpp index 6132be06..65341a73 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_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(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, @@ -106,6 +138,7 @@ main(int argc, char *argv[]) } if (0 == file_size) { + exit_status = EXIT_SUCCESS; goto close_dst; } @@ -113,35 +146,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 = 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(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, 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(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(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(file_size))) { + fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno)); } exit_status = EXIT_SUCCESS; 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 diff --git a/src/amd_detail/backend/fastpath.cpp b/src/amd_detail/backend/fastpath.cpp index 257bd29e..1238b778 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/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::get()->hipAmdFileRead(handle, devptr, size, file_offset); diff --git a/test/amd_detail/fastpath.cpp b/test/amd_detail/fastpath.cpp index 05efe19b..3e4564b0 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, IoSizeIsTruncatedToMaxRWCount) +{ + 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 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"