From 8c12c10f12e2516701726bb40cd3f31b83525eb3 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Thu, 7 Aug 2025 15:50:31 -0400 Subject: [PATCH 01/27] Wrap HIP functions needed for async. hipLaunchHostFunc and hipLaunchKernel are needed to submit CPU-run functions and GPU-run functions to the stream. hipDeviceGetAttribute is needed to query the maximum number of threads per block on the GPU. hipStreamGetDevice is used to determine the GPU the stream is associated with. --- src/amd_detail/hip.cpp | 30 ++++++++++++++++++++++++++++++ src/amd_detail/hip.h | 5 +++++ test/amd_detail/mhip.h | 8 ++++++++ 3 files changed, 43 insertions(+) diff --git a/src/amd_detail/hip.cpp b/src/amd_detail/hip.cpp index e891b44a..fbc6a7ef 100644 --- a/src/amd_detail/hip.cpp +++ b/src/amd_detail/hip.cpp @@ -162,4 +162,34 @@ Hip::hipAmdFileWrite(hipAmdFileHandle_t handle, void *devicePtr, uint64_t size, return bytes_written; } +void +Hip::hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void *user_data) const +{ + (void)throwOnHipError(::hipLaunchHostFunc(stream, fn, user_data)); +} + +void +Hip::hipLaunchKernel(const void *function_address, dim3 numBlocks, dim3 dimBlocks, void **args, + size_t sharedMemBytes, hipStream_t stream) const +{ + (void)throwOnHipError( + ::hipLaunchKernel(function_address, numBlocks, dimBlocks, args, sharedMemBytes, stream)); +} + +int +Hip::hipDeviceGetAttribute(hipDeviceAttribute_t attr, int device_id) const +{ + int attr_value; + (void)throwOnHipError(::hipDeviceGetAttribute(&attr_value, attr, device_id)); + return attr_value; +} + +hipDevice_t +Hip::hipStreamGetDevice(hipStream_t stream) const +{ + hipDevice_t device_id; + (void)throwOnHipError(::hipStreamGetDevice(stream, &device_id)); + return device_id; +} + } diff --git a/src/amd_detail/hip.h b/src/amd_detail/hip.h index 11dc0638..8b399d3b 100644 --- a/src/amd_detail/hip.h +++ b/src/amd_detail/hip.h @@ -69,6 +69,11 @@ struct Hip { virtual uint64_t hipAmdFileWrite(hipAmdFileHandle_t handle, void *devicePtr, uint64_t size, int64_t file_offset) const; virtual HipMemAddressRange hipMemGetAddressRange(hipDeviceptr_t dptr) const; + virtual void hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void *user_data) const; + virtual void hipLaunchKernel(const void *function_address, dim3 numBlocks, dim3 dimBlocks, void **args, + size_t sharedMemBytes, hipStream_t stream) const; + virtual int hipDeviceGetAttribute(hipDeviceAttribute_t attr, int device_id) const; + virtual hipDevice_t hipStreamGetDevice(hipStream_t stream) const; struct RuntimeError : public std::runtime_error { hipError_t error; diff --git a/test/amd_detail/mhip.h b/test/amd_detail/mhip.h index 871ac1c5..f8d0d630 100644 --- a/test/amd_detail/mhip.h +++ b/test/amd_detail/mhip.h @@ -40,6 +40,14 @@ struct MHip : Hip { (hipAmdFileHandle_t handle, void *devicePtr, uint64_t size, int64_t file_offset), (const, override)); MOCK_METHOD(HipMemAddressRange, hipMemGetAddressRange, (hipDeviceptr_t dptr), (const, override)); + MOCK_METHOD(void, hipLaunchHostFunc, (hipStream_t stream, hipHostFn_t fn, void *user_data), + (const, override)); + MOCK_METHOD(void, hipLaunchKernel, + (const void *function_address, dim3 numBlocks, dim3 dimBlocks, void **args, + size_t sharedMemBytes, hipStream_t stream), + (const, override)); + MOCK_METHOD(int, hipDeviceGetAttribute, (hipDeviceAttribute_t attr, int device_id), (const, override)); + MOCK_METHOD(hipDevice_t, hipStreamGetDevice, (hipStream_t stream), (const, override)); }; } From 78b34d1d93933366d8b4c089e39fc334cd413401 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Sun, 25 Jan 2026 22:22:43 -0800 Subject: [PATCH 02/27] test: Defer unlinking of file in Tmpfile to destruction time when platform is NVIDIA. cuFile can try to reopen file using the path. --- test/common/test-common.h | 26 +++++++++++++------------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/test/common/test-common.h b/test/common/test-common.h index fdd3a55a..a257446d 100644 --- a/test/common/test-common.h +++ b/test/common/test-common.h @@ -53,32 +53,32 @@ rfill(void *buffer, uint64_t len, uint32_t seed = 97) } struct Tmpfile { - int fd; + int fd; + std::string path; - Tmpfile() + Tmpfile() : Tmpfile{"."} { - char path_template[] = "hipFile.XXXXXX"; - if ((fd = mkstemp(path_template)) == -1) { - throw std::runtime_error("Could not create temporary file"); - } - if (unlink(path_template) == -1) { - throw std::runtime_error("Could not unlink temporary file)"); - } } - Tmpfile(std::string directory) + Tmpfile(std::string directory) : path{directory} { - directory += "/hipFile.XXXXXX"; - if ((fd = mkstemp(directory.data())) == -1) { + path += "/hipFile.XXXXXX"; + if ((fd = mkstemp(path.data())) == -1) { throw std::runtime_error("Could not create temporary file"); } - if (unlink(directory.c_str()) == -1) { + +#ifdef __HIP_PLATFORM_AMD__ + if (unlink(path.c_str()) == -1) { throw std::runtime_error("Could not unlink temporary file)"); } +#endif } ~Tmpfile() { +#ifdef __HIP_PLATFORM_NVIDIA__ + unlink(path.c_str()); +#endif close(fd); } }; From aa15793662baffb4e0b50bdb4da530a6162322d4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Mon, 17 Nov 2025 16:56:04 -0500 Subject: [PATCH 03/27] Add submission_size to AsyncOpFallback. Will be used to make sure that the size hasn't increased when async parameters are bound. --- src/amd_detail/backend/asyncop-fallback.cpp | 4 ++-- src/amd_detail/backend/asyncop-fallback.h | 1 + 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/src/amd_detail/backend/asyncop-fallback.cpp b/src/amd_detail/backend/asyncop-fallback.cpp index 7c1976d2..e4023158 100644 --- a/src/amd_detail/backend/asyncop-fallback.cpp +++ b/src/amd_detail/backend/asyncop-fallback.cpp @@ -42,8 +42,8 @@ AsyncOpFallback::AsyncOpFallback(IoType _io_type, std::shared_ptr _file, size_t *_size, hoff_t *_file_offset, hoff_t *_buffer_offset, ssize_t *_bytes_transferred) : AsyncOp{_io_type, _file, _buffer, _stream, _size, _file_offset, _buffer_offset, _bytes_transferred}, - bytes_transferred_internal{0}, gpu_buffer{buffer->getBuffer()}, bounce_buffer_dev_ptr{nullptr}, - bounce_buffer{nullptr, [](void *addr) { (void)addr; }} + submitted_size{*_size}, bytes_transferred_internal{0}, gpu_buffer{buffer->getBuffer()}, + bounce_buffer_dev_ptr{nullptr}, bounce_buffer{nullptr, [](void *addr) { (void)addr; }} { void *host_ptr = Context::get()->hipHostMalloc(*_size, 0); std::unique_ptr _bounce_buffer{host_ptr, hipHostDeleter}; diff --git a/src/amd_detail/backend/asyncop-fallback.h b/src/amd_detail/backend/asyncop-fallback.h index 64953fe9..c46180d6 100644 --- a/src/amd_detail/backend/asyncop-fallback.h +++ b/src/amd_detail/backend/asyncop-fallback.h @@ -26,6 +26,7 @@ enum class IoType; namespace hipFile { struct AsyncOpFallback : AsyncOp { + size_t submitted_size; ssize_t bytes_transferred_internal; void *const gpu_buffer; void *bounce_buffer_dev_ptr; From 761fcb3bb88b6c4003e37ae9064f3efcd9206ca4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Mon, 17 Nov 2025 16:56:04 -0500 Subject: [PATCH 04/27] hipFile: Add lock to stream. Lock will be taken while adding operations to the stream for each IO so that each IO is queued sequentially. --- src/amd_detail/stream.cpp | 6 ++++++ src/amd_detail/stream.h | 24 ++++++++++++++---------- test/amd_detail/mstream.h | 3 +++ 3 files changed, 23 insertions(+), 10 deletions(-) diff --git a/src/amd_detail/stream.cpp b/src/amd_detail/stream.cpp index 2815b447..af20df31 100644 --- a/src/amd_detail/stream.cpp +++ b/src/amd_detail/stream.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -53,6 +54,11 @@ Stream::pageAligned() const { return page_aligned; } +std::unique_lock +Stream::getLock() +{ + return std::unique_lock{mutex}; +} void StreamMap::registerStream(hipStream_t hip_stream, uint32_t flags) diff --git a/src/amd_detail/stream.h b/src/amd_detail/stream.h index 4762eb9f..8467e58d 100644 --- a/src/amd_detail/stream.h +++ b/src/amd_detail/stream.h @@ -7,6 +7,7 @@ #include #include #include +#include #include template class PassKey; @@ -17,11 +18,12 @@ class IStream { public: virtual ~IStream() = default; - virtual hipStream_t getHipStream() const = 0; - virtual bool fixedBufferOffset() const = 0; - virtual bool fixedFileOffset() const = 0; - virtual bool fixedIOSize() const = 0; - virtual bool pageAligned() const = 0; + virtual hipStream_t getHipStream() const = 0; + virtual bool fixedBufferOffset() const = 0; + virtual bool fixedFileOffset() const = 0; + virtual bool fixedIOSize() const = 0; + virtual bool pageAligned() const = 0; + virtual std::unique_lock getLock() = 0; }; class StreamMap; @@ -30,11 +32,12 @@ class Stream : public IStream { public: virtual ~Stream() override = default; - virtual hipStream_t getHipStream() const override; - virtual bool fixedBufferOffset() const override; - virtual bool fixedFileOffset() const override; - virtual bool fixedIOSize() const override; - virtual bool pageAligned() const override; + virtual hipStream_t getHipStream() const override; + virtual bool fixedBufferOffset() const override; + virtual bool fixedFileOffset() const override; + virtual bool fixedIOSize() const override; + virtual bool pageAligned() const override; + virtual std::unique_lock getLock() override; Stream(const hipStream_t hip_stream, uint32_t flags, const PassKey &k); @@ -49,6 +52,7 @@ class Stream : public IStream { bool fixed_file_offset; bool fixed_io_size; bool page_aligned; + std::mutex mutex; }; class StreamMap { diff --git a/test/amd_detail/mstream.h b/test/amd_detail/mstream.h index 081e94b3..f10675e9 100644 --- a/test/amd_detail/mstream.h +++ b/test/amd_detail/mstream.h @@ -6,6 +6,8 @@ #include "stream.h" +#include + namespace hipFile { class MStream : public IStream { @@ -15,6 +17,7 @@ class MStream : public IStream { MOCK_METHOD(bool, fixedFileOffset, (), (const, override)); MOCK_METHOD(bool, fixedIOSize, (), (const, override)); MOCK_METHOD(bool, pageAligned, (), (const, override)); + MOCK_METHOD(std::unique_lock, getLock, (), (override)); }; } From 3507106033a2b5732ecaa4cfd4ebbd2c8b75278d Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Mon, 17 Nov 2025 16:55:31 -0500 Subject: [PATCH 05/27] hipFile: Put paramsValid check in io.h. --- src/amd_detail/CMakeLists.txt | 1 + src/amd_detail/io.cpp | 30 ++++++++++++++++++++++++++++++ src/amd_detail/io.h | 10 ++++++++++ 3 files changed, 41 insertions(+) create mode 100644 src/amd_detail/io.cpp diff --git a/src/amd_detail/CMakeLists.txt b/src/amd_detail/CMakeLists.txt index 247f4e3f..c3e18a1f 100644 --- a/src/amd_detail/CMakeLists.txt +++ b/src/amd_detail/CMakeLists.txt @@ -19,6 +19,7 @@ set(HIPFILE_SOURCES file-descriptor.cpp hip.cpp hipfile.cpp + io.cpp mountinfo.cpp state.cpp stream.cpp diff --git a/src/amd_detail/io.cpp b/src/amd_detail/io.cpp new file mode 100644 index 00000000..874688ea --- /dev/null +++ b/src/amd_detail/io.cpp @@ -0,0 +1,30 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#include "buffer.h" +#include "io.h" + +namespace hipFile { + +bool +paramsValid(const std::shared_ptr &buffer, size_t size, hoff_t file_offset, hoff_t buffer_offset) +{ + size_t buffer_length = buffer->getLength(); + if (file_offset < 0) { + return false; + } + if (buffer_offset < 0) { + return false; + } + if (buffer_length < static_cast(buffer_offset)) { + return false; + } + if (buffer_length - static_cast(buffer_offset) < size) { + return false; + } + return true; +} + +} diff --git a/src/amd_detail/io.h b/src/amd_detail/io.h index 4ed0cc51..3c1edc64 100644 --- a/src/amd_detail/io.h +++ b/src/amd_detail/io.h @@ -4,11 +4,21 @@ */ #pragma once +#include "hipfile.h" + +#include +#include + namespace hipFile { +class IBuffer; + enum class IoType { Read, Write, }; +bool paramsValid(const std::shared_ptr &buffer, size_t size, hoff_t file_offset, + hoff_t buffer_offset); + } From b377eb5a543782968063846fc5ee8da6042ce446 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 7 Jan 2026 16:09:39 -0500 Subject: [PATCH 06/27] hipFile: Add device_id to Stream. Will be used to verify GPU buffer belongs to same GPU as the stream being used. --- src/amd_detail/stream.cpp | 19 +++++++++++++++---- src/amd_detail/stream.h | 5 ++++- test/amd_detail/mstream.h | 1 + test/amd_detail/stream.cpp | 8 ++++++++ 4 files changed, 28 insertions(+), 5 deletions(-) diff --git a/src/amd_detail/stream.cpp b/src/amd_detail/stream.cpp index af20df31..149e659a 100644 --- a/src/amd_detail/stream.cpp +++ b/src/amd_detail/stream.cpp @@ -3,6 +3,7 @@ * SPDX-License-Identifier: MIT */ #include "context.h" +#include "hip.h" #include "hipfile.h" #include "passkey.h" #include "stream.h" @@ -17,8 +18,10 @@ namespace hipFile { -Stream::Stream(const hipStream_t _hip_stream, uint32_t flags, const PassKey &) - : hip_stream{_hip_stream}, fixed_buf_offset{(flags & HIPFILE_STREAM_FIXED_BUF_OFFSET) != 0}, +Stream::Stream(const hipStream_t _hip_stream, hipDevice_t _device_id, uint32_t flags, + const PassKey &) + : hip_stream{_hip_stream}, device_id{_device_id}, + fixed_buf_offset{(flags & HIPFILE_STREAM_FIXED_BUF_OFFSET) != 0}, fixed_file_offset{(flags & HIPFILE_STREAM_FIXED_FILE_OFFSET) != 0}, fixed_io_size{(flags & HIPFILE_STREAM_FIXED_FILE_SIZE) != 0}, page_aligned{(flags & HIPFILE_STREAM_PAGE_ALIGNED_INPUTS) != 0} @@ -34,6 +37,11 @@ Stream::getHipStream() const { return hip_stream; } +hipDevice_t +Stream::getHipDevice() const +{ + return device_id; +} bool Stream::fixedBufferOffset() const { @@ -67,7 +75,9 @@ StreamMap::registerStream(hipStream_t hip_stream, uint32_t flags) throw std::invalid_argument("Stream is already registered"); } - auto stream = std::shared_ptr(new Stream(hip_stream, flags, PassKey{})); + hipDevice_t device_id = Context::get()->hipStreamGetDevice(hip_stream); + + auto stream = std::shared_ptr(new Stream(hip_stream, device_id, flags, PassKey{})); StreamMap::from_ptr[hip_stream] = stream; } @@ -87,7 +97,8 @@ StreamMap::getStream(hipStream_t hip_stream) { auto itr = StreamMap::from_ptr.find(hip_stream); if (StreamMap::from_ptr.end() == itr) { - return std::shared_ptr(new Stream(hip_stream, 0, PassKey{})); + hipDevice_t device_id = Context::get()->hipStreamGetDevice(hip_stream); + return std::shared_ptr(new Stream(hip_stream, device_id, 0, PassKey{})); } return itr->second; diff --git a/src/amd_detail/stream.h b/src/amd_detail/stream.h index 8467e58d..45e3f2d9 100644 --- a/src/amd_detail/stream.h +++ b/src/amd_detail/stream.h @@ -19,6 +19,7 @@ class IStream { virtual ~IStream() = default; virtual hipStream_t getHipStream() const = 0; + virtual hipDevice_t getHipDevice() const = 0; virtual bool fixedBufferOffset() const = 0; virtual bool fixedFileOffset() const = 0; virtual bool fixedIOSize() const = 0; @@ -33,13 +34,14 @@ class Stream : public IStream { virtual ~Stream() override = default; virtual hipStream_t getHipStream() const override; + virtual hipDevice_t getHipDevice() const override; virtual bool fixedBufferOffset() const override; virtual bool fixedFileOffset() const override; virtual bool fixedIOSize() const override; virtual bool pageAligned() const override; virtual std::unique_lock getLock() override; - Stream(const hipStream_t hip_stream, uint32_t flags, const PassKey &k); + Stream(const hipStream_t hip_stream, hipDevice_t device_id, uint32_t flags, const PassKey &k); private: Stream(const Stream &) = delete; @@ -48,6 +50,7 @@ class Stream : public IStream { Stream &operator=(const Stream &&) = delete; hipStream_t hip_stream; + hipDevice_t device_id; bool fixed_buf_offset; bool fixed_file_offset; bool fixed_io_size; diff --git a/test/amd_detail/mstream.h b/test/amd_detail/mstream.h index f10675e9..c53ed440 100644 --- a/test/amd_detail/mstream.h +++ b/test/amd_detail/mstream.h @@ -13,6 +13,7 @@ namespace hipFile { class MStream : public IStream { public: MOCK_METHOD(hipStream_t, getHipStream, (), (const, override)); + MOCK_METHOD(hipDevice_t, getHipDevice, (), (const, override)); MOCK_METHOD(bool, fixedBufferOffset, (), (const, override)); MOCK_METHOD(bool, fixedFileOffset, (), (const, override)); MOCK_METHOD(bool, fixedIOSize, (), (const, override)); diff --git a/test/amd_detail/stream.cpp b/test/amd_detail/stream.cpp index 18e8da84..9a9b4668 100644 --- a/test/amd_detail/stream.cpp +++ b/test/amd_detail/stream.cpp @@ -60,6 +60,7 @@ struct HipFileStreamValidParams TEST_P(HipFileStreamValidParams, register_stream_valid_flags_internal) { + EXPECT_CALL(mhip, hipStreamGetDevice); stream_map.registerStream(nonnull_stream, flags); auto stream = stream_map.getStream(nonnull_stream); ASSERT_EQ(stream->getHipStream(), nonnull_stream); @@ -69,30 +70,35 @@ INSTANTIATE_TEST_SUITE_P(StreamSuite, HipFileStreamValidParams, hipFileFlagsPowe TEST_F(HipFileStream, get_stream_with_unregistered_stream_works) { + EXPECT_CALL(mhip, hipStreamGetDevice); auto stream = stream_map.getStream(nonnull_stream); ASSERT_EQ(nonnull_stream, stream->getHipStream()); } TEST_F(HipFileStream, register_with_invalid_flags_throws) { + EXPECT_CALL(mhip, hipStreamGetDevice); ASSERT_THROW(stream_map.registerStream(nonnull_stream, HIPFILE_STREAM_FLAGS_MASK + 1), std::invalid_argument); } TEST_F(HipFileStream, register_twice_throws) { + EXPECT_CALL(mhip, hipStreamGetDevice); stream_map.registerStream(nonnull_stream, 0); ASSERT_THROW(stream_map.registerStream(nonnull_stream, 0), std::invalid_argument); } TEST_F(HipFileStream, deregister_with_registered_stream_works) { + EXPECT_CALL(mhip, hipStreamGetDevice); stream_map.registerStream(nonnull_stream, 0); stream_map.deregisterStream(nonnull_stream); } TEST_F(HipFileStream, deregister_twice_throws) { + EXPECT_CALL(mhip, hipStreamGetDevice); stream_map.registerStream(nonnull_stream, 0); stream_map.deregisterStream(nonnull_stream); ASSERT_THROW(stream_map.deregisterStream(nonnull_stream), std::invalid_argument); @@ -110,6 +116,7 @@ TEST(HipFileStreamDestructor, destructor_with_streams_in_use_logs) std::shared_ptr stream; { StreamMap stream_map; + EXPECT_CALL(mhip, hipStreamGetDevice); stream_map.registerStream(reinterpret_cast(1), 0); EXPECT_CALL(msys, syslog); stream = stream_map.getStream(reinterpret_cast(1)); @@ -128,6 +135,7 @@ struct HipFileStreamExternal : public HipFileOpened { TEST_F(HipFileStreamExternal, register_and_deregister_with_valid_stream_works) { + EXPECT_CALL(mhip, hipStreamGetDevice); ASSERT_EQ(hipFileStreamRegister(nonnull_stream, 0), HIPFILE_SUCCESS); ASSERT_EQ(hipFileStreamDeregister(nonnull_stream), HIPFILE_SUCCESS); } From f77f49c57b5960bfd1a8e2e285f57d0cbf64028c Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 7 Jan 2026 16:16:44 -0500 Subject: [PATCH 07/27] hipFile: Add async memcpy kernel. Copies data between CPU bounce buffer and GPU buffer. --- src/amd_detail/CMakeLists.txt | 1 + src/amd_detail/backend/memcpy-kernel.h | 13 +++ src/amd_detail/backend/memcpy-kernel.hip | 103 +++++++++++++++++++++++ 3 files changed, 117 insertions(+) create mode 100644 src/amd_detail/backend/memcpy-kernel.h create mode 100644 src/amd_detail/backend/memcpy-kernel.hip diff --git a/src/amd_detail/CMakeLists.txt b/src/amd_detail/CMakeLists.txt index c3e18a1f..75b28d2e 100644 --- a/src/amd_detail/CMakeLists.txt +++ b/src/amd_detail/CMakeLists.txt @@ -8,6 +8,7 @@ set(HIPFILE_SOURCES "${HIPFILE_SRC_COMMON_PATH}/hipfile-common.cpp" async.cpp backend/asyncop-fallback.cpp + backend/memcpy-kernel.hip backend/fallback.cpp backend/fastpath.cpp batch/batch.cpp diff --git a/src/amd_detail/backend/memcpy-kernel.h b/src/amd_detail/backend/memcpy-kernel.h new file mode 100644 index 00000000..8261df72 --- /dev/null +++ b/src/amd_detail/backend/memcpy-kernel.h @@ -0,0 +1,13 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ +#pragma once + +#include "backend/asyncop-fallback.h" + +#include + +namespace hipFile { +__global__ void hipFileMemcpyKernel(AsyncOpFallback *op); +} diff --git a/src/amd_detail/backend/memcpy-kernel.hip b/src/amd_detail/backend/memcpy-kernel.hip new file mode 100644 index 00000000..4f050e35 --- /dev/null +++ b/src/amd_detail/backend/memcpy-kernel.hip @@ -0,0 +1,103 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#include "hipfile.h" +#include "io.h" +#include "memcpy-kernel.h" + +#include +#include + +namespace hipFile { + +__global__ void +hipFileMemcpyKernel(AsyncOpFallback *op) +{ + const size_t start_idx = threadIdx.x + blockDim.x * blockIdx.x; + const ssize_t bytes_transferred_internal = op->bytes_transferred_internal; + + if (bytes_transferred_internal < 0) + return; + + IoType io_type = op->io_type; + + size_t *size_p = std::get_if(&op->size); + if (!size_p) { + if (start_idx == 0) { + op->bytes_transferred_internal = -hipFileInvalidValue; + } + return; + } + const size_t size = io_type == IoType::Read ? static_cast(bytes_transferred_internal) : *size_p; + + const hoff_t *buffer_offset_p = std::get_if(&op->buffer_offset); + if (!buffer_offset_p) { + if (start_idx == 0) { + op->bytes_transferred_internal = -hipFileInvalidValue; + } + return; + } + const hoff_t buffer_offset = *buffer_offset_p; + + size_t *const cpu_buffer = reinterpret_cast(op->bounce_buffer_dev_ptr); + if (!cpu_buffer) { + if (start_idx == 0) { + op->bytes_transferred_internal = -hipFileInvalidValue; + } + return; + } + + void *const gpu_buffer = op->gpu_buffer; + if (!gpu_buffer) { + if (start_idx == 0) { + op->bytes_transferred_internal = -hipFileInvalidValue; + } + return; + } + size_t *const gpu_buffer_region = reinterpret_cast(reinterpret_cast(op->gpu_buffer) + + static_cast(buffer_offset)); + const size_t num_chunks = size / sizeof(size_t); + + switch (io_type) { + case IoType::Read: + for (size_t idx = start_idx; idx < num_chunks; idx += blockDim.x) { + gpu_buffer_region[idx] = cpu_buffer[idx]; + } + break; + case IoType::Write: + for (size_t idx = start_idx; idx < num_chunks; idx += blockDim.x) { + cpu_buffer[idx] = gpu_buffer_region[idx]; + } + break; + default: + if (start_idx == 0) + op->bytes_transferred_internal = -hipFileInvalidValue; + return; + } + + if (start_idx == 0) { + uint8_t *const cpu_buffer_bytes = reinterpret_cast(cpu_buffer); + uint8_t *const gpu_buffer_region_bytes = reinterpret_cast(gpu_buffer_region); + const size_t remaining_start = num_chunks * sizeof(size_t); + + switch (io_type) { + case IoType::Read: + for (size_t idx = remaining_start; idx < size; idx++) { + gpu_buffer_region_bytes[idx] = cpu_buffer_bytes[idx]; + } + break; + case IoType::Write: + for (size_t idx = remaining_start; idx < size; idx++) { + cpu_buffer_bytes[idx] = gpu_buffer_region_bytes[idx]; + } + break; + default: + op->bytes_transferred_internal = -hipFileInvalidValue; + return; + } + } +} + +} From 1b49f0a93f790d47c33b95899149dc1a73de5229 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 20 Aug 2025 22:27:39 -0400 Subject: [PATCH 08/27] hipFile: Add fallback async operation. Each IO will have several functions enqueued into the stream. 1. Host function binding IO parameters (if any parameter not fixed). If read: 2. Host function reading from file to CPU bounce buffer. 3. GPU function reading from CPU bounce buffer into GPU buffer. If write: 2. GPU function writing from GPU buffer into CPU bounce buffer. 3. Host function writing from CPU bounce buffer to file. 4. Host function to cleanup completed IO data. --- src/amd_detail/backend/fallback.cpp | 165 ++++++++++++++++++++++++++++ src/amd_detail/backend/fallback.h | 12 ++ src/amd_detail/hipfile.cpp | 75 +++++++------ src/amd_detail/util.h | 46 ++++++++ 4 files changed, 267 insertions(+), 31 deletions(-) create mode 100644 src/amd_detail/util.h diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index 379b1a25..cb0d4e08 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -3,23 +3,32 @@ * SPDX-License-Identifier: MIT */ +#include "async.h" #include "backend.h" #include "buffer.h" +#include "backend/asyncop-fallback.h" +#include "backend/memcpy-kernel.h" #include "context.h" #include "fallback.h" #include "file.h" #include "hip.h" +#include "hipfile.h" #include "io.h" #include "sys.h" +#include "stream.h" +#include "util.h" #include #include #include #include #include +#include #include #include +#include #include +#include using namespace hipFile; @@ -121,3 +130,159 @@ Fallback::io(IoType io_type, shared_ptr file, shared_ptr buffer, return total_io_bytes; } + +void +Fallback::async_io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t *size_p, + hoff_t *file_offset_p, hoff_t *buffer_offset_p, ssize_t *bytes_transferred_p, + std::shared_ptr stream) +{ + if (!paramsValid(buffer, *size_p, *file_offset_p, *buffer_offset_p)) { + throw std::invalid_argument("The selected file or buffer region is invalid"); + } + + if (buffer->getGpuId() != stream->getHipDevice()) { + throw std::invalid_argument("Buffer GPU ID does not match Stream GPU ID"); + } + + *bytes_transferred_p = 0; + + if (*size_p == 0) { + return; + } + + auto op = std::shared_ptr(new AsyncOpFallback( + type, file, buffer, stream, size_p, file_offset_p, buffer_offset_p, bytes_transferred_p)); + Context::get()->addOp(op); + auto op_dev_ptr = op->devPtr(); + void *kernel_args[1] = {&op_dev_ptr}; + + try { + int max_threads_per_block = Context::get()->hipDeviceGetAttribute( + hipDeviceAttributeMaxThreadsPerBlock, buffer->getGpuId()); + auto stream_lock = stream->getLock(); + + // Launch a host function to bind parameters if anything not fixed + if (!op->stream->fixedBufferOffset() || !op->stream->fixedFileOffset() || + !op->stream->fixedIOSize()) { + Context::get()->hipLaunchHostFunc(op->stream->getHipStream(), async_io_bind_params, + op.get()); + } + + switch (op->io_type) { + case IoType::Read: + Context::get()->hipLaunchHostFunc(op->stream->getHipStream(), async_io_cpu_copy, + op.get()); + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), + kernel_args, 0, op->stream->getHipStream()); + break; + case IoType::Write: + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), + kernel_args, 0, op->stream->getHipStream()); + Context::get()->hipLaunchHostFunc(op->stream->getHipStream(), async_io_cpu_copy, + op.get()); + break; + default: + throw std::runtime_error("Invalid IO type"); + } + + Context::get()->hipLaunchHostFunc(op->stream->getHipStream(), async_io_cleanup, op.get()); + } + catch (...) { + // If something threw, still try to enqueue cleanup function + try { + Context::get()->hipLaunchHostFunc(op->stream->getHipStream(), async_io_cleanup, op.get()); + } + catch (...) { + Context::get()->syslog(LOG_CRIT, + "Unable to enqueue async cleanup function. This will leak memory."); + } + throw; + } +} + +extern "C" { +void +async_io_bind_params(void *userargs) +{ + auto op = static_cast(userargs); + // Bind params. Will maintain same value if already bound. + const hoff_t *buffer_offset = get_variant_ptr(op->buffer_offset); + op->buffer_offset.emplace(*buffer_offset); + const hoff_t *file_offset = get_variant_ptr(op->file_offset); + op->file_offset.emplace(*file_offset); + const size_t *size = get_variant_ptr(op->size); + op->size = *size; + + if (std::get(op->size) > op->submitted_size) { + op->bytes_transferred_internal = -hipFileInvalidValue; + return; + } + if (!paramsValid(op->buffer, std::get(op->size), std::get(op->file_offset), + std::get(op->buffer_offset))) { + op->bytes_transferred_internal = -hipFileInvalidValue; + } +} + +void +async_io_cleanup(void *userargs) +{ + auto op = static_cast(userargs); + ssize_t *bytes_transferred = op->bytes_transferred; + try { + Context::get()->completeOp(op); + } + catch (std::invalid_argument &e) { + *bytes_transferred = -hipFileInternalError; + return; + } + *bytes_transferred = op->bytes_transferred_internal; +} + +void +async_io_cpu_copy(void *userargs) +{ + auto op = static_cast(userargs); + size_t bytes_transferred = 0; + const size_t size = std::get(op->size); + const hoff_t file_offset = std::get(op->file_offset); + ssize_t ret = 0; + + if (op->bytes_transferred_internal < 0) { + return; + } + + try { + while (bytes_transferred < size) { + void *cur_buf_position = reinterpret_cast( + reinterpret_cast(op->bounceBufferHostPtr()) + bytes_transferred); + hoff_t cur_file_offset = file_offset + static_cast(bytes_transferred); + size_t remaining_bytes = size - bytes_transferred; + switch (op->io_type) { + case IoType::Read: + ret = Context::get()->pread(op->file->getBufferedFd(), cur_buf_position, + remaining_bytes, cur_file_offset); + break; + case IoType::Write: + ret = Context::get()->pwrite(op->file->getBufferedFd(), cur_buf_position, + remaining_bytes, cur_file_offset); + break; + default: + throw std::runtime_error("Invalid IO type"); + } + if (ret == 0) { + break; + } + bytes_transferred += static_cast(ret); + } + op->bytes_transferred_internal = static_cast(bytes_transferred); + } + catch (std::system_error &e) { + op->bytes_transferred_internal = -1; + } + catch (Hip::RuntimeError &e) { + op->bytes_transferred_internal = -hipFileDriverError; + } +} +} diff --git a/src/amd_detail/backend/fallback.h b/src/amd_detail/backend/fallback.h index b8554f0b..14e3b483 100644 --- a/src/amd_detail/backend/fallback.h +++ b/src/amd_detail/backend/fallback.h @@ -5,6 +5,8 @@ #pragma once +#include "async.h" +#include "backend/asyncop-fallback.h" #include "backend.h" #include "hipfile.h" @@ -32,6 +34,10 @@ struct Fallback : public Backend { ssize_t io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, hoff_t file_offset, hoff_t buffer_offset) override; + void async_io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t *size_p, + hoff_t *file_offset_p, hoff_t *buffer_offset_p, ssize_t *bytes_transferred_p, + std::shared_ptr stream); + // Once we can import gtest.h and make test suites or test friends everything // below here should be made protected. // protected: @@ -41,3 +47,9 @@ struct Fallback : public Backend { }; } + +extern "C" { +void async_io_bind_params(void *userargs); +void async_io_cleanup(void *userargs); +void async_io_cpu_copy(void *userargs); +} diff --git a/src/amd_detail/hipfile.cpp b/src/amd_detail/hipfile.cpp index 7c683c57..75c06757 100644 --- a/src/amd_detail/hipfile.cpp +++ b/src/amd_detail/hipfile.cpp @@ -4,6 +4,7 @@ */ #include "backend.h" +#include "backend/fallback.h" #include "batch/batch.h" #include "buffer.h" #include "context.h" @@ -46,6 +47,16 @@ catch (...) { return {hipFileInternalError, hipSuccess}; } +static void +checkNull(std::initializer_list ptrs) +{ + for (auto ptr : ptrs) { + if (!ptr) { + throw hipFileError_t{hipFileInvalidValue, hipSuccess}; + } + } +} + // Ensures that the driver is initialized without incrementing the reference // count if the driver is already initialized. Used by hipFile to ensure its // behaviour is consistent on AMD and NVIDIA @@ -400,52 +411,54 @@ catch (...) { return; } -hipFileError_t -hipFileReadAsync(hipFileHandle_t fh, void *buffer_base, size_t *size_p, hoff_t *file_offset_p, - hoff_t *buffer_offset_p, ssize_t *bytes_read_p, hipStream_t stream) +static hipFileError_t +hipFileIOAsync(IoType io_type, hipFileHandle_t fh, void *buffer_base, size_t *size_p, hoff_t *file_offset_p, + hoff_t *buffer_offset_p, ssize_t *bytes_transferred_p, hipStream_t hipStream) try { if (Context::get()->getRefCount() == 0) { - // Match cuFile behaviour ensureDriverInit(); return {hipFileInvalidValue, hipSuccess}; } - (void)fh; - (void)buffer_base; - (void)size_p; - (void)file_offset_p; - (void)buffer_offset_p; - (void)bytes_read_p; - (void)stream; + checkNull({buffer_base, size_p, file_offset_p, buffer_offset_p, bytes_transferred_p}); - throw std::runtime_error("Not Implemented"); + auto [file, buffer, stream] = + Context::get()->getFileBufferAndStream(fh, buffer_base, *size_p, 0, hipStream); + Fallback().async_io(io_type, file, buffer, size_p, file_offset_p, buffer_offset_p, bytes_transferred_p, + stream); + + return {hipFileSuccess, hipSuccess}; +} +catch (hipFileError_t e) { + return e; +} +catch (const std::invalid_argument &) { + return {hipFileInvalidValue, hipSuccess}; +} +catch (const std::bad_alloc &) { + return {hipFileHipDriverError, hipErrorOutOfMemory}; +} +catch (const FileNotRegistered &) { + return {hipFileHandleNotRegistered, hipSuccess}; } catch (...) { return handle_exception(); } +hipFileError_t +hipFileReadAsync(hipFileHandle_t fh, void *buffer_base, size_t *size_p, hoff_t *file_offset_p, + hoff_t *buffer_offset_p, ssize_t *bytes_read_p, hipStream_t stream) +{ + return hipFileIOAsync(IoType::Read, fh, buffer_base, size_p, file_offset_p, buffer_offset_p, bytes_read_p, + stream); +} + hipFileError_t hipFileWriteAsync(hipFileHandle_t fh, void *buffer_base, size_t *size_p, hoff_t *file_offset_p, hoff_t *buffer_offset_p, ssize_t *bytes_written_p, hipStream_t stream) -try { - if (Context::get()->getRefCount() == 0) { - // Match cuFile behaviour - ensureDriverInit(); - return {hipFileInvalidValue, hipSuccess}; - } - - (void)fh; - (void)buffer_base; - (void)size_p; - (void)file_offset_p; - (void)buffer_offset_p; - (void)bytes_written_p; - (void)stream; - - throw std::runtime_error("Not Implemented"); -} -catch (...) { - return handle_exception(); +{ + return hipFileIOAsync(IoType::Write, fh, buffer_base, size_p, file_offset_p, buffer_offset_p, + bytes_written_p, stream); } hipFileError_t diff --git a/src/amd_detail/util.h b/src/amd_detail/util.h new file mode 100644 index 00000000..068b16d6 --- /dev/null +++ b/src/amd_detail/util.h @@ -0,0 +1,46 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ +#pragma once + +#include +#include + +namespace hipFile { + +template inline constexpr bool is_variant_of_T_and_T_ptr = false; + +template +inline constexpr bool is_variant_of_T_and_T_ptr> = + (!std::is_pointer_v && std::is_pointer_v && + std::is_same_v, std::remove_pointer_t>) || + (std::is_pointer_v && !std::is_pointer_v && + std::is_same_v, std::remove_const_t>); + +/// @brief Retrieves a pointer to the current set value +/// @param v A variant +/// @details Takes a variant of the form or the reverse. T may be const. Will return a pointer to the +/// non-pointer entry if it is currently held, otherwise it will return the pointer entry. +template >>> +auto +get_variant_ptr(std::variant &v) +{ + if constexpr (!std::is_pointer_v) { + T1 *ptr = std::get_if(&v); + if (!ptr) { + ptr = std::get(v); + } + return ptr; + } + else { + T2 *ptr = std::get_if(&v); + if (!ptr) { + ptr = std::get(v); + } + return ptr; + } +} + +} From 95d7272c15a7abcf39ccb3c051600da76bf9c8e1 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Sat, 10 Jan 2026 21:51:50 -0500 Subject: [PATCH 09/27] hipFile: Add user-defined literals for sizes. --- cmake/AISGNUCompilerOptions.cmake | 3 ++- shared/hipfile-literals.h | 38 +++++++++++++++++++++++++++++++ 2 files changed, 40 insertions(+), 1 deletion(-) create mode 100644 shared/hipfile-literals.h diff --git a/cmake/AISGNUCompilerOptions.cmake b/cmake/AISGNUCompilerOptions.cmake index 3c0aaad0..f122f364 100644 --- a/cmake/AISGNUCompilerOptions.cmake +++ b/cmake/AISGNUCompilerOptions.cmake @@ -52,7 +52,8 @@ function(get_ais_gnu_warning_flags outvar compiler_version) -Wformat-truncation=2 -Wformat-y2k -Winvalid-pch - -Wlong-long + # This is a warning for when using + +constexpr size_t +operator""_KiB(unsigned long long int x) +{ + return 1024ULL * x; +} + +constexpr size_t +operator""_MiB(unsigned long long int x) +{ + return 1024_KiB * x; +} + +constexpr size_t +operator""_GiB(unsigned long long int x) +{ + return 1024_MiB * x; +} + +constexpr size_t +operator""_TiB(unsigned long long int x) +{ + return 1024_GiB * x; +} + +constexpr size_t +operator""_PiB(unsigned long long int x) +{ + return 1024_TiB * x; +} From bbeb326656da8d898514b65e3be9b46afa87f315 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Tue, 6 Jan 2026 17:15:14 -0500 Subject: [PATCH 10/27] Add mock for AsyncMonitor. --- src/amd_detail/async.h | 4 ++-- test/amd_detail/masyncmonitor.h | 24 ++++++++++++++++++++++++ 2 files changed, 26 insertions(+), 2 deletions(-) create mode 100644 test/amd_detail/masyncmonitor.h diff --git a/src/amd_detail/async.h b/src/amd_detail/async.h index aeb9f951..5b565cf3 100644 --- a/src/amd_detail/async.h +++ b/src/amd_detail/async.h @@ -57,8 +57,8 @@ class AsyncMonitor { virtual ~AsyncMonitor(); AsyncMonitor(); - void addOp(std::shared_ptr op); - void completeOp(AsyncOp *op); + virtual void addOp(std::shared_ptr op); + virtual void completeOp(AsyncOp *op); private: void completion_thread(); diff --git a/test/amd_detail/masyncmonitor.h b/test/amd_detail/masyncmonitor.h new file mode 100644 index 00000000..413e11a0 --- /dev/null +++ b/test/amd_detail/masyncmonitor.h @@ -0,0 +1,24 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#pragma once + +#include "async.h" +#include "context.h" + +#include + +namespace hipFile { + +struct MAsyncMonitor : AsyncMonitor { + ContextOverride co; + MAsyncMonitor() : co{this} + { + } + MOCK_METHOD(void, addOp, (std::shared_ptr op), (override)); + MOCK_METHOD(void, completeOp, (AsyncOp * op), (override)); +}; + +} From f89b97b6160552043c1c01131a6bd54bd1422c62 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Tue, 4 Nov 2025 15:48:20 -0500 Subject: [PATCH 11/27] hipFile: Adding async unit tests. --- test/amd_detail/async.cpp | 411 ++++++++++++++++++++++++++++++++++++-- 1 file changed, 398 insertions(+), 13 deletions(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 9a0c18bc..99d09033 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -5,16 +5,25 @@ #include "async.h" #include "backend/asyncop-fallback.h" -#include "hip.h" +#include "backend/fallback.h" +#include "context.h" +#include "file.h" #include "hipfile.h" +#include "hipfile-literals.h" +#include "hipfile-test.h" #include "hipfile-warnings.h" +#include "hip.h" #include "io.h" +#include "masyncmonitor.h" #include "mbuffer.h" #include "mfile.h" #include "mhip.h" +#include "mstate.h" #include "mstream.h" #include "msys.h" +#include "state.h" +#include #include #include #include @@ -23,10 +32,11 @@ #include #include #include -#include +#include #include #include #include +#include // Put tests inside the macros to suppress the global constructor // warnings @@ -45,8 +55,6 @@ using ::testing::Throw; using ::testing::Values; using ::testing::WithParamInterface; -typedef struct stat stat_t; - struct HipFileAsyncOp : public Test { HipFileAsyncOp() : buffer{std::make_shared>()}, file{std::make_shared>()}, @@ -68,10 +76,6 @@ struct HipFileAsyncOp : public Test { shared_ptr stream; }; -struct HipFileAsyncMonitor : HipFileAsyncOp { - AsyncMonitor monitor; -}; - static auto hipfileFlagsPowerSet() { @@ -216,8 +220,8 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_delete_failure_calls_syslog) IoType::Read, file, buffer, stream, &size, &file_offset, &buffer_offset, &bytes_transferred}); } -struct HipFileAsyncOpFallbackFunctions : public HipFileAsyncOp { - HipFileAsyncOpFallbackFunctions() : bounce_buffer{new uint8_t[size]} +struct HipFileAsyncOpFallbackMethods : public HipFileAsyncOp { + HipFileAsyncOpFallbackMethods() : bounce_buffer{new uint8_t[size]} { // make_shared uses placement new, which will not use hipHostMalloc/hipHostFree for // AsyncOpFallback @@ -226,7 +230,7 @@ struct HipFileAsyncOpFallbackFunctions : public HipFileAsyncOp { op = std::make_shared(IoType::Read, file, buffer, stream, &size, &file_offset, &buffer_offset, &bytes_transferred); } - ~HipFileAsyncOpFallbackFunctions() override + ~HipFileAsyncOpFallbackMethods() override { EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); } @@ -239,18 +243,22 @@ struct HipFileAsyncOpFallbackFunctions : public HipFileAsyncOp { std::shared_ptr op; }; -TEST_F(HipFileAsyncOpFallbackFunctions, bounceBufferHostPtr_returns_pointer) +TEST_F(HipFileAsyncOpFallbackMethods, bounceBufferHostPtr_returns_pointer) { ASSERT_EQ(op->bounceBufferHostPtr(), bounce_buffer.get()); } -TEST_F(HipFileAsyncOpFallbackFunctions, devPtr_calls_hipHostGetDevicePointer) +TEST_F(HipFileAsyncOpFallbackMethods, devPtr_calls_hipHostGetDevicePointer) { void *addr = reinterpret_cast(0xABACADBA); EXPECT_CALL(mhip, hipHostGetDevicePointer).WillOnce(Return(addr)); ASSERT_EQ(op->devPtr(), addr); } +struct HipFileAsyncMonitor : HipFileAsyncOp { + AsyncMonitor monitor; +}; + TEST_F(HipFileAsyncMonitor, addOp_and_completeOp_with_valid_params_works) { size_t size = 100; @@ -281,4 +289,381 @@ TEST_F(HipFileAsyncMonitor, addOp_without_completeOp_prints_error_on_AsyncMonito EXPECT_CALL(msys, syslog); } +HIPFILE_WARN_NO_EXIT_DTOR_OFF +struct AsyncIoFunction { + hipFileError_t (*function)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; +}; +static std::array asyncIOFns{ + {{hipFileReadAsync, "hipFileReadAsync"}, {hipFileWriteAsync, "hipFileWriteAsync"}}}; +HIPFILE_WARN_NO_EXIT_DTOR_ON + +struct HipFileReadWriteAsync : public HipFileOpened, public ::testing::WithParamInterface { + void SetUp() override + { + nonnull_void = reinterpret_cast(1); + nonnull_size = reinterpret_cast(1); + nonnull_offset = reinterpret_cast(1); + nonnull_ssize = reinterpret_cast(1); + nonnull_stream = reinterpret_cast(1); + io_op = GetParam().function; + name = GetParam().name; + } + StrictMock mhip; + StrictMock msys; + void *nonnull_void; + size_t *nonnull_size; + ssize_t *nonnull_ssize; + hoff_t *nonnull_offset; + hipStream_t nonnull_stream; + hipFileError_t (*io_op)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; +}; + +TEST_P(HipFileReadWriteAsync, nullSizeReturnsError) +{ + ASSERT_EQ(io_op(nonnull_void, nonnull_void, nullptr, nonnull_offset, nonnull_offset, nonnull_ssize, + nonnull_stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipFileReadWriteAsync, nullFileOffsetReturnsError) +{ + ASSERT_EQ(io_op(nonnull_void, nonnull_void, nonnull_size, nullptr, nonnull_offset, nonnull_ssize, + nonnull_stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipFileReadWriteAsync, nullBufferOffsetReturnsError) +{ + ASSERT_EQ(io_op(nonnull_void, nonnull_void, nonnull_size, nonnull_offset, nullptr, nonnull_ssize, + nonnull_stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipFileReadWriteAsync, nullBytesTransferredReturnsError) +{ + ASSERT_EQ(io_op(nonnull_void, nonnull_void, nonnull_size, nonnull_offset, nonnull_offset, nullptr, + nonnull_stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipFileReadWriteAsync, unregisteredFileReturnsError) +{ + size_t size = 1; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + hoff_t bytes_written = 0; + + ASSERT_EQ(io_op(nonnull_void, nonnull_void, &size, &file_offset, &buffer_offset, &bytes_written, + nonnull_stream), + HipFileOpError(hipFileHandleNotRegistered)); +} + +TEST_P(HipFileReadWriteAsync, registeredBufferLengthTooLongReturnsError) +{ + size_t bufferLength = 100; + expect_buffer_registration(mhip, hipMemoryTypeDevice); + Context::get()->registerBuffer(nonnull_void, bufferLength, 0); + bufferLength++; + EXPECT_CALL(msys, statx); + EXPECT_CALL(msys, fcntl); + EXPECT_CALL(msys, open); + ASSERT_EQ(io_op(Context::get()->registerFile(0), nonnull_void, &bufferLength, nonnull_offset, + nonnull_offset, nonnull_ssize, nonnull_stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipFileReadWriteAsync, badAllocReturnsHipErrorOutOfMemory) +{ + MDriverState driver; + EXPECT_CALL(driver, getRefCount).WillOnce(Throw(std::bad_alloc())); + ASSERT_EQ(io_op(nonnull_void, nonnull_void, nonnull_size, nonnull_offset, nonnull_offset, nullptr, + nonnull_stream), + HipFileHipError(hipErrorOutOfMemory)); +} + +INSTANTIATE_TEST_SUITE_P(HipFileAsyncSuite, HipFileReadWriteAsync, ::testing::ValuesIn(asyncIOFns), + [](const testing::TestParamInfo ¶m_info) { + return param_info.param.name; + }); + +struct FallbackAsyncIO : public HipFileOpened, public ::testing::WithParamInterface { + FallbackAsyncIO() + : mfile{std::make_shared>()}, mbuffer{std::make_shared>()}, + mstream{std::make_shared>()}, size{1024 * 1024}, file_offset{0}, + buffer_offset{0}, bytes_written{0} + { + } + void SetUp() override + { + io_type = GetParam(); + } + StrictMock mhip; + StrictMock msys; + std::shared_ptr> mfile; + std::shared_ptr> mbuffer; + std::shared_ptr> mstream; + IoType io_type; + size_t size; + hoff_t file_offset; + hoff_t buffer_offset; + hoff_t bytes_written; +}; + +TEST_P(FallbackAsyncIO, bufferOffsetNegativeReturnsError) +{ + buffer_offset = -1; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + +TEST_P(FallbackAsyncIO, bufferOffsetTooLargeReturnsError) +{ + buffer_offset = 1024 * 1024 + 1; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + +TEST_P(FallbackAsyncIO, sizeTooLargeReturnsError) +{ + size = 1024 * 1024 + 1; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + +TEST_P(FallbackAsyncIO, fileOffsetNegativeReturnsError) +{ + file_offset = -1; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + +TEST_P(FallbackAsyncIO, mismatchingGpuIdsThrows) +{ + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_CALL(*mbuffer, getGpuId).WillOnce(Return(0)); + EXPECT_CALL(*mstream, getHipDevice).WillOnce(Return(1)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + +TEST_P(FallbackAsyncIO, attemptToQueueCleanupOnStreamSubmissionFailure) +{ + /* This test will log some errors about async operations being outstanding and being unable to free + * memory. The outstanding operations occur because async_io_cleanup is not run. Some memory is unable to + * be freed, due to being alloced with malloc, and attempted to be freed with hipHostFree. This happens + * because we mock the allocation of the pinned memory, but the mocks are destructed before the + * AsyncOpFallback destructor runs, so we can't mock the deallocation of the pinned memory. + */ + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + EXPECT_CALL(*mbuffer, getGpuId).Times(AnyNumber()).WillRepeatedly(Return(0)); + EXPECT_CALL(*mbuffer, getBuffer).WillOnce(Return(reinterpret_cast(0xBADBADBB))); + EXPECT_CALL(*mstream, getHipDevice).WillOnce(Return(0)); + EXPECT_CALL(*mstream, fixedIOSize).Times(AnyNumber()).WillRepeatedly(Return(false)); + EXPECT_CALL(*mstream, fixedFileOffset).Times(AnyNumber()).WillRepeatedly(Return(false)); + EXPECT_CALL(*mstream, fixedBufferOffset).Times(AnyNumber()).WillRepeatedly(Return(false)); + + auto op_data = malloc(sizeof(AsyncOpFallback)); + ASSERT_NE(op_data, nullptr); + auto bounce_buffer = malloc(size); + ASSERT_NE(bounce_buffer, nullptr); + EXPECT_CALL(mhip, hipHostMalloc).WillOnce(Return(op_data)).WillOnce(Return(bounce_buffer)); + EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer), _)) + .WillOnce(Return(reinterpret_cast(0xDEBBBBBB))); + EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(op_data), _)) + .WillOnce(Return(reinterpret_cast(0xDE000000))); + // EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); + // EXPECT_CALL(mhip, hipHostFree(Eq(op_data.get()))); + EXPECT_CALL(mhip, hipDeviceGetAttribute).WillOnce(Return(1024)); + EXPECT_CALL(*mstream, getLock); + EXPECT_CALL(*mstream, getHipStream).Times(AnyNumber()); + EXPECT_CALL(mhip, hipLaunchHostFunc).WillOnce(Throw(Hip::RuntimeError(hipErrorInvalidHandle))); + EXPECT_CALL(mhip, hipLaunchHostFunc(_, Eq(&async_io_cleanup), _)); + + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + Hip::RuntimeError); +} + +INSTANTIATE_TEST_SUITE_P(FallbackAsyncIOSuite, FallbackAsyncIO, + ::testing::Values(IoType::Read, IoType::Write)); + +struct AsyncIoOp : public ::testing::Test { + + AsyncIoOp() + : mfile{std::make_shared>()}, mbuffer{std::make_shared>()}, + mstream{std::make_shared>()} + { + } + void SetUp() override + { + op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + bounce_buffer = std::shared_ptr(new uint8_t[size]); + EXPECT_CALL(mhip, hipHostMalloc) + .WillOnce(Return(op_data.get())) + .WillOnce(Return(bounce_buffer.get())); + EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); + EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); + EXPECT_CALL(mhip, hipHostFree(Eq(op_data.get()))); + EXPECT_CALL(*mstream, fixedBufferOffset) + .Times(AnyNumber()) + .WillRepeatedly(Return(fixed_buffer_offset)); + EXPECT_CALL(*mstream, fixedFileOffset).Times(AnyNumber()).WillRepeatedly(Return(fixed_file_offset)); + EXPECT_CALL(*mstream, fixedIOSize).Times(AnyNumber()).WillRepeatedly(Return(fixed_io_size)); + EXPECT_CALL(*mstream, pageAligned).Times(AnyNumber()).WillRepeatedly(Return(true)); + EXPECT_CALL(*mbuffer, getBuffer); + op = std::shared_ptr(new AsyncOpFallback( + io_type, mfile, mbuffer, mstream, &size, &file_offset, &buffer_offset, &bytes_transferred)); + } + StrictMock mhip; + StrictMock msys; + StrictMock masync_monitor; + std::shared_ptr> mfile; + std::shared_ptr> mbuffer; + std::shared_ptr> mstream; + IoType io_type = IoType::Read; + size_t size = 1_MiB; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + hoff_t bytes_transferred = 0; + bool fixed_buffer_offset = false; + bool fixed_file_offset = false; + bool fixed_io_size = false; + std::shared_ptr op_data; + std::shared_ptr bounce_buffer; + std::shared_ptr op; +}; + +TEST_F(AsyncIoOp, bindIncreasingSizeReturnsError) +{ + size += 1; + async_io_bind_params(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileInvalidValue); +} + +TEST_F(AsyncIoOp, bindInvalidRegionReturnsError) +{ + buffer_offset = 1; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(size)); + async_io_bind_params(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileInvalidValue); +} + +TEST_F(AsyncIoOp, bindAllParamsAreFixedAfter) +{ + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(size)); + async_io_bind_params(op.get()); + auto current_buffer_offset = std::get(op->buffer_offset); + (void)current_buffer_offset; + auto current_file_offset = std::get(op->file_offset); + (void)current_file_offset; + auto current_io_size = std::get(op->size); + (void)current_io_size; +} + +struct AsyncIoOpCleanup : public AsyncIoOp { + void SetUp() override + { + fixed_buffer_offset = true; + fixed_file_offset = true; + fixed_io_size = true; + AsyncIoOp::SetUp(); + } +}; + +TEST_F(AsyncIoOpCleanup, cleanupBytesTransferredIsUpdated) +{ + op->bytes_transferred_internal = 1000; + EXPECT_CALL(masync_monitor, completeOp); + async_io_cleanup(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, *op->bytes_transferred); +} + +TEST_F(AsyncIoOpCleanup, cleanupInvalidOpSetsError) +{ + op->bytes_transferred_internal = 1000; + EXPECT_CALL(masync_monitor, completeOp).WillOnce(Throw(std::invalid_argument("error"))); + async_io_cleanup(op.get()); + ASSERT_EQ(*op->bytes_transferred, -hipFileInternalError); +} + +struct AsyncIoOpBindParams { + IoType io_type; + bool fixed_buffer_offset; + bool fixed_file_offset; + bool fixed_io_size; +}; + +struct AsyncIoOpWithParams : public AsyncIoOp, public ::testing::WithParamInterface { + void SetUp() override + { + auto [_io_type, _fixed_buffer_offset, _fixed_file_offset, _fixed_io_size] = GetParam(); + io_type = _io_type; + fixed_buffer_offset = _fixed_buffer_offset; + fixed_file_offset = _fixed_file_offset; + fixed_io_size = _fixed_io_size; + AsyncIoOp::SetUp(); + } +}; + +TEST_P(AsyncIoOpWithParams, cpuCopyWillReturnEarlyOnPreviousError) +{ + op->bytes_transferred_internal = -1; + // No other mocks will be called + async_io_cpu_copy(op.get()); +} + +TEST_P(AsyncIoOpWithParams, cpuIOReturnsFullSize) +{ + if (io_type == IoType::Read) { + EXPECT_CALL(msys, pread).WillOnce(Return(size)); + } + else { + EXPECT_CALL(msys, pwrite).WillOnce(Return(size)); + } + EXPECT_CALL(*mfile, getBufferedFd); + async_io_cpu_copy(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, size); +} + +TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadErrorReturnsError) +{ + if (io_type == IoType::Read) { + EXPECT_CALL(msys, pread).WillOnce(Throw(std::system_error(3, std::generic_category()))); + } + else { + EXPECT_CALL(msys, pwrite).WillOnce(Throw(std::system_error(3, std::generic_category()))); + } + EXPECT_CALL(*mfile, getBufferedFd); + async_io_cpu_copy(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, -1); +} + +TEST_P(AsyncIoOpWithParams, cpuCopyReadHipErrorReturnsError) +{ + if (io_type == IoType::Read) { + EXPECT_CALL(msys, pread).WillOnce(Throw(Hip::RuntimeError(hipErrorInvalidHandle))); + } + else { + EXPECT_CALL(msys, pwrite).WillOnce(Throw(Hip::RuntimeError(hipErrorInvalidHandle))); + } + EXPECT_CALL(*mfile, getBufferedFd); + async_io_cpu_copy(op.get()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileDriverError); +} + +INSTANTIATE_TEST_SUITE_P(AsyncIoOpWithParamsSuite, AsyncIoOpWithParams, + ::testing::Values(AsyncIoOpBindParams{IoType::Read, true, true, true}, + AsyncIoOpBindParams{IoType::Write, true, true, true})); + HIPFILE_WARN_NO_GLOBAL_CTOR_ON From 0252f3ac8def049779a813b28b0931cc375082f6 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Tue, 28 Oct 2025 16:41:50 -0400 Subject: [PATCH 12/27] hipFile: Adding async system tests. --- test/CMakeLists.txt | 1 + test/system/async.cpp | 872 ++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 873 insertions(+) create mode 100644 test/system/async.cpp diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index cb605659..e79aa0c9 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -19,6 +19,7 @@ set(UNIT_TEST_SOURCE_FILES # They will not pass if a HIP capable GPU is not present # and configured on the system. set(SYSTEM_TEST_SOURCE_FILES + system/async.cpp system/buffer.cpp system/config.cpp system/driver.cpp diff --git a/test/system/async.cpp b/test/system/async.cpp new file mode 100644 index 00000000..fb2012e8 --- /dev/null +++ b/test/system/async.cpp @@ -0,0 +1,872 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#ifdef __HIP_PLATFORM_AMD__ +#include "backend/asyncop-fallback.h" +#include "backend/memcpy-kernel.h" +#include "buffer.h" +#include "context.h" +#include "hip.h" +#include "io.h" +#include "state.h" +#include "stream.h" +#endif + +#include "hipfile-literals.h" +#include "hipfile-warnings.h" +#include "hipfile.h" +#include "test-common.h" +#include "test-options.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#ifdef __HIP_PLATFORM_NVIDIA__ +#include +#endif + +#ifdef __HIP_PLATFORM_AMD__ +using namespace hipFile; +#endif + +HIPFILE_WARN_NO_GLOBAL_CTOR_OFF +extern SystemTestOptions test_env; + +HIPFILE_WARN_NO_EXIT_DTOR_OFF +struct AsyncIoFunction { + hipFileError_t (*function)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; +}; +static std::array asyncIOFns{ + {{hipFileReadAsync, "hipFileReadAsync"}, {hipFileWriteAsync, "hipFileWriteAsync"}}}; +HIPFILE_WARN_NO_EXIT_DTOR_ON + +static bool +isGpuMemory(void *mem) +{ + hipPointerAttribute_t attrs; + assert(hipPointerGetAttributes(&attrs, mem) == hipSuccess); + return attrs.type == hipMemoryTypeDevice; +} + +static std::vector +copyGpuMemory(void *gpu_mem, hoff_t gpu_mem_offset, size_t region_size) +{ + std::vector mem_region(region_size); + assert(hipMemcpy(mem_region.data(), + reinterpret_cast(reinterpret_cast(gpu_mem) + + static_cast(gpu_mem_offset)), + region_size, hipMemcpyDeviceToHost) == hipSuccess); + assert(hipStreamSynchronize(nullptr) == hipSuccess); + return mem_region; +} + +static void +assertMemoryRegionsMatch(void *mem1, hoff_t mem1_offset, void *mem2, hoff_t mem2_offset, size_t region_size) +{ + std::vector mem1_v; + std::vector mem2_v; + if (isGpuMemory(mem1)) { + mem1_v = copyGpuMemory(mem1, mem1_offset, region_size); + mem1 = mem1_v.data(); + } + if (isGpuMemory(mem2)) { + mem2_v = copyGpuMemory(mem2, mem2_offset, region_size); + mem2 = mem2_v.data(); + } + assert(memcmp(mem1, mem2, region_size) == 0); +} + +static void +assertFileAndMemoryRegionsMatch(void *mem, hoff_t mem_offset, int fd, hoff_t fd_offset, size_t region_size) +{ + assert(fd_offset >= 0); + auto file_region = std::vector(region_size); + + ssize_t rv = pread(fd, file_region.data(), region_size, fd_offset); + assert(rv > 0 && static_cast(rv) == region_size); + + assertMemoryRegionsMatch(file_region.data(), 0, mem, mem_offset, region_size); +} + +static void +assertZeroedMemRegion(void *mem, hoff_t mem_offset, size_t region_size) +{ + std::vector mem_v; + if (isGpuMemory(mem)) { + mem_v = copyGpuMemory(mem, mem_offset, region_size); + mem = mem_v.data(); + } + for (size_t i = 0; i < region_size; ++i) { + assert(reinterpret_cast(mem)[i] == 0); + } +} + +static void +assertZeroedFileRegion(int fd, hoff_t fd_offset, size_t region_size) +{ + assert(fd_offset >= 0); + auto file_region = std::vector(region_size); + ssize_t rv = pread(fd, file_region.data(), region_size, fd_offset); + assert(rv > 0 && static_cast(rv) == region_size); + for (size_t i = 0; i < region_size; ++i) { + assert(file_region.data()[i] == 0); + } +} + +static void +randomizeMemoryRegion(void *mem, hoff_t offset, size_t region_size) +{ + ssize_t rv; + int rand_fd = open("/dev/urandom", O_RDONLY); + assert(rand_fd != -1); + if (isGpuMemory(mem)) { + std::vector mem_v(region_size); + rv = read(rand_fd, mem_v.data(), region_size); + assert(rv > 0 && static_cast(rv) == region_size); + assert(hipMemcpy( + reinterpret_cast(reinterpret_cast(mem) + static_cast(offset)), + mem_v.data(), region_size, hipMemcpyHostToDevice) == hipSuccess); + assert(hipStreamSynchronize(nullptr) == hipSuccess); + } + else { + rv = read(rand_fd, + reinterpret_cast(reinterpret_cast(mem) + static_cast(offset)), + region_size); + assert(rv > 0 && static_cast(rv) == region_size); + } + assert(close(rand_fd) == 0); +} + +static void +zeroMemoryRegion(void *mem, hoff_t offset, size_t region_size) +{ + if (isGpuMemory(mem)) { + assert(hipMemset( + reinterpret_cast(reinterpret_cast(mem) + static_cast(offset)), + 0, region_size) == hipSuccess); + assert(hipStreamSynchronize(nullptr) == hipSuccess); + } + else { + memset(reinterpret_cast(reinterpret_cast(mem) + static_cast(offset)), 0, + region_size); + } +} + +static void +zeroFileRegion(int fd, size_t size, hoff_t offset = 0) +{ + auto vec = std::vector(size, 0); + ssize_t rv = pwrite(fd, vec.data(), size, offset); + assert(rv > 0 && static_cast(rv) == size); +} + +static void +randomizeFileRegion(int fd, size_t size, hoff_t offset = 0) +{ + auto vec = std::vector(size, 0); + randomizeMemoryRegion(vec.data(), 0, size); + ssize_t rv = pwrite(fd, vec.data(), size, offset); + assert(rv > 0 && static_cast(rv) == size); +} + +#ifdef __HIP_PLATFORM_AMD__ +class HipAsyncMemcpyKernel : public ::testing::Test { +public: + HipAsyncMemcpyKernel() : tf{test_env.ais_capable_dir} + { + } + void SetUp() override + { + ASSERT_EQ(hipFileDriverOpen(), HIPFILE_SUCCESS); + ASSERT_EQ(hipMalloc(&dev_ptr, buffer_size), hipSuccess); + ASSERT_EQ(hipFileBufRegister(dev_ptr, buffer_size, 0), HIPFILE_SUCCESS); + hipFileDescr_t d{hipFileHandleTypeOpaqueFD, {tf.fd}, nullptr}; + ASSERT_EQ(hipFileHandleRegister(&fh, &d), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamCreateWithFlags(&hip_stream, hipStreamNonBlocking), hipSuccess); + ASSERT_EQ(hipFileStreamRegister(hip_stream, 0xf), HIPFILE_SUCCESS); + auto [_file, _buffer, _stream] = + Context::get()->getFileBufferAndStream(fh, dev_ptr, buffer_size, 0, hip_stream); + file = _file; + buffer = _buffer; + stream = _stream; + op = std::shared_ptr(new AsyncOpFallback( + io_type, file, buffer, stream, &io_size, &file_offset, &buffer_offset, &bytes_transferred)); + bytes_transferred = static_cast(io_size); + if (io_type == IoType::Read) { + // For a read, bytes_transferred_internal needs to be set to simulate that the read from disk + // occurred. We fill the source (CPU bounce buffer) with random data and memset the destination + // (GPU buffer) to zero. + randomizeMemoryRegion(op->bounceBufferHostPtr(), 0, io_size); + op->bytes_transferred_internal = static_cast(io_size); + ASSERT_EQ(hipMemset(op->gpu_buffer, 0, buffer_size), hipSuccess); + ASSERT_EQ(hipStreamSynchronize(nullptr), hipSuccess); + } + else { + // For a write, we fill the source (GPU bounce buffer) with random data and memset the destination + // (CPU bounce buffer) to zero. + randomizeMemoryRegion(op->gpu_buffer, 0, buffer_size); + memset(op->bounceBufferHostPtr(), 0, io_size); + } + op_dev_ptr = op->devPtr(); + kernel_args[0] = {&op_dev_ptr}; + max_threads_per_block = Context::get()->hipDeviceGetAttribute( + hipDeviceAttributeMaxThreadsPerBlock, buffer->getGpuId()); + } + void TearDown() override + { + ASSERT_EQ(hipFileStreamDeregister(hip_stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamDestroy(hip_stream), hipSuccess); + ASSERT_EQ(hipFree(dev_ptr), hipSuccess); + } + void *dev_ptr; + Tmpfile tf; + hipFileHandle_t fh; + IoType io_type = IoType::Read; + size_t file_size = 4_KiB; + size_t buffer_size = 4_KiB; + size_t io_size = 4_KiB; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + hipStream_t hip_stream; + std::shared_ptr file; + std::shared_ptr buffer; + std::shared_ptr stream; + std::shared_ptr op; + void *op_dev_ptr; + void *kernel_args[1]; + int max_threads_per_block; +}; + +TEST_F(HipAsyncMemcpyKernel, unfixedSizeReturnsInval) +{ + op->size = &io_size; + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), kernel_args, 0, + op->stream->getHipStream()); + Context::get()->hipStreamSynchronize(op->stream->getHipStream()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileInvalidValue); +} + +TEST_F(HipAsyncMemcpyKernel, unfixedBufferOffsetReturnsInval) +{ + op->buffer_offset = &buffer_offset; + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), kernel_args, 0, + op->stream->getHipStream()); + Context::get()->hipStreamSynchronize(op->stream->getHipStream()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileInvalidValue); +} + +TEST_F(HipAsyncMemcpyKernel, nullCpuBufferDevPointerReturnsInval) +{ + op->bounce_buffer_dev_ptr = nullptr; + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), kernel_args, 0, + op->stream->getHipStream()); + Context::get()->hipStreamSynchronize(op->stream->getHipStream()); + ASSERT_EQ(op->bytes_transferred_internal, -hipFileInvalidValue); +} + +struct AsyncMemcpyKernelTestParams { + IoType io_type; + size_t file_size; + size_t buffer_size; + size_t io_size; + hoff_t file_offset; + hoff_t buffer_offset; + ssize_t bytes_transferred_start; + ssize_t bytes_transferred_result; +}; + +class HipAsyncMemcpyKernelWithParams : public HipAsyncMemcpyKernel, + public ::testing::WithParamInterface { +public: + void SetUp() override + { + auto [_io_type, _file_size, _buffer_size, _io_size, _file_offset, _buffer_offset, + _bytes_transferred_start, _bytes_transferred_result] = GetParam(); + io_type = _io_type; + file_size = _file_size; + buffer_size = _buffer_size; + io_size = _io_size; + file_offset = _file_offset; + buffer_offset = _buffer_offset; + bytes_transferred_start = _bytes_transferred_start; + bytes_transferred_result = _bytes_transferred_result; + HipAsyncMemcpyKernel::SetUp(); + op->bytes_transferred_internal = bytes_transferred_start; + } + void TearDown() override + { + HipAsyncMemcpyKernel::TearDown(); + } + ssize_t bytes_transferred_start; + ssize_t bytes_transferred_result; +}; + +TEST_P(HipAsyncMemcpyKernelWithParams, verifyIoRegions) +{ + Context::get()->hipLaunchKernel(reinterpret_cast(hipFileMemcpyKernel), dim3(1), + dim3(static_cast(max_threads_per_block)), kernel_args, 0, + op->stream->getHipStream()); + Context::get()->hipStreamSynchronize(op->stream->getHipStream()); + ASSERT_EQ(op->bytes_transferred_internal, bytes_transferred_result); + if (op->bytes_transferred_internal > 0) { + if (io_type == IoType::Read) { + assertZeroedMemRegion(op->gpu_buffer, 0, static_cast(buffer_offset)); + } + assertMemoryRegionsMatch(op->bounceBufferHostPtr(), 0, op->gpu_buffer, buffer_offset, + static_cast(op->bytes_transferred_internal)); + if (io_type == IoType::Read) { + size_t end_length = buffer_size - (static_cast(buffer_offset) + + static_cast(op->bytes_transferred_internal)); + assertZeroedMemRegion(op->gpu_buffer, buffer_offset + op->bytes_transferred_internal, end_length); + } + } +} + +INSTANTIATE_TEST_SUITE_P( + HipAsyncMemcpyKernelSuite, HipAsyncMemcpyKernelWithParams, + testing::Values( + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB, 0, 0, 1_MiB, 1_MiB}, + AsyncMemcpyKernelTestParams{IoType::Write, 1_MiB, 1_MiB, 1_MiB, 0, 0, 1_MiB, 1_MiB}, + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB - 8_KiB, 0, 4_KiB, 1_MiB - 8_KiB, + 1_MiB - 8_KiB}, + AsyncMemcpyKernelTestParams{IoType::Write, 1_MiB, 1_MiB, 1_MiB - 8_KiB, 0, 4_KiB, 1_MiB - 8_KiB, + 1_MiB - 8_KiB}, + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB - 1, 0, 0, 1_MiB - 1, 1_MiB - 1}, + AsyncMemcpyKernelTestParams{IoType::Write, 1_MiB, 1_MiB, 1_MiB - 1, 0, 0, 1_MiB - 1, 1_MiB - 1}, + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB - 1, 0, 1, 1_MiB - 1, 1_MiB - 1}, + AsyncMemcpyKernelTestParams{IoType::Write, 1_MiB, 1_MiB, 1_MiB - 1, 0, 1, 1_MiB - 1, 1_MiB - 1}, + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB - 2, 0, 1, 1_MiB - 2, 1_MiB - 2}, + AsyncMemcpyKernelTestParams{IoType::Write, 1_MiB, 1_MiB, 1_MiB - 2, 0, 1, 1_MiB - 2, 1_MiB - 2}, + AsyncMemcpyKernelTestParams{IoType::Read, 1_MiB, 1_MiB, 1_MiB, 0, 0, 512_KiB, 512_KiB}), + [](const testing::TestParamInfo ¶m_info) { + auto params = param_info.param; + + std::stringstream label; + if (params.io_type == IoType::Read) { + label << "read"; + } + else { + label << "write"; + } + label << "_" << params.file_size << "_" << params.buffer_size << "_" << params.io_size << "_" + << params.file_offset << "_" << params.buffer_offset << "_" << params.bytes_transferred_start + << "_" << params.bytes_transferred_result; + return label.str(); + }); +#endif + +class HipAsync : public ::testing::Test { +public: + HipAsync() : tf{test_env.ais_capable_dir} + { + } + void SetUp() override + { + ASSERT_EQ(hipFileDriverOpen(), HIPFILE_SUCCESS); + ASSERT_EQ(hipMalloc(&dev_ptr, buffer_size), hipSuccess); + + randomizeFileRegion(tf.fd, file_size); + + ASSERT_EQ(hipFileBufRegister(dev_ptr, buffer_size, 0), HIPFILE_SUCCESS); + hipFileDescr_t d{hipFileHandleTypeOpaqueFD, {tf.fd}, nullptr}; + ASSERT_EQ(hipFileHandleRegister(&fh, &d), HIPFILE_SUCCESS); + } + + void TearDown() override + { + ASSERT_EQ(hipFree(dev_ptr), hipSuccess); + } + void *dev_ptr; + Tmpfile tf; + hipFileHandle_t fh; + size_t io_size = 1_MiB; + size_t file_size = 1_MiB; + size_t buffer_size = 1_MiB; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; +}; + +class HipAsyncStreamFixed : public HipAsync { +public: + void SetUp() override + { + HipAsync::SetUp(); + zeroFileRegion(tf.fd, file_size); + ASSERT_EQ(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), hipSuccess); + ASSERT_EQ(hipFileStreamRegister(stream, 0xf), HIPFILE_SUCCESS); + } + void TearDown() override + { + ASSERT_EQ(hipFileStreamDeregister(stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamDestroy(stream), hipSuccess); + HipAsync::TearDown(); + } + hipStream_t stream; +}; + +TEST_F(HipAsyncStreamFixed, readRegionPastEndOfFile) +{ + file_offset += 4_KiB; + ASSERT_EQ( + hipFileReadAsync(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, file_size - 4_KiB); + ASSERT_EQ(bytes_transferred, file_size - 4_KiB); +} + +class HipAsyncReadWriteStreamFixedWithParams : public HipAsync, + public ::testing::WithParamInterface { +public: + void SetUp() override + { + HipAsync::SetUp(); + auto params = GetParam(); + io_op = params.function; + name = params.name; + ASSERT_EQ(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), hipSuccess); + ASSERT_EQ(hipFileStreamRegister(stream, 0xf), HIPFILE_SUCCESS); + } + void TearDown() override + { + ASSERT_EQ(hipFileStreamDeregister(stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamDestroy(stream), hipSuccess); + HipAsync::TearDown(); + } + hipFileError_t (*io_op)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; + hipStream_t stream; +}; + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullBufferBaseReturnsError) +{ + ASSERT_EQ(io_op(fh, nullptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullSizeReturnsError) +{ + ASSERT_EQ(io_op(fh, dev_ptr, nullptr, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullFileOffsetReturnsError) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, nullptr, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullBufferOffsetReturnsError) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, nullptr, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullBytesReadReturnsError) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, nullptr, stream), + HipFileOpError(hipFileInvalidValue)); +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, ioSizeTooLargeForBuffer) +{ + io_size += 4096; +#ifdef __HIP_PLATFORM_AMD__ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +#else + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); +#endif +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, fileOffsetNegative) +{ + file_offset = -4096; +#ifdef __HIP_PLATFORM_AMD__ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +#else + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); +#endif +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetNegative) +{ + buffer_offset = -4096; +#ifdef __HIP_PLATFORM_AMD__ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +#else + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); +#endif +} + +TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetTooLarge) +{ + buffer_offset = static_cast(buffer_size) + 4096; +#ifdef __HIP_PLATFORM_AMD__ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HipFileOpError(hipFileInvalidValue)); +#else + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); +#endif +} + +INSTANTIATE_TEST_SUITE_P( + HipAsyncReadWriteStreamFixedWithParamsSuite, HipAsyncReadWriteStreamFixedWithParams, + testing::ValuesIn(asyncIOFns), + [](const testing::TestParamInfo ¶m_info) { + auto params = param_info.param; + return params.name; + }); + +class HipAsyncReadWriteParamsUnchanged + : public HipAsync, + public ::testing::WithParamInterface> { +public: + void SetUp() override + { + HipAsync::SetUp(); + auto params = GetParam(); + auto op_info = std::get<0>(params); + io_op = op_info.function; + name = op_info.name; + if (name == "hipFileReadAsync") { + zeroMemoryRegion(dev_ptr, 0, buffer_size); + randomizeFileRegion(tf.fd, file_size); + } + else { + zeroFileRegion(tf.fd, file_size); + randomizeMemoryRegion(dev_ptr, 0, buffer_size); + } + auto flags = std::get<1>(params) | std::get<2>(params) | std::get<3>(params); + ASSERT_EQ(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), hipSuccess); + ASSERT_EQ(hipFileStreamRegister(stream, flags), HIPFILE_SUCCESS); + } + void TearDown() override + { + ASSERT_EQ(hipFileStreamDeregister(stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamDestroy(stream), hipSuccess); + HipAsync::TearDown(); + } + hipFileError_t (*io_op)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; + hipStream_t stream; +}; + +TEST_P(HipAsyncReadWriteParamsUnchanged, zeroOffsets) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); +} + +TEST_P(HipAsyncReadWriteParamsUnchanged, ioSizeUnaligned) +{ + io_size -= 1; + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, static_cast(io_size), 1); + } + else { + assertZeroedFileRegion(tf.fd, static_cast(io_size), 1); + } +} + +TEST_P(HipAsyncReadWriteParamsUnchanged, fileOffsetAndIoSizeUnaligned) +{ + file_offset = 1; + io_size -= 1; + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, static_cast(io_size), 1); + } + else { + assertZeroedFileRegion(tf.fd, 0, 1); + } +} + +TEST_P(HipAsyncReadWriteParamsUnchanged, bufferOffsetAndIoSizeUnaligned) +{ + buffer_offset = 1; + io_size -= 1; + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, 0, 1); + } + else { + assertZeroedFileRegion(tf.fd, static_cast(io_size), 1); + } +} + +TEST_P(HipAsyncReadWriteParamsUnchanged, zeroSized) +{ + io_size = 0; + bytes_transferred = 1; + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + ASSERT_EQ(bytes_transferred, 0); +} + +TEST_P(HipAsyncReadWriteParamsUnchanged, multipleOpsOnSameStreamAreSequential) +{ + + size_t num_ios = (io_size / 8_KiB) - 1; + io_size = 16_KiB; + std::vector bytes_transferred_v(num_ios, 0); + std::vector buffer_offset_v(num_ios, 0); + std::vector file_offset_v(num_ios, 0); + if (name == "hipFileReadAsync") { + for (size_t i = 0; i < num_ios; ++i) { + buffer_offset_v[i] = static_cast(8_KiB) * static_cast(i); + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset_v[i], &bytes_transferred_v[i], + stream), + HIPFILE_SUCCESS); + } + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + + for (size_t i = 0; i < num_ios; ++i) { + ASSERT_EQ(bytes_transferred_v[i], io_size); + size_t region_size = 8_KiB; + if (i == num_ios - 1) { + region_size = 16_KiB; + } + assertFileAndMemoryRegionsMatch(dev_ptr, static_cast(i * 8_KiB), tf.fd, 0, region_size); + } + } + else { + for (size_t i = 0; i < num_ios; ++i) { + file_offset_v[i] = static_cast(8_KiB) * static_cast(i); + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset_v[i], &buffer_offset, &bytes_transferred_v[i], + stream), + HIPFILE_SUCCESS); + } + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + + for (size_t i = 0; i < num_ios; ++i) { + ASSERT_EQ(bytes_transferred_v[i], io_size); + size_t region_size = 8_KiB; + if (i == num_ios - 1) { + region_size = 16_KiB; + } + assertFileAndMemoryRegionsMatch(dev_ptr, 0, tf.fd, static_cast(i * 8_KiB), region_size); + } + } +} + +static auto +hipFileAsyncIOFlagsPowerSet() +{ + return ::testing::Combine(::testing::ValuesIn(asyncIOFns), + ::testing::Values(0, HIPFILE_STREAM_FIXED_BUF_OFFSET), + ::testing::Values(0, HIPFILE_STREAM_FIXED_FILE_OFFSET), + ::testing::Values(0, HIPFILE_STREAM_FIXED_FILE_SIZE)); +} + +INSTANTIATE_TEST_SUITE_P( + HipAsyncStreamUnchanged, HipAsyncReadWriteParamsUnchanged, hipFileAsyncIOFlagsPowerSet(), + [](const testing::TestParamInfo ¶m_info) { + auto params = param_info.param; + auto flags = std::get<1>(params) | std::get<2>(params) | std::get<3>(params); + return std::get<0>(param_info.param).name + "_" + std::to_string(flags); + }); + +#ifdef __HIP_PLATFORM_AMD__ + +class HipAsyncStreamUnfixedPaused : public HipAsync { +public: + void SetUp() override + { + HipAsync::SetUp(); + int attr = 0; + ASSERT_EQ(hipDeviceGetAttribute(&attr, hipDeviceAttributeCanUseStreamWaitValue, 0), hipSuccess); + ASSERT_EQ(attr, 1); + ASSERT_EQ(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), hipSuccess); + ASSERT_EQ(hipFileStreamRegister(stream, 0), HIPFILE_SUCCESS); + ASSERT_EQ( + hipExtMallocWithFlags(reinterpret_cast(&flag), sizeof(uint64_t), hipMallocSignalMemory), + hipSuccess); + updateFlag(0); + ASSERT_EQ(hipStreamWaitValue64(stream, flag, 1, hipStreamWaitValueEq), hipSuccess); + } + void TearDown() override + { + ASSERT_EQ(hipFileStreamDeregister(stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamDestroy(stream), hipSuccess); + HipAsync::TearDown(); + } + + void updateFlag(uint64_t new_flag) + { + ASSERT_EQ(hipMemcpy(flag, &new_flag, sizeof(uint64_t), hipMemcpyHostToDevice), hipSuccess); + } + hipStream_t stream; + uint64_t *flag; +}; + +class HipAsyncStreamUnfixedPausedReadWrite : public HipAsyncStreamUnfixedPaused, + public ::testing::WithParamInterface { +public: + void SetUp() override + { + HipAsyncStreamUnfixedPaused::SetUp(); + io_op = GetParam().function; + name = GetParam().name; + if (name == "hipFileReadAsync") { + zeroMemoryRegion(dev_ptr, 0, buffer_size); + randomizeFileRegion(tf.fd, file_size); + } + else { + randomizeMemoryRegion(dev_ptr, 0, buffer_size); + zeroFileRegion(tf.fd, file_size); + } + } + void TearDown() override + { + HipAsyncStreamUnfixedPaused::TearDown(); + } + hipFileError_t (*io_op)(hipFileHandle_t, void *, size_t *, hoff_t *, hoff_t *, ssize_t *, hipStream_t); + std::string name; +}; + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, smallerIOSizeIsValid) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + io_size -= 4_KiB; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, static_cast(io_size), 4_KiB); + } + else { + assertZeroedFileRegion(tf.fd, static_cast(io_size), 4_KiB); + } +} + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, increasedIoSizeIsInvalid) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + io_size += 4096; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidValue); +} + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, changedBufferOffsetCanCauseInvalidIoSize) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + buffer_offset += 4096; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidValue); +} + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, changedIOSizeAndFileOffset) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + io_size -= 4096; + file_offset += 4096; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, static_cast(io_size), 4_KiB); + } + else { + assertZeroedFileRegion(tf.fd, 0, 4_KiB); + } +} + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, changedIOSizeAndBufferOffset) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + io_size -= 4096; + buffer_offset += 4096; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, 0, 4_KiB); + } + else { + assertZeroedFileRegion(tf.fd, static_cast(io_size), 4_KiB); + } +} + +TEST_P(HipAsyncStreamUnfixedPausedReadWrite, changedIOSizeAndBufferOffsetAndFileOffset) +{ + ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), + HIPFILE_SUCCESS); + io_size -= 4096; + buffer_offset += 4096; + file_offset += 4096; + updateFlag(1); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, io_size); + assertFileAndMemoryRegionsMatch(dev_ptr, buffer_offset, tf.fd, file_offset, io_size); + if (name == "hipFileReadAsync") { + assertZeroedMemRegion(dev_ptr, 0, 4_KiB); + } + else { + assertZeroedFileRegion(tf.fd, 0, 4_KiB); + } +} + +INSTANTIATE_TEST_SUITE_P( + HipAsyncStreamUnfixed, HipAsyncStreamUnfixedPausedReadWrite, ::testing::ValuesIn(asyncIOFns), + [](const testing::TestParamInfo ¶m_info) { + return param_info.param.name; + }); + +#endif + +HIPFILE_WARN_NO_GLOBAL_CTOR_ON From 90b878b6f1f9c6b56a7a846d64a33db85a379554 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 21 Jan 2026 13:15:18 -0500 Subject: [PATCH 13/27] Build GPU targets listed in ROCm compatibility matrix for 7.1.1 --- .github/workflows/build-ais.yml | 2 ++ 1 file changed, 2 insertions(+) diff --git a/.github/workflows/build-ais.yml b/.github/workflows/build-ais.yml index c718a615..65059477 100644 --- a/.github/workflows/build-ais.yml +++ b/.github/workflows/build-ais.yml @@ -19,6 +19,7 @@ env: AIS_INPUT_ROCM_VERSION: ${{ inputs.rocm_version }} # Code coverage report only vetted to work for amdclang++ on Ubuntu AIS_USE_CODE_COVERAGE: ${{ inputs.cxx_compiler == 'amdclang++' && inputs.platform == 'ubuntu' }} + AIS_HIP_ARCHITECTURES: gfx950;gfx1201;gfx1200;gfx1101;gfx1100;gfx1030;gfx942;gfx90a;gfx908 on: workflow_call: inputs: @@ -141,6 +142,7 @@ jobs: cmake \ -DCMAKE_CXX_COMPILER="${_AIS_INPUT_CXX_COMPILER}" \ -DCMAKE_CXX_FLAGS="-Werror" \ + -DCMAKE_HIP_ARCHITECTURES="${{ env.AIS_HIP_ARCHITECTURES }}" \ -DCMAKE_HIP_PLATFORM=amd \ -DAIS_BUILD_DOCS=ON \ -DAIS_USE_CODE_COVERAGE=${{ env.AIS_USE_CODE_COVERAGE == 'true' && 'ON' || 'OFF' }} \ From 3a3dd82c98dd401e29198137a0e3163ef5ab349a Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Thu, 29 Jan 2026 16:43:03 -0500 Subject: [PATCH 14/27] fixup! hipFile: Add device_id to Stream. review: Move hipStreamGetDevice call to Stream constructor. --- src/amd_detail/stream.cpp | 15 ++++++--------- src/amd_detail/stream.h | 2 +- test/amd_detail/stream.cpp | 1 - 3 files changed, 7 insertions(+), 11 deletions(-) diff --git a/src/amd_detail/stream.cpp b/src/amd_detail/stream.cpp index 149e659a..aef4e934 100644 --- a/src/amd_detail/stream.cpp +++ b/src/amd_detail/stream.cpp @@ -18,10 +18,8 @@ namespace hipFile { -Stream::Stream(const hipStream_t _hip_stream, hipDevice_t _device_id, uint32_t flags, - const PassKey &) - : hip_stream{_hip_stream}, device_id{_device_id}, - fixed_buf_offset{(flags & HIPFILE_STREAM_FIXED_BUF_OFFSET) != 0}, +Stream::Stream(const hipStream_t _hip_stream, uint32_t flags, const PassKey &) + : hip_stream{_hip_stream}, device_id{0}, fixed_buf_offset{(flags & HIPFILE_STREAM_FIXED_BUF_OFFSET) != 0}, fixed_file_offset{(flags & HIPFILE_STREAM_FIXED_FILE_OFFSET) != 0}, fixed_io_size{(flags & HIPFILE_STREAM_FIXED_FILE_SIZE) != 0}, page_aligned{(flags & HIPFILE_STREAM_PAGE_ALIGNED_INPUTS) != 0} @@ -30,6 +28,8 @@ Stream::Stream(const hipStream_t _hip_stream, hipDevice_t _device_id, uint32_t f if ((flags & HIPFILE_STREAM_FLAGS_MASK) != flags) { throw std::invalid_argument("Invalid flags for stream"); } + + device_id = Context::get()->hipStreamGetDevice(hip_stream); } hipStream_t @@ -75,9 +75,7 @@ StreamMap::registerStream(hipStream_t hip_stream, uint32_t flags) throw std::invalid_argument("Stream is already registered"); } - hipDevice_t device_id = Context::get()->hipStreamGetDevice(hip_stream); - - auto stream = std::shared_ptr(new Stream(hip_stream, device_id, flags, PassKey{})); + auto stream = std::shared_ptr(new Stream(hip_stream, flags, PassKey{})); StreamMap::from_ptr[hip_stream] = stream; } @@ -97,8 +95,7 @@ StreamMap::getStream(hipStream_t hip_stream) { auto itr = StreamMap::from_ptr.find(hip_stream); if (StreamMap::from_ptr.end() == itr) { - hipDevice_t device_id = Context::get()->hipStreamGetDevice(hip_stream); - return std::shared_ptr(new Stream(hip_stream, device_id, 0, PassKey{})); + return std::shared_ptr(new Stream(hip_stream, 0, PassKey{})); } return itr->second; diff --git a/src/amd_detail/stream.h b/src/amd_detail/stream.h index 45e3f2d9..6ef04877 100644 --- a/src/amd_detail/stream.h +++ b/src/amd_detail/stream.h @@ -41,7 +41,7 @@ class Stream : public IStream { virtual bool pageAligned() const override; virtual std::unique_lock getLock() override; - Stream(const hipStream_t hip_stream, hipDevice_t device_id, uint32_t flags, const PassKey &k); + Stream(const hipStream_t hip_stream, uint32_t flags, const PassKey &k); private: Stream(const Stream &) = delete; diff --git a/test/amd_detail/stream.cpp b/test/amd_detail/stream.cpp index 9a9b4668..d137f3e1 100644 --- a/test/amd_detail/stream.cpp +++ b/test/amd_detail/stream.cpp @@ -77,7 +77,6 @@ TEST_F(HipFileStream, get_stream_with_unregistered_stream_works) TEST_F(HipFileStream, register_with_invalid_flags_throws) { - EXPECT_CALL(mhip, hipStreamGetDevice); ASSERT_THROW(stream_map.registerStream(nonnull_stream, HIPFILE_STREAM_FLAGS_MASK + 1), std::invalid_argument); } From f21c275ffc49889dce919b7018f4a02e32aada1f Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Thu, 29 Jan 2026 16:49:08 -0500 Subject: [PATCH 15/27] fixup! hipFile: Put paramsValid check in io.h. review: Fix error in paramsValid check and use it in fallback and fastpath io functions. --- src/amd_detail/backend/fallback.cpp | 18 ++---------------- src/amd_detail/backend/fastpath.cpp | 17 ++--------------- src/amd_detail/io.cpp | 2 +- 3 files changed, 5 insertions(+), 32 deletions(-) diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index cb0d4e08..cac01fd0 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -60,24 +60,10 @@ ssize_t Fallback::io(IoType io_type, shared_ptr file, shared_ptr buffer, size_t size, hoff_t file_offset, hoff_t buffer_offset, size_t chunk_size) { - size_t buflen{buffer->getLength()}; - size = min(size, hipFile::MAX_RW_COUNT); - if (file_offset < 0) { - throw std::invalid_argument("Negative file offset"); - } - - if (buffer_offset < 0) { - throw std::invalid_argument("Negative buffer offset"); - } - - if (buflen <= static_cast(buffer_offset)) { - throw std::invalid_argument("Buffer offset larger than buffer length"); - } - - if (buflen - static_cast(buffer_offset) < size) { - throw std::invalid_argument("IO could overflow buffer"); + if (!paramsValid(buffer, size, file_offset, buffer_offset)) { + throw std::invalid_argument("The selected file or buffer region is invalid"); } auto ptr = Context::get()->mmap(nullptr, chunk_size, PROT_READ | PROT_WRITE, diff --git a/src/amd_detail/backend/fastpath.cpp b/src/amd_detail/backend/fastpath.cpp index 1238b778..0bad9f1c 100644 --- a/src/amd_detail/backend/fastpath.cpp +++ b/src/amd_detail/backend/fastpath.cpp @@ -157,24 +157,11 @@ Fastpath::io(IoType type, shared_ptr file, shared_ptr buffer, si void *devptr{reinterpret_cast(reinterpret_cast(buffer->getBuffer()) + buffer_offset)}; hipAmdFileHandle_t handle{}; size_t nbytes{}; - size_t buflen{buffer->getLength()}; handle.fd = file->getUnbufferedFd().value(); - if (file_offset < 0) { - throw std::invalid_argument("Negative file offset"); - } - - if (buffer_offset < 0) { - throw std::invalid_argument("Negative buffer offset"); - } - - if (buflen <= static_cast(buffer_offset)) { - throw std::invalid_argument("Invalid buffer offset"); - } - - if (buflen - static_cast(buffer_offset) < size) { - throw std::invalid_argument("IO could overflow buffer"); + if (!paramsValid(buffer, size, file_offset, buffer_offset)) { + throw std::invalid_argument("The selected file or buffer region is invalid"); } // Currently, when IO sizes > MAX_RW_COUNT are submitted to amdgpu/kfd an diff --git a/src/amd_detail/io.cpp b/src/amd_detail/io.cpp index 874688ea..70d1b44c 100644 --- a/src/amd_detail/io.cpp +++ b/src/amd_detail/io.cpp @@ -18,7 +18,7 @@ paramsValid(const std::shared_ptr &buffer, size_t size, hoff_t file_off if (buffer_offset < 0) { return false; } - if (buffer_length < static_cast(buffer_offset)) { + if (buffer_length <= static_cast(buffer_offset)) { return false; } if (buffer_length - static_cast(buffer_offset) < size) { From 29225149a8bd3dd4e68decbccd6d73d068cf71ae Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Thu, 29 Jan 2026 16:53:29 -0500 Subject: [PATCH 16/27] fixup! hipFile: Add async memcpy kernel. review: Remove blockIdx.x and document that we expect a single block. --- src/amd_detail/backend/memcpy-kernel.hip | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/amd_detail/backend/memcpy-kernel.hip b/src/amd_detail/backend/memcpy-kernel.hip index 4f050e35..f70efe51 100644 --- a/src/amd_detail/backend/memcpy-kernel.hip +++ b/src/amd_detail/backend/memcpy-kernel.hip @@ -15,7 +15,8 @@ namespace hipFile { __global__ void hipFileMemcpyKernel(AsyncOpFallback *op) { - const size_t start_idx = threadIdx.x + blockDim.x * blockIdx.x; + // Expects to be called with a single block + const size_t start_idx = threadIdx.x; const ssize_t bytes_transferred_internal = op->bytes_transferred_internal; if (bytes_transferred_internal < 0) From a5bbbf80f4f13e21ddb18e17700ee2d07c3e9960 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 4 Feb 2026 13:17:59 -0500 Subject: [PATCH 17/27] fixup! hipFile: Add fallback async operation. review: Add EINTR handling. Remove impossible HIP error path. --- src/amd_detail/backend/fallback.cpp | 33 +++++++++++++++-------------- 1 file changed, 17 insertions(+), 16 deletions(-) diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index cac01fd0..4fe31baa 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -239,12 +239,12 @@ async_io_cpu_copy(void *userargs) return; } - try { - while (bytes_transferred < size) { - void *cur_buf_position = reinterpret_cast( - reinterpret_cast(op->bounceBufferHostPtr()) + bytes_transferred); - hoff_t cur_file_offset = file_offset + static_cast(bytes_transferred); - size_t remaining_bytes = size - bytes_transferred; + while (bytes_transferred < size) { + void *cur_buf_position = reinterpret_cast( + reinterpret_cast(op->bounceBufferHostPtr()) + bytes_transferred); + hoff_t cur_file_offset = file_offset + static_cast(bytes_transferred); + size_t remaining_bytes = size - bytes_transferred; + try { switch (op->io_type) { case IoType::Read: ret = Context::get()->pread(op->file->getBufferedFd(), cur_buf_position, @@ -257,18 +257,19 @@ async_io_cpu_copy(void *userargs) default: throw std::runtime_error("Invalid IO type"); } - if (ret == 0) { - break; + } + catch (const std::system_error &e) { + if (e.code().value() == EINTR) { + continue; } - bytes_transferred += static_cast(ret); + op->bytes_transferred_internal = -1; + return; } - op->bytes_transferred_internal = static_cast(bytes_transferred); - } - catch (std::system_error &e) { - op->bytes_transferred_internal = -1; - } - catch (Hip::RuntimeError &e) { - op->bytes_transferred_internal = -hipFileDriverError; + if (ret == 0) { + break; + } + bytes_transferred += static_cast(ret); } + op->bytes_transferred_internal = static_cast(bytes_transferred); } } From 180b0f4732ea7efb288a7ea7ba6db630d03921e4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 4 Feb 2026 13:18:10 -0500 Subject: [PATCH 18/27] fixup! hipFile: Adding async unit tests. review: Remove test for impossible HIP path. --- test/amd_detail/async.cpp | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 99d09033..512947b2 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -649,19 +649,6 @@ TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadErrorReturnsError) ASSERT_EQ(op->bytes_transferred_internal, -1); } -TEST_P(AsyncIoOpWithParams, cpuCopyReadHipErrorReturnsError) -{ - if (io_type == IoType::Read) { - EXPECT_CALL(msys, pread).WillOnce(Throw(Hip::RuntimeError(hipErrorInvalidHandle))); - } - else { - EXPECT_CALL(msys, pwrite).WillOnce(Throw(Hip::RuntimeError(hipErrorInvalidHandle))); - } - EXPECT_CALL(*mfile, getBufferedFd); - async_io_cpu_copy(op.get()); - ASSERT_EQ(op->bytes_transferred_internal, -hipFileDriverError); -} - INSTANTIATE_TEST_SUITE_P(AsyncIoOpWithParamsSuite, AsyncIoOpWithParams, ::testing::Values(AsyncIoOpBindParams{IoType::Read, true, true, true}, AsyncIoOpBindParams{IoType::Write, true, true, true})); From e701954f4958d2e00115a4505a1640b35d729842 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 4 Feb 2026 14:02:05 -0500 Subject: [PATCH 19/27] review: Limit IO to MAX_RW_COUNT. --- src/amd_detail/backend/asyncop-fallback.cpp | 8 ++-- src/amd_detail/backend/fallback.cpp | 6 ++- test/amd_detail/async.cpp | 47 +++++++++++++++++++++ 3 files changed, 56 insertions(+), 5 deletions(-) diff --git a/src/amd_detail/backend/asyncop-fallback.cpp b/src/amd_detail/backend/asyncop-fallback.cpp index e4023158..e1306947 100644 --- a/src/amd_detail/backend/asyncop-fallback.cpp +++ b/src/amd_detail/backend/asyncop-fallback.cpp @@ -5,6 +5,7 @@ #include "async.h" #include "asyncop-fallback.h" +#include "backend.h" #include "buffer.h" #include "context.h" #include "hip.h" @@ -42,10 +43,11 @@ AsyncOpFallback::AsyncOpFallback(IoType _io_type, std::shared_ptr _file, size_t *_size, hoff_t *_file_offset, hoff_t *_buffer_offset, ssize_t *_bytes_transferred) : AsyncOp{_io_type, _file, _buffer, _stream, _size, _file_offset, _buffer_offset, _bytes_transferred}, - submitted_size{*_size}, bytes_transferred_internal{0}, gpu_buffer{buffer->getBuffer()}, - bounce_buffer_dev_ptr{nullptr}, bounce_buffer{nullptr, [](void *addr) { (void)addr; }} + submitted_size{std::min(*_size, hipFile::MAX_RW_COUNT)}, bytes_transferred_internal{0}, + gpu_buffer{buffer->getBuffer()}, bounce_buffer_dev_ptr{nullptr}, + bounce_buffer{nullptr, [](void *addr) { (void)addr; }} { - void *host_ptr = Context::get()->hipHostMalloc(*_size, 0); + void *host_ptr = Context::get()->hipHostMalloc(submitted_size, 0); std::unique_ptr _bounce_buffer{host_ptr, hipHostDeleter}; std::swap(bounce_buffer, _bounce_buffer); void *dev_ptr = Context::get()->hipHostGetDevicePointer(bounce_buffer.get(), 0); diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index 4fe31baa..6745fbc2 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -122,7 +122,9 @@ Fallback::async_io(IoType type, std::shared_ptr file, std::shared_ptr stream) { - if (!paramsValid(buffer, *size_p, *file_offset_p, *buffer_offset_p)) { + size_t limited_size = min(*size_p, hipFile::MAX_RW_COUNT); + + if (!paramsValid(buffer, limited_size, *file_offset_p, *buffer_offset_p)) { throw std::invalid_argument("The selected file or buffer region is invalid"); } @@ -199,7 +201,7 @@ async_io_bind_params(void *userargs) const hoff_t *file_offset = get_variant_ptr(op->file_offset); op->file_offset.emplace(*file_offset); const size_t *size = get_variant_ptr(op->size); - op->size = *size; + op->size = std::min(*size, hipFile::MAX_RW_COUNT); if (std::get(op->size) > op->submitted_size) { op->bytes_transferred_internal = -hipFileInvalidValue; diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 512947b2..3f0005df 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -24,6 +24,7 @@ #include "state.h" #include +#include #include #include #include @@ -153,6 +154,24 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_new_uses_pinned_host_memory) IoType::Read, file, buffer, stream, &size, &file_offset, &buffer_offset, &bytes_transferred}); } +TEST_F(HipFileAsyncOp, AsyncOpFallbackLimitsMaxIoSize) +{ + size_t size = 4_GiB; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + hoff_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + auto bounce_buffer = std::shared_ptr(new uint8_t[1_KiB]); + EXPECT_CALL(mhip, hipHostMalloc(hipFile::MAX_RW_COUNT, _)).WillOnce(Return(bounce_buffer.get())); + EXPECT_CALL(mhip, hipHostMalloc(sizeof(AsyncOpFallback), _)).WillOnce(Return(op_data.get())); + EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); + EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); + EXPECT_CALL(mhip, hipHostFree(Eq(op_data.get()))); + auto op = std::shared_ptr(new AsyncOpFallback{ + IoType::Read, file, buffer, stream, &size, &file_offset, &buffer_offset, &bytes_transferred}); + ASSERT_EQ(op->submitted_size, hipFile::MAX_RW_COUNT); +} + TEST_F(HipFileAsyncOp, AsyncOpFallback_new_failure_throws_bad_alloc) { size_t size = 100; @@ -438,6 +457,19 @@ TEST_P(FallbackAsyncIO, sizeTooLargeReturnsError) std::invalid_argument); } +TEST_P(FallbackAsyncIO, paramsValidUsesLimitedSize) +{ + size = 4_GiB; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(2_GiB)); + // Use GPU mismatch to exit function early. These functions are only called because the size was limited + // and now does not exceed the buffer size. + EXPECT_CALL(*mbuffer, getGpuId).WillOnce(Return(0)); + EXPECT_CALL(*mstream, getHipDevice).WillOnce(Return(1)); + EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, + &bytes_written, mstream), + std::invalid_argument); +} + TEST_P(FallbackAsyncIO, fileOffsetNegativeReturnsError) { file_offset = -1; @@ -597,6 +629,21 @@ TEST_F(AsyncIoOpCleanup, cleanupInvalidOpSetsError) ASSERT_EQ(*op->bytes_transferred, -hipFileInternalError); } +struct AsyncIoOpLimitedSize : public AsyncIoOp { + void SetUp() override + { + size = hipFile::MAX_RW_COUNT + 1; + AsyncIoOp::SetUp(); + } +}; + +TEST_F(AsyncIoOpLimitedSize, bindLimitsSize) +{ + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(hipFile::MAX_RW_COUNT)); + async_io_bind_params(op.get()); + ASSERT_EQ(std::get(op->size), hipFile::MAX_RW_COUNT); +} + struct AsyncIoOpBindParams { IoType io_type; bool fixed_buffer_offset; From a880bd9ef6efff99da26f221e02bca7f297905a4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 4 Feb 2026 17:09:16 -0500 Subject: [PATCH 20/27] Add EINTR unit test. --- test/amd_detail/async.cpp | 21 ++++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 3f0005df..8b32671a 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -25,6 +25,7 @@ #include #include +#include #include #include #include @@ -683,7 +684,7 @@ TEST_P(AsyncIoOpWithParams, cpuIOReturnsFullSize) ASSERT_EQ(op->bytes_transferred_internal, size); } -TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadErrorReturnsError) +TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadPwriteErrorReturnsError) { if (io_type == IoType::Read) { EXPECT_CALL(msys, pread).WillOnce(Throw(std::system_error(3, std::generic_category()))); @@ -696,6 +697,24 @@ TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadErrorReturnsError) ASSERT_EQ(op->bytes_transferred_internal, -1); } +TEST_P(AsyncIoOpWithParams, cpuCopyReadPreadPwriteRetriesOnEINTR) +{ + if (io_type == IoType::Read) { + EXPECT_CALL(msys, pread) + .WillOnce(Throw(std::system_error(EINTR, std::generic_category()))) + .WillOnce(Return(size)); + } + else { + EXPECT_CALL(msys, pwrite) + .WillOnce(Throw(std::system_error(EINTR, std::generic_category()))) + .WillOnce(Return(size)); + } + EXPECT_CALL(*mfile, getBufferedFd).Times(2); + async_io_cpu_copy(op.get()); + + ASSERT_EQ(op->bytes_transferred_internal, size); +} + INSTANTIATE_TEST_SUITE_P(AsyncIoOpWithParamsSuite, AsyncIoOpWithParams, ::testing::Values(AsyncIoOpBindParams{IoType::Read, true, true, true}, AsyncIoOpBindParams{IoType::Write, true, true, true})); From a7d13ce77e9d3d26a2333aefb53dd22614a789a4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Wed, 4 Feb 2026 17:23:08 -0500 Subject: [PATCH 21/27] fixup! hipFile: Adding async unit tests. review: Can use same buffer_offset and buffer length for too large offset test. --- test/amd_detail/async.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 8b32671a..7a9125e2 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -442,8 +442,8 @@ TEST_P(FallbackAsyncIO, bufferOffsetNegativeReturnsError) TEST_P(FallbackAsyncIO, bufferOffsetTooLargeReturnsError) { - buffer_offset = 1024 * 1024 + 1; - EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(1024 * 1024)); + buffer_offset = 1024 * 1024; + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(buffer_offset)); EXPECT_THROW(Fallback().async_io(io_type, mfile, mbuffer, &size, &file_offset, &buffer_offset, &bytes_written, mstream), std::invalid_argument); From c3bfa834844f2cee2d15d06787f757fd187836c4 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Thu, 5 Feb 2026 12:49:58 -0500 Subject: [PATCH 22/27] Fix isGpuMemory check on NVIDIA. --- test/system/async.cpp | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/test/system/async.cpp b/test/system/async.cpp index fb2012e8..7a9c4329 100644 --- a/test/system/async.cpp +++ b/test/system/async.cpp @@ -61,7 +61,15 @@ static bool isGpuMemory(void *mem) { hipPointerAttribute_t attrs; - assert(hipPointerGetAttributes(&attrs, mem) == hipSuccess); + hipError_t err = hipPointerGetAttributes(&attrs, mem); + +#ifdef __HIP_PLATFORM_NVIDIA__ + // NVIDIA doesn't support pointers allocated outside of runtime + if (err == hipErrorInvalidValue) { + return false; + } +#endif + assert(err == hipSuccess); return attrs.type == hipMemoryTypeDevice; } From f29f3f38767f73d28bebb356fd54168bcf6c6b6c Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Fri, 6 Feb 2026 15:33:48 -0500 Subject: [PATCH 23/27] review: Add comments to async parameter tests and show cuFile error returned later. --- test/system/async.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/test/system/async.cpp b/test/system/async.cpp index 7a9c4329..c2de014f 100644 --- a/test/system/async.cpp +++ b/test/system/async.cpp @@ -498,6 +498,7 @@ TEST_P(HipAsyncReadWriteStreamFixedWithParams, nullBytesReadReturnsError) HipFileOpError(hipFileInvalidValue)); } +// cuFile accepts operation and returns error through bytes_transferred TEST_P(HipAsyncReadWriteStreamFixedWithParams, ioSizeTooLargeForBuffer) { io_size += 4096; @@ -507,9 +508,12 @@ TEST_P(HipAsyncReadWriteStreamFixedWithParams, ioSizeTooLargeForBuffer) #else ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidMappingSize); #endif } +// cuFile accepts operation and returns error through bytes_transferred TEST_P(HipAsyncReadWriteStreamFixedWithParams, fileOffsetNegative) { file_offset = -4096; @@ -519,9 +523,12 @@ TEST_P(HipAsyncReadWriteStreamFixedWithParams, fileOffsetNegative) #else ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidValue); #endif } +// cuFile accepts operation and returns error through bytes_transferred TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetNegative) { buffer_offset = -4096; @@ -531,9 +538,12 @@ TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetNegative) #else ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInternalError); #endif } +// cuFile accepts operation and returns error through bytes_transferred TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetTooLarge) { buffer_offset = static_cast(buffer_size) + 4096; @@ -543,6 +553,8 @@ TEST_P(HipAsyncReadWriteStreamFixedWithParams, bufferOffsetTooLarge) #else ASSERT_EQ(io_op(fh, dev_ptr, &io_size, &file_offset, &buffer_offset, &bytes_transferred, stream), HIPFILE_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidMappingSize); #endif } From a1576e5267a78765fce60f9abbbcbdc1c974a2fb Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Fri, 6 Feb 2026 16:03:55 -0500 Subject: [PATCH 24/27] Fix backend.h header include. --- test/amd_detail/async.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index 7a9125e2..ac53a849 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -4,6 +4,7 @@ */ #include "async.h" +#include "backend.h" #include "backend/asyncop-fallback.h" #include "backend/fallback.h" #include "context.h" @@ -24,7 +25,6 @@ #include "state.h" #include -#include #include #include #include From d2a2e067e7e28321f35886afdc00f367d1b8e0a6 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Fri, 6 Feb 2026 16:04:50 -0500 Subject: [PATCH 25/27] Fix accidental bracket in error message. --- test/common/test-common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/common/test-common.h b/test/common/test-common.h index a257446d..95debab6 100644 --- a/test/common/test-common.h +++ b/test/common/test-common.h @@ -69,7 +69,7 @@ struct Tmpfile { #ifdef __HIP_PLATFORM_AMD__ if (unlink(path.c_str()) == -1) { - throw std::runtime_error("Could not unlink temporary file)"); + throw std::runtime_error("Could not unlink temporary file"); } #endif } From b50d9b3b6223dc01230f42eeb335760cd8def109 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Tue, 10 Feb 2026 16:42:24 -0500 Subject: [PATCH 26/27] Remove unused variable. --- src/amd_detail/backend/fallback.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index 6745fbc2..aa386f31 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -221,7 +221,7 @@ async_io_cleanup(void *userargs) try { Context::get()->completeOp(op); } - catch (std::invalid_argument &e) { + catch (const std::invalid_argument &) { *bytes_transferred = -hipFileInternalError; return; } From feaec9777d1a836d2544a4a732497e3ec885fb16 Mon Sep 17 00:00:00 2001 From: Jordan Patterson Date: Tue, 10 Feb 2026 17:03:57 -0500 Subject: [PATCH 27/27] Fix bytes transferred type in unit tests. --- test/amd_detail/async.cpp | 118 +++++++++++++++++++------------------- 1 file changed, 59 insertions(+), 59 deletions(-) diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index ac53a849..39930bc5 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -110,12 +110,12 @@ struct HipFileAsyncOpStreamParams TEST_P(HipFileAsyncOpStreamParams, asyncOp_construction_has_correct_variants) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op = std::make_shared(IoType::Read, file, buffer, stream, &size, &file_offset, - &buffer_offset, &bytes_transferred); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op = std::make_shared(IoType::Read, file, buffer, stream, &size, &file_offset, + &buffer_offset, &bytes_transferred); // Unfixed flags will be pointers if (flags & HIPFILE_STREAM_FIXED_BUF_OFFSET) { @@ -141,12 +141,12 @@ INSTANTIATE_TEST_SUITE_P(StreamSuite, HipFileAsyncOpStreamParams, hipfileFlagsPo TEST_F(HipFileAsyncOp, AsyncOpFallback_new_uses_pinned_host_memory) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); - auto bounce_buffer = std::shared_ptr(new uint8_t[size]); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + auto bounce_buffer = std::shared_ptr(new uint8_t[size]); EXPECT_CALL(mhip, hipHostMalloc).WillOnce(Return(op_data.get())).WillOnce(Return(bounce_buffer.get())); EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); @@ -157,12 +157,12 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_new_uses_pinned_host_memory) TEST_F(HipFileAsyncOp, AsyncOpFallbackLimitsMaxIoSize) { - size_t size = 4_GiB; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); - auto bounce_buffer = std::shared_ptr(new uint8_t[1_KiB]); + size_t size = 4_GiB; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + auto bounce_buffer = std::shared_ptr(new uint8_t[1_KiB]); EXPECT_CALL(mhip, hipHostMalloc(hipFile::MAX_RW_COUNT, _)).WillOnce(Return(bounce_buffer.get())); EXPECT_CALL(mhip, hipHostMalloc(sizeof(AsyncOpFallback), _)).WillOnce(Return(op_data.get())); EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); @@ -175,11 +175,11 @@ TEST_F(HipFileAsyncOp, AsyncOpFallbackLimitsMaxIoSize) TEST_F(HipFileAsyncOp, AsyncOpFallback_new_failure_throws_bad_alloc) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); EXPECT_CALL(mhip, hipHostMalloc).WillOnce(Throw(Hip::RuntimeError(hipErrorOutOfMemory))); EXPECT_THROW(std::shared_ptr(new AsyncOpFallback{IoType::Read, file, buffer, stream, &size, &file_offset, &buffer_offset, @@ -189,11 +189,11 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_new_failure_throws_bad_alloc) TEST_F(HipFileAsyncOp, AsyncOpFallback_bounce_alloc_failure_throws) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); EXPECT_CALL(mhip, hipHostMalloc) .WillOnce(Return(op_data.get())) .WillOnce(Throw(Hip::RuntimeError(hipErrorOutOfMemory))); @@ -206,12 +206,12 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_bounce_alloc_failure_throws) TEST_F(HipFileAsyncOp, AsyncOpFallback_bounce_buffer_deleter_failure_calls_syslog) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); - auto bounce_buffer = std::shared_ptr(new uint8_t[size]); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + auto bounce_buffer = std::shared_ptr(new uint8_t[size]); EXPECT_CALL(mhip, hipHostMalloc).WillOnce(Return(op_data.get())).WillOnce(Return(bounce_buffer.get())); EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))) @@ -224,12 +224,12 @@ TEST_F(HipFileAsyncOp, AsyncOpFallback_bounce_buffer_deleter_failure_calls_syslo TEST_F(HipFileAsyncOp, AsyncOpFallback_delete_failure_calls_syslog) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); - auto bounce_buffer = std::shared_ptr(new uint8_t[size]); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op_data = std::shared_ptr(new uint8_t[sizeof(AsyncOpFallback)]); + auto bounce_buffer = std::shared_ptr(new uint8_t[size]); EXPECT_CALL(mhip, hipHostMalloc).WillOnce(Return(op_data.get())).WillOnce(Return(bounce_buffer.get())); EXPECT_CALL(mhip, hipHostGetDevicePointer(Eq(bounce_buffer.get()), _)); EXPECT_CALL(mhip, hipHostFree(Eq(bounce_buffer.get()))); @@ -258,7 +258,7 @@ struct HipFileAsyncOpFallbackMethods : public HipFileAsyncOp { size_t size = 100; hoff_t file_offset = 0; hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; + ssize_t bytes_transferred = 0; std::shared_ptr bounce_buffer; std::shared_ptr op; }; @@ -281,12 +281,12 @@ struct HipFileAsyncMonitor : HipFileAsyncOp { TEST_F(HipFileAsyncMonitor, addOp_and_completeOp_with_valid_params_works) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op = std::make_shared(IoType::Read, file, buffer, stream, &size, &file_offset, - &buffer_offset, &bytes_transferred); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op = std::make_shared(IoType::Read, file, buffer, stream, &size, &file_offset, + &buffer_offset, &bytes_transferred); monitor.addOp(op); EXPECT_NO_THROW(monitor.completeOp(op.get())); @@ -299,12 +299,12 @@ TEST_F(HipFileAsyncMonitor, completeOp_with_invalid_op_throws) TEST_F(HipFileAsyncMonitor, addOp_without_completeOp_prints_error_on_AsyncMonitor_destruction) { - size_t size = 100; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; - auto op = std::make_unique(IoType::Read, file, buffer, stream, &size, &file_offset, - &buffer_offset, &bytes_transferred); + size_t size = 100; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_transferred = 0; + auto op = std::make_unique(IoType::Read, file, buffer, stream, &size, &file_offset, + &buffer_offset, &bytes_transferred); monitor.addOp(std::move(op)); EXPECT_CALL(msys, syslog); } @@ -370,10 +370,10 @@ TEST_P(HipFileReadWriteAsync, nullBytesTransferredReturnsError) TEST_P(HipFileReadWriteAsync, unregisteredFileReturnsError) { - size_t size = 1; - hoff_t file_offset = 0; - hoff_t buffer_offset = 0; - hoff_t bytes_written = 0; + size_t size = 1; + hoff_t file_offset = 0; + hoff_t buffer_offset = 0; + ssize_t bytes_written = 0; ASSERT_EQ(io_op(nonnull_void, nonnull_void, &size, &file_offset, &buffer_offset, &bytes_written, nonnull_stream), @@ -428,7 +428,7 @@ struct FallbackAsyncIO : public HipFileOpened, public ::testing::WithParamInterf size_t size; hoff_t file_offset; hoff_t buffer_offset; - hoff_t bytes_written; + ssize_t bytes_written; }; TEST_P(FallbackAsyncIO, bufferOffsetNegativeReturnsError) @@ -568,7 +568,7 @@ struct AsyncIoOp : public ::testing::Test { size_t size = 1_MiB; hoff_t file_offset = 0; hoff_t buffer_offset = 0; - hoff_t bytes_transferred = 0; + ssize_t bytes_transferred = 0; bool fixed_buffer_offset = false; bool fixed_file_offset = false; bool fixed_io_size = false;