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' }} \ 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; +} diff --git a/src/amd_detail/CMakeLists.txt b/src/amd_detail/CMakeLists.txt index 247f4e3f..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 @@ -19,6 +20,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/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/src/amd_detail/backend/asyncop-fallback.cpp b/src/amd_detail/backend/asyncop-fallback.cpp index 7c1976d2..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}, - bytes_transferred_internal{0}, gpu_buffer{buffer->getBuffer()}, bounce_buffer_dev_ptr{nullptr}, + 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/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; diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index 379b1a25..aa386f31 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; @@ -51,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, @@ -121,3 +116,162 @@ 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) +{ + 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"); + } + + 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 = std::min(*size, hipFile::MAX_RW_COUNT); + + 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 (const std::invalid_argument &) { + *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; + } + + 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, + 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"); + } + } + catch (const std::system_error &e) { + if (e.code().value() == EINTR) { + continue; + } + op->bytes_transferred_internal = -1; + return; + } + if (ret == 0) { + break; + } + bytes_transferred += static_cast(ret); + } + op->bytes_transferred_internal = static_cast(bytes_transferred); +} +} 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/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/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..f70efe51 --- /dev/null +++ b/src/amd_detail/backend/memcpy-kernel.hip @@ -0,0 +1,104 @@ +/* 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) +{ + // 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) + 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; + } + } +} + +} 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/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/io.cpp b/src/amd_detail/io.cpp new file mode 100644 index 00000000..70d1b44c --- /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); + } diff --git a/src/amd_detail/stream.cpp b/src/amd_detail/stream.cpp index 2815b447..aef4e934 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" @@ -10,6 +11,7 @@ #include #include +#include #include #include #include @@ -17,7 +19,7 @@ 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}, + : 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} @@ -26,6 +28,8 @@ Stream::Stream(const hipStream_t _hip_stream, uint32_t flags, const PassKey::get()->hipStreamGetDevice(hip_stream); } hipStream_t @@ -33,6 +37,11 @@ Stream::getHipStream() const { return hip_stream; } +hipDevice_t +Stream::getHipDevice() const +{ + return device_id; +} bool Stream::fixedBufferOffset() const { @@ -53,6 +62,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..6ef04877 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,13 @@ 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 hipDevice_t getHipDevice() 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 +33,13 @@ 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 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); @@ -45,10 +50,12 @@ 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; bool page_aligned; + std::mutex mutex; }; class StreamMap { 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; + } +} + +} 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/amd_detail/async.cpp b/test/amd_detail/async.cpp index 9a0c18bc..39930bc5 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -4,17 +4,28 @@ */ #include "async.h" +#include "backend.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 #include @@ -23,10 +34,11 @@ #include #include #include -#include +#include #include #include #include +#include // Put tests inside the macros to suppress the global constructor // warnings @@ -45,8 +57,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 +78,6 @@ struct HipFileAsyncOp : public Test { shared_ptr stream; }; -struct HipFileAsyncMonitor : HipFileAsyncOp { - AsyncMonitor monitor; -}; - static auto hipfileFlagsPowerSet() { @@ -104,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) { @@ -135,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()))); @@ -149,13 +155,31 @@ 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; + 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()), _)); + 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; - 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, @@ -165,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))); @@ -182,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()))) @@ -200,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()))); @@ -216,8 +240,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 +250,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()))); } @@ -234,31 +258,35 @@ struct HipFileAsyncOpFallbackFunctions : 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; }; -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; - 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())); @@ -271,14 +299,424 @@ 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); } +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; + ssize_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; + ssize_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; + 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); +} + +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, 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; + 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; + ssize_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 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; + 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, cpuCopyReadPreadPwriteErrorReturnsError) +{ + 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, 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})); + HIPFILE_WARN_NO_GLOBAL_CTOR_ON 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)); +}; + +} 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)); }; } diff --git a/test/amd_detail/mstream.h b/test/amd_detail/mstream.h index 081e94b3..c53ed440 100644 --- a/test/amd_detail/mstream.h +++ b/test/amd_detail/mstream.h @@ -6,15 +6,19 @@ #include "stream.h" +#include + 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)); MOCK_METHOD(bool, pageAligned, (), (const, override)); + MOCK_METHOD(std::unique_lock, getLock, (), (override)); }; } diff --git a/test/amd_detail/stream.cpp b/test/amd_detail/stream.cpp index 18e8da84..d137f3e1 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,6 +70,7 @@ 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()); } @@ -81,18 +83,21 @@ TEST_F(HipFileStream, register_with_invalid_flags_throws) 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 +115,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 +134,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); } diff --git a/test/common/test-common.h b/test/common/test-common.h index fdd3a55a..95debab6 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) { - throw std::runtime_error("Could not unlink temporary file)"); + +#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); } }; diff --git a/test/system/async.cpp b/test/system/async.cpp new file mode 100644 index 00000000..c2de014f --- /dev/null +++ b/test/system/async.cpp @@ -0,0 +1,892 @@ +/* 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; + 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; +} + +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)); +} + +// cuFile accepts operation and returns error through bytes_transferred +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); + 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; +#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); + 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; +#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); + 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; +#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); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(bytes_transferred, -hipFileInvalidMappingSize); +#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