diff --git a/CHANGELOG.md b/CHANGELOG.md index f1119aa2..60624355 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -12,6 +12,7 @@ * The CMake namespace was changed from `roc::` to `hip::` * `AIS_BUILD_EXAMPLES` has been renamed to `AIS_INSTALL_EXAMPLES` * `AIS_USE_SANITIZERS` now also enables the following sanitizers: integer, float-divide-by-zero, local-bounds, vptr, nullability (in addition to address, leak, and undefined). Sanitizers should also now emit usable stack trace info. +* The AIS optimized IO path will automatically fallback to the POSIX IO path if a failure occurs and the compatability mode has not been disabled. ### Removed * The rocFile library has been completely removed and the code is now a part of hipFile. diff --git a/src/amd_detail/CMakeLists.txt b/src/amd_detail/CMakeLists.txt index c264f630..1153b87a 100644 --- a/src/amd_detail/CMakeLists.txt +++ b/src/amd_detail/CMakeLists.txt @@ -7,6 +7,7 @@ set(HIPFILE_SOURCES "${HIPFILE_SRC_COMMON_PATH}/hipfile-common.cpp" async.cpp + backend.cpp backend/asyncop-fallback.cpp backend/memcpy-kernel.hip backend/fallback.cpp diff --git a/src/amd_detail/backend.cpp b/src/amd_detail/backend.cpp new file mode 100644 index 00000000..c9fc5725 --- /dev/null +++ b/src/amd_detail/backend.cpp @@ -0,0 +1,89 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#include "backend.h" +#include "buffer.h" +#include "file.h" +#include "io.h" + +#include +#include +#include +#include +#include +#include + +using namespace hipFile; + +ssize_t +Backend::io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset) +{ + ssize_t nbytes = _io_impl(type, file, buffer, size, file_offset, buffer_offset); + switch (type) { + case (IoType::Read): + update_read_stats(nbytes); + break; + case (IoType::Write): + update_write_stats(nbytes); + break; + default: + break; + } + return nbytes; +} + +ssize_t +BackendWithFallback::io(IoType type, std::shared_ptr file, std::shared_ptr buffer, + size_t size, hoff_t file_offset, hoff_t buffer_offset) +{ + ssize_t nbytes{0}; + try { + nbytes = _io_impl(type, file, buffer, size, file_offset, buffer_offset); + if (nbytes < 0) { + // Typically we should not reach this point. But in case we do, throw + // an exception to use the fallback backend. + throw std::system_error(-static_cast(nbytes), std::generic_category()); + } + } + catch (...) { + std::exception_ptr e_ptr = std::current_exception(); + if (is_fallback_eligible(e_ptr, nbytes, file, buffer, size, file_offset, buffer_offset)) { + nbytes = fallback_backend->io(type, file, buffer, size, file_offset, buffer_offset); + } + else { + throw; + } + return nbytes; + } + switch (type) { + case (IoType::Read): + update_read_stats(nbytes); + break; + case (IoType::Write): + update_write_stats(nbytes); + break; + default: + break; + } + return nbytes; +} + +bool +BackendWithFallback::is_fallback_eligible(std::exception_ptr e_ptr, ssize_t nbytes, + std::shared_ptr file, std::shared_ptr buffer, + size_t size, hoff_t file_offset, hoff_t buffer_offset) const +{ + (void)e_ptr; + (void)nbytes; + return static_cast(fallback_backend) && + fallback_backend->score(file, buffer, size, file_offset, buffer_offset) >= 0; +} + +void +BackendWithFallback::register_fallback_backend(std::shared_ptr backend) noexcept +{ + fallback_backend = backend; +} diff --git a/src/amd_detail/backend.h b/src/amd_detail/backend.h index c2e040df..cfc0d289 100644 --- a/src/amd_detail/backend.h +++ b/src/amd_detail/backend.h @@ -12,6 +12,7 @@ #include "sys.h" #include +#include #include #include #include @@ -51,11 +52,74 @@ struct Backend { /// @param file_offset Offset from the start of the file /// @param buffer_offset Offset from the start of the buffer /// - /// @return Number of bytes transferred, negative on error + /// @return Number of bytes transferred /// /// @throws Hip::RuntimeError Sys::RuntimeError virtual ssize_t io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, - hoff_t file_offset, hoff_t buffer_offset) = 0; + hoff_t file_offset, hoff_t buffer_offset); + + /// @brief Update the read stats for this Backend + /// + /// @param nbytes Number of bytes read + virtual void update_read_stats(ssize_t nbytes) = 0; + + /// @brief Update the write stats for this Backend + /// + /// @param nbytes Number of bytes written + virtual void update_write_stats(ssize_t nbytes) = 0; + +protected: + /// @brief Perform a read or write operation + /// + /// @note Provides a common target across all Backends that provides the + /// implementation for running IO. + /// @param type IO type (read/write) + /// @param file File to read from or write to + /// @param buffer Buffer to write to or read from + /// @param size Number of bytes to transfer + /// @param file_offset Offset from the start of the file + /// @param buffer_offset Offset from the start of the buffer + /// + /// @return Number of bytes transferred + /// + /// @throws Hip::RuntimeError Sys::RuntimeError + virtual ssize_t _io_impl(IoType type, std::shared_ptr file, std::shared_ptr buffer, + size_t size, hoff_t file_offset, hoff_t buffer_offset) = 0; +}; + +// BackendWithFallback allows for an IO to be retried automatically with a +// different Backend in the event of an error. +struct BackendWithFallback : 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 final; + + /// @brief Check if a failed IO operation can be re-issued to the fallback Backend. + /// + /// @param e_ptr exception_ptr to the thrown exception from the failed IO + /// @param nbytes Return value from `_io_impl`, or 0 if an exception was thrown. + /// @param file File to read from or write to + /// @param buffer Buffer to write to or read from + /// @param size Number of bytes to transfer + /// @param file_offset Offset from the start of the file + /// @param buffer_offset Offset from the start of the buffer + /// + /// @note By default, BackendWithFallback checks if a Backend has been + /// registered for retrying an IO, and that fallback backend supports + /// the request. + /// @note The parameters from the original IO request are passed to this function. + /// + /// @return True if this BackendWithFallback can retry the IO, else False. + virtual bool is_fallback_eligible(std::exception_ptr e_ptr, ssize_t nbytes, std::shared_ptr file, + std::shared_ptr buffer, size_t size, hoff_t file_offset, + hoff_t buffer_offset) const; + + /// @brief Register a Backend to retry a failed IO operation. + /// + /// @param backend Backend to retry a failed IO operation. + void register_fallback_backend(std::shared_ptr backend) noexcept; + +protected: + std::shared_ptr fallback_backend; }; } diff --git a/src/amd_detail/backend/fallback.cpp b/src/amd_detail/backend/fallback.cpp index 1d2caa73..2e33e05a 100644 --- a/src/amd_detail/backend/fallback.cpp +++ b/src/amd_detail/backend/fallback.cpp @@ -58,8 +58,33 @@ Fallback::io(IoType type, std::shared_ptr file, std::shared_ptr } ssize_t -Fallback::io(IoType io_type, shared_ptr file, shared_ptr buffer, size_t size, +Fallback::io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, hoff_t file_offset, hoff_t buffer_offset, size_t chunk_size) +{ + ssize_t nbytes = _io_impl(type, file, buffer, size, file_offset, buffer_offset, chunk_size); + switch (type) { + case (IoType::Read): + update_read_stats(nbytes); + break; + case (IoType::Write): + update_write_stats(nbytes); + break; + default: + break; + } + return nbytes; +} + +ssize_t +Fallback::_io_impl(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset) +{ + return _io_impl(type, file, buffer, size, file_offset, buffer_offset, DefaultChunkSize); +} + +ssize_t +Fallback::_io_impl(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 = min(size, hipFile::MAX_RW_COUNT); @@ -115,19 +140,21 @@ Fallback::io(IoType io_type, shared_ptr file, shared_ptr buffer, } } while (static_cast(total_io_bytes) < size); - switch (io_type) { - case IoType::Read: - statsAddFallbackPathRead(static_cast(total_io_bytes)); - break; - case IoType::Write: - statsAddFallbackPathWrite(static_cast(total_io_bytes)); - break; - default: - break; - } return total_io_bytes; } +void +Fallback::update_read_stats(ssize_t nbytes) +{ + statsAddFallbackPathRead(static_cast(nbytes)); +} + +void +Fallback::update_write_stats(ssize_t nbytes) +{ + statsAddFallbackPathWrite(static_cast(nbytes)); +} + 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, diff --git a/src/amd_detail/backend/fallback.h b/src/amd_detail/backend/fallback.h index 14e3b483..d35c068a 100644 --- a/src/amd_detail/backend/fallback.h +++ b/src/amd_detail/backend/fallback.h @@ -34,16 +34,24 @@ 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 update_read_stats(ssize_t nbytes) override; + + void update_write_stats(ssize_t nbytes) 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: - ssize_t io(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, hoff_t file_offset, hoff_t buffer_offset, size_t chunk_size); + +protected: + ssize_t _io_impl(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset) override; + ssize_t _io_impl(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset, size_t chunk_size); }; } diff --git a/src/amd_detail/backend/fastpath.cpp b/src/amd_detail/backend/fastpath.cpp index aba587c8..fa06eeac 100644 --- a/src/amd_detail/backend/fastpath.cpp +++ b/src/amd_detail/backend/fastpath.cpp @@ -154,9 +154,21 @@ Fastpath::score(shared_ptr file, shared_ptr buffer, size_t size, return accept_io ? 100 : -1; } +void +Fastpath::update_read_stats(ssize_t nbytes) +{ + statsAddFastPathRead(static_cast(nbytes)); +} + +void +Fastpath::update_write_stats(ssize_t nbytes) +{ + statsAddFastPathWrite(static_cast(nbytes)); +} + ssize_t -Fastpath::io(IoType type, shared_ptr file, shared_ptr buffer, size_t size, hoff_t file_offset, - hoff_t buffer_offset) +Fastpath::_io_impl(IoType type, shared_ptr file, shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset) { void *devptr{reinterpret_cast(reinterpret_cast(buffer->getBuffer()) + buffer_offset)}; hipAmdFileHandle_t handle{}; @@ -184,15 +196,5 @@ Fastpath::io(IoType type, shared_ptr file, shared_ptr buffer, si default: throw std::runtime_error("Invalid IoType"); } - switch (type) { - case IoType::Read: - statsAddFastPathRead(nbytes); - break; - case IoType::Write: - statsAddFastPathWrite(nbytes); - break; - default: - break; - } return static_cast(nbytes); } diff --git a/src/amd_detail/backend/fastpath.h b/src/amd_detail/backend/fastpath.h index 61f0915e..87e72433 100644 --- a/src/amd_detail/backend/fastpath.h +++ b/src/amd_detail/backend/fastpath.h @@ -9,6 +9,7 @@ #include "hipfile.h" #include +#include #include namespace hipFile { @@ -23,14 +24,17 @@ enum class IoType; namespace hipFile { -struct Fastpath : public Backend { +struct Fastpath : public BackendWithFallback { virtual ~Fastpath() override = default; - int score(std::shared_ptr file, std::shared_ptr buffer, size_t size, hoff_t file_offset, - hoff_t buffer_offset) const override; + int score(std::shared_ptr file, std::shared_ptr buffer, size_t size, hoff_t file_offset, + hoff_t buffer_offset) const override; + void update_read_stats(ssize_t nbytes) override; + void update_write_stats(ssize_t nbytes) override; - 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; +protected: + ssize_t _io_impl(IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset) override; }; } diff --git a/src/amd_detail/state.cpp b/src/amd_detail/state.cpp index fd4aacfa..19cb29aa 100644 --- a/src/amd_detail/state.cpp +++ b/src/amd_detail/state.cpp @@ -275,11 +275,18 @@ std::vector> DriverState::getBackends() const { static bool once = [&]() { - if (Context::get()->fastpath()) { - backends.emplace_back(new Fastpath{}); - } + std::shared_ptr fallback_backend; if (Context::get()->fallback()) { - backends.emplace_back(new Fallback{}); + fallback_backend = std::make_shared(); + backends.push_back(fallback_backend); + } + + if (Context::get()->fastpath()) { + auto new_backend = std::make_shared(); + if (fallback_backend) { + new_backend->register_fallback_backend(fallback_backend); + } + backends.push_back(new_backend); } return true; }(); diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 0a0869be..cf508bad 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -24,6 +24,7 @@ set(SYSTEM_TEST_SOURCE_FILES system/buffer.cpp system/config.cpp system/driver.cpp + system/fastpath.cpp system/io.cpp system/main.cpp system/version.cpp @@ -32,6 +33,7 @@ set(SYSTEM_TEST_SOURCE_FILES set(TEST_SYSINCLS ${HIPFILE_INCLUDE_PATH} ${HIPFILE_TEST_COMMON_PATH} + ${HIPFILE_AMD_TEST_PATH} # For mhip ) if(AIS_BUILD_NVIDIA_DETAIL) @@ -66,6 +68,7 @@ ais_add_executable( SYSINCLS ${TEST_SYSINCLS} ) target_link_libraries(hipfile_system_tests PRIVATE GTest::gtest) +target_link_libraries(hipfile_system_tests PRIVATE GTest::gmock) target_link_libraries(hipfile_system_tests PRIVATE Boost::program_options) ais_gtest_discover_tests( diff --git a/test/amd_detail/CMakeLists.txt b/test/amd_detail/CMakeLists.txt index c44514f5..e42f17a9 100644 --- a/test/amd_detail/CMakeLists.txt +++ b/test/amd_detail/CMakeLists.txt @@ -12,6 +12,7 @@ set(SHARED_SOURCE_FILES set(TEST_SOURCE_FILES async.cpp + backend.cpp batch/batch.cpp configuration.cpp context.cpp diff --git a/test/amd_detail/async.cpp b/test/amd_detail/async.cpp index f1801c6c..45f09b2c 100644 --- a/test/amd_detail/async.cpp +++ b/test/amd_detail/async.cpp @@ -23,6 +23,7 @@ #include "mstream.h" #include "msys.h" #include "state.h" +#include "test-common.h" #include #include @@ -386,7 +387,7 @@ TEST_P(HipFileReadWriteAsync, badAllocReturnsHipErrorOutOfMemory) 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)); + HipFileDriverError(hipErrorOutOfMemory)); } INSTANTIATE_TEST_SUITE_P(HipFileAsyncSuite, HipFileReadWriteAsync, ::testing::ValuesIn(asyncIOFns), diff --git a/test/amd_detail/backend.cpp b/test/amd_detail/backend.cpp new file mode 100644 index 00000000..d0c00d0f --- /dev/null +++ b/test/amd_detail/backend.cpp @@ -0,0 +1,103 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#include "io.h" +#include "hipfile-test.h" +#include "hipfile-warnings.h" +#include "mbackend.h" +#include "mbuffer.h" +#include "mfile.h" + +#include +#include +#include + +using namespace hipFile; + +using ::testing::AnyNumber; +using ::testing::Return; +using ::testing::StrictMock; +using ::testing::Throw; + +// Put tests inside the macros to suppress the global constructor +// warnings +HIPFILE_WARN_NO_GLOBAL_CTOR_OFF + +struct DummyBackendWithFallback : public MBackendWithFallback {}; + +struct DummyFallbackBackend : MBackend {}; + +struct HipFileBackendWithFallback : public HipFileUnopened, ::testing::WithParamInterface { + std::shared_ptr> mock_buffer; + std::shared_ptr> mock_file; + + std::shared_ptr> default_backend; + std::shared_ptr> fallback_backend; + + IoType io_type; + ssize_t successful_io_size = 0x1234; + + void SetUp() override + { + mock_buffer = std::make_shared>(); + mock_file = std::make_shared>(); + + default_backend = std::make_shared>(); + EXPECT_CALL(*default_backend, _io_impl).Times(AnyNumber()); + EXPECT_CALL(*default_backend, update_read_stats).Times(AnyNumber()); + EXPECT_CALL(*default_backend, update_write_stats).Times(AnyNumber()); + + fallback_backend = std::make_shared>(); + EXPECT_CALL(*fallback_backend, io).Times(AnyNumber()); + EXPECT_CALL(*fallback_backend, score).WillRepeatedly(Return(0)); + + io_type = GetParam(); + } + + HipFileBackendWithFallback() + { + } +}; + +TEST_P(HipFileBackendWithFallback, IOSuccess) +{ + EXPECT_CALL(*default_backend, _io_impl).WillOnce(Return(successful_io_size)); + + ssize_t nbytes = default_backend->io(io_type, mock_file, mock_buffer, 0, 0, 0); + + ASSERT_EQ(nbytes, successful_io_size); +} + +TEST_P(HipFileBackendWithFallback, IOFailureNoFallback) +{ + EXPECT_CALL(*default_backend, _io_impl).WillOnce(Throw(std::runtime_error("IO failure"))); + + EXPECT_THROW(default_backend->io(io_type, mock_file, mock_buffer, 0, 0, 0), std::runtime_error); +} + +TEST_P(HipFileBackendWithFallback, IOFailureWithIneligibleRetry) +{ + // The Backend has registered a fallback, but has determined that the + // IO error should not be retried. + default_backend->register_fallback_backend(fallback_backend); + EXPECT_CALL(*default_backend, _io_impl).WillOnce(Throw(std::runtime_error("IO failure"))); + EXPECT_CALL(*fallback_backend, score).WillOnce(Return(-1)); + + EXPECT_THROW(default_backend->io(io_type, mock_file, mock_buffer, 0, 0, 0), std::runtime_error); +} + +TEST_P(HipFileBackendWithFallback, IOFailureWithGoodFallback) +{ + default_backend->register_fallback_backend(fallback_backend); + EXPECT_CALL(*default_backend, _io_impl).WillOnce(Throw(std::runtime_error("IO failure"))); + EXPECT_CALL(*fallback_backend, io).WillOnce(Return(successful_io_size)); + + ssize_t nbytes = default_backend->io(io_type, mock_file, mock_buffer, 0, 0, 0); + ASSERT_EQ(nbytes, successful_io_size); +} + +INSTANTIATE_TEST_SUITE_P(, HipFileBackendWithFallback, ::testing::Values(IoType::Read, IoType::Write)); + +HIPFILE_WARN_NO_GLOBAL_CTOR_ON diff --git a/test/amd_detail/buffer.cpp b/test/amd_detail/buffer.cpp index 91ac41f2..9ba45954 100644 --- a/test/amd_detail/buffer.cpp +++ b/test/amd_detail/buffer.cpp @@ -11,6 +11,7 @@ #include "hipfile-warnings.h" #include "mhip.h" #include "state.h" +#include "test-common.h" #include #include @@ -97,7 +98,7 @@ TEST_F(HipFileBuffer, register_hip_pointer_get_attributes_error) { StrictMock mhip; EXPECT_CALL(mhip, hipPointerGetAttributes).WillOnce(testing::Throw(Hip::RuntimeError(hipErrorUnknown))); - ASSERT_EQ(hipFileBufRegister(nonnull_ptr, 0, 0), HipFileHipError(hipErrorUnknown)); + ASSERT_EQ(hipFileBufRegister(nonnull_ptr, 0, 0), HipFileDriverError(hipErrorUnknown)); // hipErrorInvalidValue is handled differently to match the behaviour of cufile EXPECT_CALL(mhip, hipPointerGetAttributes) diff --git a/test/amd_detail/driver.cpp b/test/amd_detail/driver.cpp index 02e2377b..a271d2a2 100644 --- a/test/amd_detail/driver.cpp +++ b/test/amd_detail/driver.cpp @@ -9,6 +9,7 @@ #include "mhip.h" #include "mmountinfo.h" #include "msys.h" +#include "test-common.h" #include #include diff --git a/test/amd_detail/fallback.cpp b/test/amd_detail/fallback.cpp index dffa530c..ebebadae 100644 --- a/test/amd_detail/fallback.cpp +++ b/test/amd_detail/fallback.cpp @@ -19,6 +19,7 @@ #include "mmountinfo.h" #include "msys.h" #include "state.h" +#include "test-common.h" #include #include diff --git a/test/amd_detail/fastpath.cpp b/test/amd_detail/fastpath.cpp index 258f29d0..f0d4c061 100644 --- a/test/amd_detail/fastpath.cpp +++ b/test/amd_detail/fastpath.cpp @@ -3,6 +3,7 @@ * SPDX-License-Identifier: MIT */ +#include "backend/fallback.h" #include "backend/fastpath.h" #include "hip.h" #include "hipfile.h" @@ -12,10 +13,12 @@ #include "mbuffer.h" #include "mfile.h" #include "mhip.h" +#include "msys.h" #include #include #include +#include #include #include #include @@ -25,6 +28,7 @@ #include #include #include +#include using namespace hipFile; using namespace testing; @@ -555,4 +559,128 @@ TEST_P(FastpathIoParam, IoSizeIsTruncatedToMaxRWCount) INSTANTIATE_TEST_SUITE_P(FastpathTest, FastpathIoParam, Values(IoType::Read, IoType::Write)); +struct FastpathIoParamWithFallback : public FastpathTestBase, + public TestWithParam> { + inline IoType _get_param_io_type() const + { + return std::get<0>(GetParam()); + } + + inline const std::exception_ptr _get_param_exc_ptr() const + { + return std::get<1>(GetParam()); + } +}; + +// The Fastpath can throw a few different kinds of derived std::runtime_errors. +TEST_P(FastpathIoParamWithFallback, IntegrationRunWithFallback) +{ + StrictMock mhip; + StrictMock msys; + + auto fallback_backend = std::make_shared>(); + auto fastpath_backend = std::make_shared>(); + fastpath_backend->register_fallback_backend(fallback_backend); + + const int DEFAULT_BUFFERED_FD = DEFAULT_UNBUFFERED_FD.value() + 1; + + // Called by both Fastpath and Fallback + EXPECT_CALL(*mbuffer, getBuffer).WillRepeatedly(Return(DEFAULT_BUFFER_ADDR)); + EXPECT_CALL(*mbuffer, getLength).Times(2).WillRepeatedly(Return(DEFAULT_BUFFER_LENGTH)); + // Called only by Fastpath + EXPECT_CALL(*mfile, getUnbufferedFd).WillOnce(Return(DEFAULT_UNBUFFERED_FD)); + // Called only by Fallback + EXPECT_CALL(*mbuffer, getType).WillOnce(Return(hipMemoryTypeDevice)); + EXPECT_CALL(*mfile, getBufferedFd).WillRepeatedly(Return(DEFAULT_BUFFERED_FD)); + EXPECT_CALL(mhip, hipMemcpy).WillRepeatedly(Return()); + EXPECT_CALL(msys, mmap).WillOnce(Return(reinterpret_cast(0x12345678))); + EXPECT_CALL(msys, munmap).WillOnce(Return()); + switch (_get_param_io_type()) { + case IoType::Read: + // Called by Fastpath + EXPECT_CALL(mhip, hipAmdFileRead).WillOnce(Rethrow(_get_param_exc_ptr())); + // Called by Fallback + EXPECT_CALL(msys, pread).WillRepeatedly(ReturnArg<2>()); + break; + case IoType::Write: + // Called by Fastpath + EXPECT_CALL(mhip, hipAmdFileWrite).WillOnce(Rethrow(_get_param_exc_ptr())); + // Called by Fallback + EXPECT_CALL(mhip, hipStreamSynchronize).WillRepeatedly(Return()); + EXPECT_CALL(msys, fdatasync).WillRepeatedly(Return()); + EXPECT_CALL(msys, pwrite).WillRepeatedly(ReturnArg<2>()); + break; + default: + FAIL() << "Invalid IoType"; + } + + ssize_t num_bytes = fastpath_backend->io(_get_param_io_type(), mfile, mbuffer, DEFAULT_IO_SIZE, 0, 0); + ASSERT_EQ(num_bytes, DEFAULT_IO_SIZE); +} + +// If the fallback backend rejects the IO, the original exception from +// Fastpath should be raised. +TEST_P(FastpathIoParamWithFallback, IntegrationFallbackRejectsIO) +{ + StrictMock mhip; + StrictMock msys; + + auto fallback_backend = std::make_shared>(); + auto fastpath_backend = std::make_shared>(); + fastpath_backend->register_fallback_backend(fallback_backend); + + // Called only by Fastpath + EXPECT_CALL(*mbuffer, getBuffer).WillOnce(Return(DEFAULT_BUFFER_ADDR)); + EXPECT_CALL(*mbuffer, getLength).WillOnce(Return(DEFAULT_BUFFER_LENGTH)); + EXPECT_CALL(*mfile, getUnbufferedFd).WillOnce(Return(DEFAULT_UNBUFFERED_FD)); + // Called only by Fallback - should fail the score() check. + EXPECT_CALL(*mbuffer, getType).WillOnce(Return(hipMemoryTypeHost)); + switch (_get_param_io_type()) { + case IoType::Read: + // Called by Fastpath + EXPECT_CALL(mhip, hipAmdFileRead).WillOnce(Rethrow(_get_param_exc_ptr())); + break; + case IoType::Write: + // Called by Fastpath + EXPECT_CALL(mhip, hipAmdFileWrite).WillOnce(Rethrow(_get_param_exc_ptr())); + break; + default: + FAIL() << "Invalid IoType"; + } + + // Have to rethrow the exception_ptr to be able to access the exception + // This looks ugly, but is better than the alternative of trying to preserve the + // exception type when setting Throw(*std::shared_ptr>). + try { + std::rethrow_exception(_get_param_exc_ptr()); + } + catch (const std::exception &expected_exc) { + // Can't use EXPECT_THROW due to the thrown exception type only being known at runtime + try { + fastpath_backend->io(_get_param_io_type(), mfile, mbuffer, DEFAULT_IO_SIZE, 0, 0); + FAIL() << "io() was expected to throw, but it returned normally"; + } + catch (const std::exception &actual_exc) { + // Verify that the propagated exception has the same dynamic type and message + // as the one stored in the original std::exception_ptr, without relying on + // pointer identity of the underlying exception object. + ASSERT_EQ(typeid(expected_exc), typeid(actual_exc)); + ASSERT_STREQ(expected_exc.what(), actual_exc.what()); + } + catch (...) { + FAIL() << "io() threw something other than a std::exception"; + } + } +} + +// Using std::exception_ptr is more straightforward here than storing a pointer +// to a derived std::exception type, which would require careful handling when +// setting expectations. Note that Throw() does not accept std::exception_ptr, +// but the public (though undocumented) Rethrow() action does support it. +INSTANTIATE_TEST_SUITE_P( + FastpathTest, FastpathIoParamWithFallback, + Combine(Values(IoType::Read, IoType::Write), + Values(std::make_exception_ptr(Hip::RuntimeError(hipErrorNoDevice)), + std::make_exception_ptr(std::system_error(make_error_code(errc::no_such_device)))))); + HIPFILE_WARN_NO_GLOBAL_CTOR_ON diff --git a/test/amd_detail/handle.cpp b/test/amd_detail/handle.cpp index a067a5c1..cb1bb632 100644 --- a/test/amd_detail/handle.cpp +++ b/test/amd_detail/handle.cpp @@ -12,6 +12,7 @@ #include "mmountinfo.h" #include "mountinfo.h" #include "state.h" +#include "test-common.h" #include #include diff --git a/test/amd_detail/hipfile-api.cpp b/test/amd_detail/hipfile-api.cpp index 73367ffe..43e30a2f 100644 --- a/test/amd_detail/hipfile-api.cpp +++ b/test/amd_detail/hipfile-api.cpp @@ -29,6 +29,7 @@ #include "mstate.h" #include "msys.h" #include "state.h" +#include "test-common.h" #include #include diff --git a/test/amd_detail/hipfile-test.h b/test/amd_detail/hipfile-test.h index b6a72cf8..714c1691 100644 --- a/test/amd_detail/hipfile-test.h +++ b/test/amd_detail/hipfile-test.h @@ -12,6 +12,7 @@ #include "mhip.h" #include "mmountinfo.h" #include "msys.h" +#include "test-common.h" #include #include @@ -19,56 +20,6 @@ #include #include -// *********************************************************************** -// ERRORS AND ERROR HANDLING -// *********************************************************************** - -// Set a particular hipFile error -constexpr hipFileError_t -HipFileHipError(hipError_t err) -{ - return {hipFileHipDriverError, err}; -} - -// Set a particular HIP error -constexpr hipFileError_t -HipFileOpError(hipFileOpError_t err) -{ - return {err, hipSuccess}; -} - -// == overload for hipFileError_t values -inline bool -operator==(const hipFileError_t &lhs, const hipFileError_t &rhs) -{ - return lhs.err == rhs.err && lhs.hip_drv_err == rhs.hip_drv_err; -} - -// != overload for hipFileError_t values -inline bool -operator!=(const hipFileError_t &lhs, const hipFileError_t &rhs) -{ - return lhs.err != rhs.err || lhs.hip_drv_err != rhs.hip_drv_err; -} - -// << overload for hipFileError_t values -// -// Unused in the test code, but kept here for iostream debugging -#ifndef NDEBUG -#include -inline std::ostream & -operator<<(std::ostream &os, const hipFileError_t &rfe) -{ - return os << "hipFileError_t{ err: " << rfe.err << ", hip_drv_err: " << rfe.hip_drv_err << " }"; -} -#endif - -// Convenience "success" value -inline constexpr hipFileError_t HIPFILE_SUCCESS{hipFileSuccess, hipSuccess}; - -// Convenience "invalid argument" value -inline constexpr hipFileError_t HIPFILE_INVALID_VALUE{hipFileInvalidValue, hipSuccess}; - // *********************************************************************** // BASE ERROR CLASSES // *********************************************************************** diff --git a/test/amd_detail/mbackend.h b/test/amd_detail/mbackend.h index f4c395df..d544e0b2 100644 --- a/test/amd_detail/mbackend.h +++ b/test/amd_detail/mbackend.h @@ -18,6 +18,22 @@ struct MBackend : Backend { (hipFile::IoType type, std::shared_ptr, std::shared_ptr, size_t, hoff_t, hoff_t), (override)); + MOCK_METHOD(void, update_read_stats, (ssize_t nbytes), (override)); + MOCK_METHOD(void, update_write_stats, (ssize_t nbytes), (override)); + MOCK_METHOD(ssize_t, _io_impl, + (hipFile::IoType type, std::shared_ptr, std::shared_ptr, size_t, hoff_t, + hoff_t), + (override)); }; +struct MBackendWithFallback : BackendWithFallback { + MOCK_METHOD(int, score, (std::shared_ptr, std::shared_ptr, size_t, hoff_t, hoff_t), + (const, override)); + MOCK_METHOD(void, update_read_stats, (ssize_t nbytes), (override)); + MOCK_METHOD(void, update_write_stats, (ssize_t nbytes), (override)); + MOCK_METHOD(ssize_t, _io_impl, + (IoType type, std::shared_ptr file, std::shared_ptr buffer, size_t size, + hoff_t file_offset, hoff_t buffer_offset), + (override)); +}; } diff --git a/test/amd_detail/mhip.h b/test/amd_detail/mhip.h index f8d0d630..8b4fc6c2 100644 --- a/test/amd_detail/mhip.h +++ b/test/amd_detail/mhip.h @@ -7,9 +7,12 @@ #include "context.h" #include "hip.h" +#include "test-common.h" #include +#define MHIP_PASSTHROUGH(func) MOCK_PASSTHROUGH(hipFile::Hip, func) + /* mhipxx (mock hip++) * * Mock implementations for Hip. Enables unit tests to mock HIP APIs. @@ -48,6 +51,25 @@ struct MHip : Hip { (const, override)); MOCK_METHOD(int, hipDeviceGetAttribute, (hipDeviceAttribute_t attr, int device_id), (const, override)); MOCK_METHOD(hipDevice_t, hipStreamGetDevice, (hipStream_t stream), (const, override)); + + void enable_passthrough() + { + MHIP_PASSTHROUGH(hipPointerGetAttributes); + MHIP_PASSTHROUGH(hipMemcpy); + MHIP_PASSTHROUGH(hipStreamSynchronize); + MHIP_PASSTHROUGH(hipHostMalloc); + MHIP_PASSTHROUGH(hipHostFree); + MHIP_PASSTHROUGH(hipHostGetDevicePointer); + MHIP_PASSTHROUGH(hipRuntimeGetVersion); + MHIP_PASSTHROUGH(hipGetProcAddress); + MHIP_PASSTHROUGH(hipAmdFileRead); + MHIP_PASSTHROUGH(hipAmdFileWrite); + MHIP_PASSTHROUGH(hipMemGetAddressRange); + MHIP_PASSTHROUGH(hipLaunchHostFunc); + MHIP_PASSTHROUGH(hipLaunchKernel); + MHIP_PASSTHROUGH(hipDeviceGetAttribute); + MHIP_PASSTHROUGH(hipStreamGetDevice); + } }; } diff --git a/test/amd_detail/stream.cpp b/test/amd_detail/stream.cpp index d137f3e1..3a6906f8 100644 --- a/test/amd_detail/stream.cpp +++ b/test/amd_detail/stream.cpp @@ -9,6 +9,7 @@ #include "mhip.h" #include "msys.h" #include "stream.h" +#include "test-common.h" #include #include diff --git a/test/common/test-common.h b/test/common/test-common.h index 95debab6..d225ea51 100644 --- a/test/common/test-common.h +++ b/test/common/test-common.h @@ -14,33 +14,58 @@ #include #include +// ON_CALL is not compatible with StrictMock's +#define MOCK_PASSTHROUGH(base_class, func) \ + ON_CALL(*this, func).WillByDefault( \ + [this](auto&&... args) { \ + return this->base_class::func(std::forward(args)...); \ + } \ + ) + +// Set a particular hipfile error constexpr hipFileError_t HipFileOpError(hipFileOpError_t err) { return {err, hipSuccess}; } +// Set a particular HIP error constexpr hipFileError_t HipFileDriverError(hipError_t err) { return {hipFileHipDriverError, err}; } +// == overload for hipFileError_t values inline bool operator==(const hipFileError_t &lhs, const hipFileError_t &rhs) { return lhs.err == rhs.err && lhs.hip_drv_err == rhs.hip_drv_err; } -static std::ostream & +// != overload for hipFileError_t values +inline bool +operator!=(const hipFileError_t &lhs, const hipFileError_t &rhs) +{ + return lhs.err != rhs.err || lhs.hip_drv_err != rhs.hip_drv_err; +} + +// Unused in the test code, but kept here for iostream debugging +#ifndef NDEBUG +#include +inline std::ostream & operator<<(std::ostream &os, const hipFileError_t &hfe) { return os << "hipFileError_t{ err: " << hfe.err << ", hip_drv_err: " << hfe.hip_drv_err << " }"; } +#endif // Convenience "success" value inline constexpr hipFileError_t HIPFILE_SUCCESS{hipFileSuccess, hipSuccess}; +// Convenience "invalid argument" value +inline constexpr hipFileError_t HIPFILE_INVALID_VALUE{hipFileInvalidValue, hipSuccess}; + inline void rfill(void *buffer, uint64_t len, uint32_t seed = 97) { diff --git a/test/system/fastpath.cpp b/test/system/fastpath.cpp new file mode 100644 index 00000000..e826a82d --- /dev/null +++ b/test/system/fastpath.cpp @@ -0,0 +1,117 @@ +/* Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + */ + +#include "hip.h" +#include "hipfile.h" +#include "io.h" +#include "mhip.h" + +#include "test-common.h" +#include "test-options.h" + +#include +#include +#include +#include +#include + +//using namespace hipFile; +using namespace testing; +using namespace std; + +extern SystemTestOptions test_env; + +HIPFILE_WARN_NO_GLOBAL_CTOR_OFF + +struct FastpathWithFallbackIO : public TestWithParam> { + + Tmpfile tmpfile; + size_t tmpfile_size; + hipFileHandle_t tmpfile_handle; + void *device_buffer; + size_t device_buffer_size; + + NiceMock mhip; + + FastpathWithFallbackIO() + : tmpfile{test_env.ais_capable_dir}, tmpfile_size{1024 * 1024}, tmpfile_handle{nullptr}, + device_buffer{nullptr}, device_buffer_size{1024 * 1024} + { + } + + void SetUp() override + { + if (unsetenv("HIPFILE_FORCE_COMPAT_MODE")) { + FAIL() << "Could not clear HIPFILE_FORCE_COMPAT_MODE"; + } + if (setenv("HIPFILE_ALLOW_COMPAT_MODE", "true", 1)) { + FAIL() << "Could not set HIPFILE_ALLOW_COMPAT_MODE=true"; + } + std::cout << "HIPFILE_ALLOW_COMPAT_MODE=" << getenv("HIPFILE_ALLOW_COMPAT_MODE") << std::endl; + // Must be called prior to any expectations on Hip set. + mhip.enable_passthrough(); + + switch (_get_param_io_type()) { + case hipFile::IoType::Read: + EXPECT_CALL(mhip, hipAmdFileRead).WillOnce(Rethrow(_get_param_exc_ptr())); + break; + case hipFile::IoType::Write: + EXPECT_CALL(mhip, hipAmdFileWrite).WillOnce(Rethrow(_get_param_exc_ptr())); + break; + default: + FAIL() << "Unsupported IoTestBackend"; + } + + ASSERT_EQ(0, ftruncate(tmpfile.fd, static_cast(tmpfile_size))); + + hipFileDescr_t descr{}; + descr.type = hipFileHandleTypeOpaqueFD; + descr.handle.fd = tmpfile.fd; + + ASSERT_EQ(HIPFILE_SUCCESS, hipFileHandleRegister(&tmpfile_handle, &descr)); + ASSERT_EQ(hipSuccess, hipMalloc(&device_buffer, device_buffer_size)); + ASSERT_EQ(HIPFILE_SUCCESS, hipFileBufRegister(device_buffer, device_buffer_size, 0)); + } + + void TearDown() override + { + ASSERT_EQ(HIPFILE_SUCCESS, hipFileBufDeregister(device_buffer)); + ASSERT_EQ(hipSuccess, hipFree(device_buffer)); + hipFileHandleDeregister(tmpfile_handle); + } + + inline hipFile::IoType _get_param_io_type() const + { + return std::get<0>(GetParam()); + } + + inline const std::exception_ptr _get_param_exc_ptr() const + { + return std::get<1>(GetParam()); + } +}; + +TEST_P(FastpathWithFallbackIO, DummyTest) +{ + // Simply checks to see if the set & teardown portions have been configured correctly. + switch (_get_param_io_type()) { + case hipFile::IoType::Read: + ASSERT_EQ(tmpfile_size, hipFileRead(tmpfile_handle, device_buffer, tmpfile_size, 0, 0)); + break; + case hipFile::IoType::Write: + ASSERT_EQ(tmpfile_size, hipFileWrite(tmpfile_handle, device_buffer, tmpfile_size, 0, 0)); + break; + default: + FAIL() << "Unsupported IoTestBackend"; + } +} + +INSTANTIATE_TEST_SUITE_P( + , FastpathWithFallbackIO, + Combine(Values(hipFile::IoType::Read, hipFile::IoType::Write), + Values(std::make_exception_ptr(hipFile::Hip::RuntimeError(hipErrorNoDevice)), + std::make_exception_ptr(std::system_error(make_error_code(errc::no_such_device)))))); + +HIPFILE_WARN_NO_GLOBAL_CTOR_ON