From 671778bdd3b3ac787d5ea2af0e6dbb70d55acc9b Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Mon, 8 Feb 2021 15:32:25 -0500 Subject: [PATCH] SWDEV-232428 - Push hostcall implementation to the device layer This change unifies the hostcall implementation for all the backends, by pushing the common logic to the device layer. This is done by replacing the use of hsa_signal_t with device::Signal (a light wrapper around it). Change-Id: I7b6fca7930b5a0b199da5d85e2e048354cc04e7b --- rocclr/CMakeLists.txt | 1 + .../{rocm/rochostcall.cpp => devhostcall.cpp} | 38 +++++------ .../{rocm/rochostcall.hpp => devhostcall.hpp} | 7 +- rocclr/device/device.hpp | 4 ++ rocclr/device/devsignal.hpp | 65 +++++++++++++++++++ rocclr/device/gpu/gpudevice.hpp | 10 +++ rocclr/device/pal/paldevice.hpp | 20 ++++-- rocclr/device/rocm/CMakeLists.txt | 2 +- rocclr/device/rocm/rocdevice.cpp | 9 ++- rocclr/device/rocm/rocdevice.hpp | 7 ++ rocclr/device/rocm/rocsignal.cpp | 53 +++++++++++++++ rocclr/device/rocm/rocsignal.hpp | 47 ++++++++++++++ rocclr/device/rocm/rocvirtual.cpp | 2 +- 13 files changed, 234 insertions(+), 31 deletions(-) rename rocclr/device/{rocm/rochostcall.cpp => devhostcall.cpp} (94%) rename rocclr/device/{rocm/rochostcall.hpp => devhostcall.hpp} (95%) create mode 100644 rocclr/device/devsignal.hpp create mode 100644 rocclr/device/rocm/rocsignal.cpp create mode 100644 rocclr/device/rocm/rocsignal.hpp diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt index 2c3466e8b8..0548e4d754 100644 --- a/rocclr/CMakeLists.txt +++ b/rocclr/CMakeLists.txt @@ -153,6 +153,7 @@ set(oclruntime_src device/devprogram.cpp device/devhcprintf.cpp device/devhcmessages.cpp + device/devhostcall.cpp platform/activity.cpp platform/kernel.cpp platform/context.cpp diff --git a/rocclr/device/rocm/rochostcall.cpp b/rocclr/device/devhostcall.cpp similarity index 94% rename from rocclr/device/rocm/rochostcall.cpp rename to rocclr/device/devhostcall.cpp index 9755b80395..7f74cc8eb9 100644 --- a/rocclr/device/rocm/rochostcall.cpp +++ b/rocclr/device/devhostcall.cpp @@ -22,8 +22,9 @@ #include "top.hpp" #include "utils/flags.hpp" -#include "rochostcall.hpp" #include "device/devhcmessages.hpp" +#include "device/devhostcall.hpp" +#include "device/devsignal.hpp" #include "os/os.hpp" #include "thread/monitor.hpp" @@ -31,10 +32,6 @@ #include "utils/debug.hpp" #include "utils/flags.hpp" -#include "rochostcall.hpp" - -#include - #include #include #include @@ -100,7 +97,7 @@ class HostcallBuffer { /** Array of packet payloads */ Payload* payloads_; /** Signal used by kernels to indicate new work */ - hsa_signal_t doorbell_; + void* doorbell_; /** Stack of free packets. Uses tagged pointers. */ uint64_t free_stack_; /** Stack of ready packets. Uses tagged pointers */ @@ -114,7 +111,7 @@ class HostcallBuffer { public: void processPackets(MessageHandler& messages); void initialize(uint32_t num_packets); - void setDoorbell(hsa_signal_t doorbell) { doorbell_ = doorbell; }; + void setDoorbell(void* doorbell) { doorbell_ = doorbell; }; }; static_assert(std::is_standard_layout::value, @@ -262,7 +259,7 @@ void HostcallBuffer::initialize(uint32_t num_packets) { */ class HostcallListener { std::set buffers_; - hsa_signal_t doorbell_; + device::Signal* doorbell_; MessageHandler messages_; class Thread : public amd::Thread { @@ -303,20 +300,19 @@ class HostcallListener { } void terminate(); - bool initialize(); + bool initialize(amd::Device &dev); }; HostcallListener* hostcallListener = nullptr; amd::Monitor listenerLock("Hostcall listener lock"); void HostcallListener::consumePackets() { - uint64_t signal_value = SIGNAL_INIT; uint64_t timeout = 1024 * 1024; + uint64_t signal_value = SIGNAL_INIT; while (true) { while (true) { - uint64_t new_value = hsa_signal_wait_scacquire(doorbell_, HSA_SIGNAL_CONDITION_NE, signal_value, timeout, - HSA_WAIT_STATE_BLOCKED); + uint64_t new_value = doorbell_->Wait(signal_value, device::Signal::Condition::Ne, timeout); if (new_value != signal_value) { signal_value = new_value; break; @@ -343,19 +339,19 @@ void HostcallListener::terminate() { return; } - hsa_signal_store_screlease(doorbell_, SIGNAL_DONE); + doorbell_->Reset(SIGNAL_DONE); // FIXME_lmoriche: fix termination handshake while (thread_.state() < Thread::FINISHED) { amd::Os::yield(); } - hsa_signal_destroy(doorbell_); + delete doorbell_; } void HostcallListener::addBuffer(HostcallBuffer* buffer) { assert(buffers_.count(buffer) == 0 && "buffer already present"); - buffer->setDoorbell(doorbell_); + buffer->setDoorbell(doorbell_->getHandle()); buffers_.insert(buffer); } @@ -364,16 +360,16 @@ void HostcallListener::removeBuffer(HostcallBuffer* buffer) { buffers_.erase(buffer); } -bool HostcallListener::initialize() { - auto status = hsa_signal_create(SIGNAL_INIT, 0, NULL, &doorbell_); - if (status != HSA_STATUS_SUCCESS) { +bool HostcallListener::initialize(amd::Device &dev) { + doorbell_ = dev.createSignal(); + if ((doorbell_ == nullptr) || !doorbell_->Init(SIGNAL_INIT, device::Signal::WaitState::Blocked)) { return false; } // If the listener thread was not successfully initialized, clean // everything up and bail out. if (thread_.state() < Thread::INITIALIZED) { - hsa_signal_destroy(doorbell_); + delete doorbell_; return false; } @@ -381,14 +377,14 @@ bool HostcallListener::initialize() { return true; } -bool enableHostcalls(void* bfr, uint32_t numPackets) { +bool enableHostcalls(amd::Device &dev, void* bfr, uint32_t numPackets) { auto buffer = reinterpret_cast(bfr); buffer->initialize(numPackets); amd::ScopedLock lock(listenerLock); if (!hostcallListener) { hostcallListener = new HostcallListener(); - if (!hostcallListener->initialize()) { + if (!hostcallListener->initialize(dev)) { ClPrint(amd::LOG_ERROR, (amd::LOG_INIT | amd::LOG_QUEUE | amd::LOG_RESOURCE), "Failed to launch hostcall listener"); delete hostcallListener; diff --git a/rocclr/device/rocm/rochostcall.hpp b/rocclr/device/devhostcall.hpp similarity index 95% rename from rocclr/device/rocm/rochostcall.hpp rename to rocclr/device/devhostcall.hpp index 09a4b6b343..bfbc107782 100644 --- a/rocclr/device/rocm/rochostcall.hpp +++ b/rocclr/device/devhostcall.hpp @@ -20,6 +20,11 @@ #pragma once +#include "top.hpp" +#include "device/device.hpp" + +#include + /** \file Support for invoking host services from the device. * * A hostcall is a fixed-size request generated by a kernel running @@ -72,5 +77,5 @@ size_t getHostcallBufferSize(uint32_t num_packets); */ uint32_t getHostcallBufferAlignment(void); -bool enableHostcalls(void* buffer, uint32_t numPackets); +bool enableHostcalls(amd::Device& dev, void* buffer, uint32_t numPackets); void disableHostcalls(void* buffer); diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index 4ac05cb2b0..0084ac18f2 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -36,6 +36,7 @@ #include "amdocl/cl_profile_amd.h" #include "acl.h" #include "hwdebug.hpp" +#include "devsignal.hpp" #include #include @@ -1599,6 +1600,9 @@ class Device : public RuntimeObject { const device::Memory& parent //!< Parent device memory object for the view ) const = 0; + ///! Allocates a device signal object + virtual device::Signal* createSignal() const = 0; + //! Return true if initialized external API interop, otherwise false virtual bool bindExternalDevice( uint flags, //!< Enum val. for ext.API type: GL, D3D10, etc. diff --git a/rocclr/device/devsignal.hpp b/rocclr/device/devsignal.hpp new file mode 100644 index 0000000000..577881a489 --- /dev/null +++ b/rocclr/device/devsignal.hpp @@ -0,0 +1,65 @@ +/* Copyright (c) 2021-present Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#pragma once + +#include "top.hpp" + +namespace device { + +// Light abstraction over HSA/PAL signals +class Signal: public amd::HeapObject { +public: + enum class Condition : uint32_t { + Eq = 0, + Ne = 1, + Lt = 2, + Gte = 3, + }; + + enum class WaitState : uint32_t { + Blocked = 0, + Active = 1, + }; + +protected: + WaitState ws_; + +public: + virtual ~Signal() {} + + virtual bool Init(uint64_t init, WaitState ws) + { return false; } + + // Blocks the current thread untill the condition c is satisfied + // or amount of time specified by timeout passes + virtual uint64_t Wait(uint64_t value, Condition c, uint64_t timeout) + { return -1; } + + // Atomically sets the current value of the signal + virtual void Reset(uint64_t value) + {} + + // Return the handle to the underlying amd_signal_t object + virtual void* getHandle() + { return nullptr; } +}; + +}; \ No newline at end of file diff --git a/rocclr/device/gpu/gpudevice.hpp b/rocclr/device/gpu/gpudevice.hpp index 1f31368f3d..30693c716f 100644 --- a/rocclr/device/gpu/gpudevice.hpp +++ b/rocclr/device/gpu/gpudevice.hpp @@ -102,6 +102,11 @@ class NullDevice : public amd::Device { return NULL; } + //! Signal object allocation + virtual device::Signal* createSignal() const { + return nullptr; + } + //! Acquire external graphics API object in the host thread //! Needed for OpenGL objects on CPU device @@ -415,6 +420,11 @@ class Device : public NullDevice, public CALGSLDevice { const device::Memory& parent //!< Parent device memory object for the view ) const; + //! Signal object allocation + virtual device::Signal* createSignal() const { + return nullptr; + } + //! Create the device program. virtual device::Program* createProgram(amd::Program& owner, amd::option::Options* options = NULL); diff --git a/rocclr/device/pal/paldevice.hpp b/rocclr/device/pal/paldevice.hpp index 281c6ea5c7..98c31dff57 100644 --- a/rocclr/device/pal/paldevice.hpp +++ b/rocclr/device/pal/paldevice.hpp @@ -71,15 +71,15 @@ class NullDevice : public amd::Device { ); //! Instantiate a new virtual device - virtual device::VirtualDevice* createVirtualDevice(amd::CommandQueue* queue = NULL) { - return NULL; + virtual device::VirtualDevice* createVirtualDevice(amd::CommandQueue* queue = nullptr) { + return nullptr; } //! Compile the given source code. - virtual device::Program* createProgram(amd::Program& owner, amd::option::Options* options = NULL); + virtual device::Program* createProgram(amd::Program& owner, amd::option::Options* options = nullptr); //! Just returns NULL for the dummy device - virtual device::Memory* createMemory(amd::Memory& owner) const { return NULL; } + virtual device::Memory* createMemory(amd::Memory& owner) const { return nullptr; } //! Sampler object allocation virtual bool createSampler(const amd::Sampler& owner, //!< abstraction layer sampler object @@ -94,7 +94,12 @@ class NullDevice : public amd::Device { amd::Memory& owner, //!< Owner memory object const device::Memory& parent //!< Parent device memory object for the view ) const { - return NULL; + return nullptr; + } + + //! Signal object allocation + virtual device::Signal* createSignal() const { + return nullptr; } //! Acquire external graphics API object in the host thread @@ -354,6 +359,11 @@ class Device : public NullDevice { const device::Memory& parent //!< Parent device memory object for the view ) const; + //! Signal object allocation + virtual device::Signal* createSignal() const { + return nullptr; + } + //! Create the device program. virtual device::Program* createProgram(amd::Program& owner, amd::option::Options* options = NULL); diff --git a/rocclr/device/rocm/CMakeLists.txt b/rocclr/device/rocm/CMakeLists.txt index cb540f1e50..dbba2279be 100644 --- a/rocclr/device/rocm/CMakeLists.txt +++ b/rocclr/device/rocm/CMakeLists.txt @@ -11,7 +11,7 @@ add_library(oclrocm OBJECT rocappprofile.cpp rocsettings.cpp rocschedcl.cpp - rochostcall.cpp + rocsignal.cpp ) target_include_directories(oclrocm diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index d726bf6304..4237f5c328 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -32,17 +32,18 @@ #include "vdi_common.hpp" #include "device/comgrctx.hpp" +#include "device/devhostcall.hpp" #include "device/rocm/rocdevice.hpp" #include "device/rocm/rocblit.hpp" #include "device/rocm/rocvirtual.hpp" #include "device/rocm/rocprogram.hpp" #include "device/rocm/rocmemory.hpp" #include "device/rocm/rocglinterop.hpp" +#include "device/rocm/rocsignal.hpp" #ifdef WITH_AMDGPU_PRO #include "pro/prodriver.hpp" #endif #include "platform/sampler.hpp" -#include "rochostcall.hpp" #include #include @@ -2677,7 +2678,7 @@ void* Device::getOrCreateHostcallBuffer(hsa_queue_t* queue, bool coop_queue, } else { coopHostcallBuffer_ = buffer; } - if (!enableHostcalls(buffer, numPackets)) { + if (!enableHostcalls(*this, buffer, numPackets)) { ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, "Failed to register hostcall buffer %p with listener", buffer); return nullptr; @@ -2855,5 +2856,9 @@ void Device::getGlobalCUMask(std::string cuMaskStr) { } } +device::Signal* Device::createSignal() const { + return new roc::Signal(); +} + } // namespace roc #endif // WITHOUT_HSA_BACKEND diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index f1d50a73eb..fd1e40d10d 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/rocclr/device/rocm/rocdevice.hpp @@ -165,6 +165,11 @@ class NullDevice : public amd::Device { return nullptr; } + virtual device::Signal* createSignal() const { + ShouldNotReachHere(); + return nullptr; + } + //! Just returns nullptr for the dummy device virtual void* svmAlloc(amd::Context& context, //!< The context used to create a buffer size_t size, //!< size of svm spaces @@ -340,6 +345,8 @@ class Device : public NullDevice { return nullptr; } + virtual device::Signal* createSignal() const; + //! Acquire external graphics API object in the host thread //! Needed for OpenGL objects on CPU device virtual bool bindExternalDevice(uint flags, void* const pDevice[], void* pContext, diff --git a/rocclr/device/rocm/rocsignal.cpp b/rocclr/device/rocm/rocsignal.cpp new file mode 100644 index 0000000000..c252075e43 --- /dev/null +++ b/rocclr/device/rocm/rocsignal.cpp @@ -0,0 +1,53 @@ +/* Copyright (c) 2021-present Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "rocsignal.hpp" + +namespace roc { + +Signal::~Signal() { + hsa_signal_destroy(signal_); +} + +bool Signal::Init(uint64_t init, device::Signal::WaitState ws) { + hsa_status_t status = hsa_signal_create(init, 0, nullptr, &signal_); + if (status != HSA_STATUS_SUCCESS) { + return false; + } + + ws_ = ws; + + return true; +} + +uint64_t Signal::Wait(uint64_t value, device::Signal::Condition c, uint64_t timeout) { + return hsa_signal_wait_scacquire( + signal_, + static_cast(c), + value, + timeout, + static_cast(ws_)); +} + +void Signal::Reset(uint64_t value) { + hsa_signal_store_screlease(signal_, value); +} + +}; \ No newline at end of file diff --git a/rocclr/device/rocm/rocsignal.hpp b/rocclr/device/rocm/rocsignal.hpp new file mode 100644 index 0000000000..1c54fce40b --- /dev/null +++ b/rocclr/device/rocm/rocsignal.hpp @@ -0,0 +1,47 @@ +/* Copyright (c) 2021-present Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#pragma once + +#include "device/devsignal.hpp" + +#include + +namespace roc { + +class Signal: public device::Signal { +private: + hsa_signal_t signal_; + +public: + ~Signal() override; + + bool Init(uint64_t init, device::Signal::WaitState ws) override; + + uint64_t Wait(uint64_t value, device::Signal::Condition c, uint64_t timeout) override; + + void Reset(uint64_t value) override; + + void* getHandle() override { + return reinterpret_cast(signal_.handle); + } +}; + +}; \ No newline at end of file diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 373a597113..414dd87191 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -18,6 +18,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include "device/devhostcall.hpp" #include "device/rocm/rocdevice.hpp" #include "device/rocm/rocvirtual.hpp" #include "device/rocm/rockernel.hpp" @@ -30,7 +31,6 @@ #include "platform/command_utils.hpp" #include "platform/memory.hpp" #include "platform/sampler.hpp" -#include "rochostcall.hpp" #include "utils/debug.hpp" #include "os/os.hpp" #include "amd_hsa_kernel_code.h"