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"