From 9eb113d9913c0ed02a604ff28d3e1ca898f08131 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Mon, 8 Mar 2021 13:48:44 -0500 Subject: [PATCH] SWDEV-232428 - PAL Hostcall support Since the majority of the Hostcall implementation now sits in the commmon layer, the PAL backend simply just needs to invoke it. One thing that is missing though is HSA signal support. The newly added pal::Signal class is a light emulaion of what HSA signals provide. The current implementation is just enough to get Hostcall working, but it can be expanded in the future if needed to fully emulate HSA signals. The major difference for now between PAL and ROCm hostcall implemenations is that PAL doesn't support blocking signals. This will be enabled in the near future. For now use active wait for PAL. Change-Id: I746557354ab9d71a7d4a31f9320fcc2fee5aee7f [ROCm/clr commit: 99e8ac55cd87ee60a3218f6617f96183240e10ca] --- projects/clr/rocclr/device/devhostcall.cpp | 13 ++- projects/clr/rocclr/device/devhostcall.hpp | 2 +- projects/clr/rocclr/device/devsignal.hpp | 6 +- projects/clr/rocclr/device/pal/CMakeLists.txt | 1 + projects/clr/rocclr/device/pal/paldevice.hpp | 3 +- projects/clr/rocclr/device/pal/palkernel.cpp | 12 +++ projects/clr/rocclr/device/pal/palsignal.cpp | 97 +++++++++++++++++++ projects/clr/rocclr/device/pal/palsignal.hpp | 50 ++++++++++ projects/clr/rocclr/device/pal/palvirtual.cpp | 49 ++++++++++ projects/clr/rocclr/device/pal/palvirtual.hpp | 4 + projects/clr/rocclr/device/rocm/rocsignal.cpp | 2 +- projects/clr/rocclr/device/rocm/rocsignal.hpp | 2 +- 12 files changed, 232 insertions(+), 9 deletions(-) create mode 100644 projects/clr/rocclr/device/pal/palsignal.cpp create mode 100644 projects/clr/rocclr/device/pal/palsignal.hpp diff --git a/projects/clr/rocclr/device/devhostcall.cpp b/projects/clr/rocclr/device/devhostcall.cpp index 8577abc01e..53cdfaa73b 100644 --- a/projects/clr/rocclr/device/devhostcall.cpp +++ b/projects/clr/rocclr/device/devhostcall.cpp @@ -300,7 +300,7 @@ class HostcallListener { } void terminate(); - bool initialize(amd::Device &dev); + bool initialize(const amd::Device &dev); }; HostcallListener* hostcallListener = nullptr; @@ -360,9 +360,14 @@ void HostcallListener::removeBuffer(HostcallBuffer* buffer) { buffers_.erase(buffer); } -bool HostcallListener::initialize(amd::Device &dev) { +bool HostcallListener::initialize(const amd::Device &dev) { doorbell_ = dev.createSignal(); - if ((doorbell_ == nullptr) || !doorbell_->Init(SIGNAL_INIT, device::Signal::WaitState::Blocked)) { +#ifdef WITH_PAL_DEVICE + auto ws = device::Signal::WaitState::Active; +#elif WITH_HSA_DEVICE + auto ws = device::Signal::WaitState::Blocked; +#endif + if ((doorbell_ == nullptr) || !doorbell_->Init(dev, SIGNAL_INIT, ws)) { return false; } @@ -377,7 +382,7 @@ bool HostcallListener::initialize(amd::Device &dev) { return true; } -bool enableHostcalls(amd::Device &dev, void* bfr, uint32_t numPackets) { +bool enableHostcalls(const amd::Device &dev, void* bfr, uint32_t numPackets) { auto buffer = reinterpret_cast(bfr); buffer->initialize(numPackets); diff --git a/projects/clr/rocclr/device/devhostcall.hpp b/projects/clr/rocclr/device/devhostcall.hpp index bfbc107782..1a3f275a9d 100644 --- a/projects/clr/rocclr/device/devhostcall.hpp +++ b/projects/clr/rocclr/device/devhostcall.hpp @@ -77,5 +77,5 @@ size_t getHostcallBufferSize(uint32_t num_packets); */ uint32_t getHostcallBufferAlignment(void); -bool enableHostcalls(amd::Device& dev, void* buffer, uint32_t numPackets); +bool enableHostcalls(const amd::Device& dev, void* buffer, uint32_t numPackets); void disableHostcalls(void* buffer); diff --git a/projects/clr/rocclr/device/devsignal.hpp b/projects/clr/rocclr/device/devsignal.hpp index 577881a489..3ce7b0766d 100644 --- a/projects/clr/rocclr/device/devsignal.hpp +++ b/projects/clr/rocclr/device/devsignal.hpp @@ -22,6 +22,10 @@ #include "top.hpp" +namespace amd { + class Device; +}; + namespace device { // Light abstraction over HSA/PAL signals @@ -45,7 +49,7 @@ protected: public: virtual ~Signal() {} - virtual bool Init(uint64_t init, WaitState ws) + virtual bool Init(const amd::Device& dev, uint64_t init, WaitState ws) { return false; } // Blocks the current thread untill the condition c is satisfied diff --git a/projects/clr/rocclr/device/pal/CMakeLists.txt b/projects/clr/rocclr/device/pal/CMakeLists.txt index 5f7c355d9d..121da30707 100644 --- a/projects/clr/rocclr/device/pal/CMakeLists.txt +++ b/projects/clr/rocclr/device/pal/CMakeLists.txt @@ -153,6 +153,7 @@ target_sources(rocclrpal PRIVATE palresource.cpp palschedcl.cpp palsettings.cpp + palsignal.cpp palthreadtrace.cpp paltimestamp.cpp palvirtual.cpp diff --git a/projects/clr/rocclr/device/pal/paldevice.hpp b/projects/clr/rocclr/device/pal/paldevice.hpp index 98c31dff57..778f9f9f86 100644 --- a/projects/clr/rocclr/device/pal/paldevice.hpp +++ b/projects/clr/rocclr/device/pal/paldevice.hpp @@ -36,6 +36,7 @@ #include "device/pal/palsettings.hpp" #include "device/pal/palappprofile.hpp" #include "device/pal/palgpuopen.hpp" +#include "device/pal/palsignal.hpp" #include "acl.h" #include "memory" @@ -361,7 +362,7 @@ class Device : public NullDevice { //! Signal object allocation virtual device::Signal* createSignal() const { - return nullptr; + return new pal::Signal(); } //! Create the device program. diff --git a/projects/clr/rocclr/device/pal/palkernel.cpp b/projects/clr/rocclr/device/pal/palkernel.cpp index 8dc74a1982..dcfaa045c9 100644 --- a/projects/clr/rocclr/device/pal/palkernel.cpp +++ b/projects/clr/rocclr/device/pal/palkernel.cpp @@ -349,6 +349,18 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const WriteAqlArgAt(const_cast
(parameters), &bufferPtr, it.size_, it.offset_); } break; + case amd::KernelParameterDescriptor::HiddenHostcallBuffer: + if (amd::IS_HIP) { + auto buffer = gpu.getOrCreateHostcallBuffer(); + if (!buffer) { + ClPrint(amd::LOG_ERROR, amd::LOG_KERN, + "Kernel expects a hostcall buffer, but none found"); + return false; + } + assert(it.size_ == sizeof(buffer) && "check the sizes"); + WriteAqlArgAt(const_cast
(parameters), &buffer, it.size_, it.offset_); + } + break; case amd::KernelParameterDescriptor::HiddenDefaultQueue: if (vmDefQueue != 0) { WriteAqlArgAt(const_cast
(parameters), &vmDefQueue, it.size_, it.offset_); diff --git a/projects/clr/rocclr/device/pal/palsignal.cpp b/projects/clr/rocclr/device/pal/palsignal.cpp new file mode 100644 index 0000000000..53a9cefc8f --- /dev/null +++ b/projects/clr/rocclr/device/pal/palsignal.cpp @@ -0,0 +1,97 @@ +/* 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 "device/device.hpp" +#include "palsignal.hpp" +#include "paldevice.hpp" +#include "os/os.hpp" + +#include + +namespace pal { + +Signal::~Signal() { + dev_->context().svmFree(amdSignal_); +} + +bool Signal::Init(const amd::Device& dev, uint64_t init, device::Signal::WaitState ws) { + dev_ = static_cast(&dev); + ws_ = ws; + + void* buffer = dev_->context().svmAlloc(sizeof(amd_signal_t), alignof(amd_signal_t), + CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS); + if (!buffer) { + ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, + "Failed to create amd_signal_t buffer"); + return false; + } + std::memset(buffer, 0, sizeof(amd_signal_t)); + + amdSignal_ = new (buffer) amd_signal_t(); + amdSignal_->value = init; + + return true; +} + +uint64_t Signal::Wait(uint64_t value, device::Signal::Condition c, uint64_t timeout) { + auto cmp = [](device::Signal::Condition c) -> std::function { + switch (c) { + case device::Signal::Condition::Eq: + return [](auto ls, auto rs) { return ls == rs; }; + case device::Signal::Condition::Ne: + return [](auto ls, auto rs) { return ls != rs; }; + case device::Signal::Condition::Lt: + return [](auto ls, auto rs) { return ls < rs; }; + case device::Signal::Condition::Gte: + return [](auto ls, auto rs) { return ls >= rs; }; + }; + ShouldNotReachHere(); + return [](auto ls, auto rs) { return false; }; + } (c); + + if (ws_ == device::Signal::WaitState::Blocked) { + guarantee(false, "Unimplemented"); + } else if (ws_ == device::Signal::WaitState::Active) { + auto start = amd::Os::timeNanos(); + while (true) { + auto end = amd::Os::timeNanos(); + auto duration = 1000 * (end - start); // convert to us + if (duration >= timeout) { + return -1; + } + + if (!cmp(amdSignal_->value, value)) { + amd::Os::yield(); + continue; + } + + std::atomic_thread_fence(std::memory_order_acquire); + return amdSignal_->value; + } + } + + return -1; +} + +void Signal::Reset(uint64_t value) { + amdSignal_->value = value; +} + +}; diff --git a/projects/clr/rocclr/device/pal/palsignal.hpp b/projects/clr/rocclr/device/pal/palsignal.hpp new file mode 100644 index 0000000000..b02a43b1ff --- /dev/null +++ b/projects/clr/rocclr/device/pal/palsignal.hpp @@ -0,0 +1,50 @@ +/* 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 pal { + +class Device; + +class Signal: public device::Signal { +private: + const Device* dev_; + amd_signal_t* amdSignal_; + +public: + ~Signal() override; + + bool Init(const amd::Device& dev, 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 amdSignal_; + } +}; + +}; \ No newline at end of file diff --git a/projects/clr/rocclr/device/pal/palvirtual.cpp b/projects/clr/rocclr/device/pal/palvirtual.cpp index 8538f5b96a..9532a3c134 100644 --- a/projects/clr/rocclr/device/pal/palvirtual.cpp +++ b/projects/clr/rocclr/device/pal/palvirtual.cpp @@ -32,6 +32,7 @@ #include "device/pal/palblit.hpp" #include "device/pal/paldebugger.hpp" #include "device/appprofile.hpp" +#include "device/devhostcall.hpp" #include "hsa.h" #include "amd_hsa_kernel_code.h" #include "amd_hsa_queue.h" @@ -1032,6 +1033,9 @@ bool VirtualGPU::create(bool profiling, uint deviceQueueSize, uint rtCUs, &dbg_vmid); } + // The hostcall buffer for this vqueue is initialized on demand. + hostcallBuffer_ = nullptr; + return true; } @@ -1142,6 +1146,14 @@ VirtualGPU::~VirtualGPU() { it->execution().unlock(); } } + + if (hostcallBuffer_ != nullptr) { + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, + "deleting hostcall buffer %p for virtual queue %p", + hostcallBuffer_, this); + disableHostcalls(hostcallBuffer_); + dev().context().svmFree(hostcallBuffer_); + } } void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& vcmd) { @@ -3781,4 +3793,41 @@ void VirtualGPU::submitTransferBufferFromFile(amd::TransferBufferFileCommand& cm } } } + +void* VirtualGPU::getOrCreateHostcallBuffer() { + if (hostcallBuffer_ != nullptr) { + return hostcallBuffer_; + } + + // The number of packets required in each buffer is at least equal to the + // maximum number of waves supported by the device. + auto wavesPerCu = dev().info().maxThreadsPerCU_ / dev().info().wavefrontWidth_; + auto numPackets = dev().info().maxComputeUnits_ * wavesPerCu; + + auto size = getHostcallBufferSize(numPackets); + auto align = getHostcallBufferAlignment(); + + hostcallBuffer_ = dev().context().svmAlloc(size, align, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS); + if (!hostcallBuffer_) { + ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, + "Failed to create hostcall buffer"); + return nullptr; + } + + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, + "Created hostcall buffer %p (numPackets == %d, size == %d, align == %d) for virtual queue %p\n", + hostcallBuffer_, + numPackets, + size, + align, + this); + + if (!enableHostcalls(dev(), hostcallBuffer_, numPackets)) { + ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, + "Failed to register hostcall buffer %p with listener", + hostcallBuffer_); + return nullptr; + } + return hostcallBuffer_; +} } // namespace pal diff --git a/projects/clr/rocclr/device/pal/palvirtual.hpp b/projects/clr/rocclr/device/pal/palvirtual.hpp index d01efda3d7..0614f1721c 100644 --- a/projects/clr/rocclr/device/pal/palvirtual.hpp +++ b/projects/clr/rocclr/device/pal/palvirtual.hpp @@ -552,6 +552,8 @@ class VirtualGPU : public device::VirtualDevice { } } + void* getOrCreateHostcallBuffer(); + protected: void profileEvent(EngineType engine, bool type) const; @@ -673,6 +675,8 @@ class VirtualGPU : public device::VirtualDevice { Queue* queues_[AllEngines]; //!< HW queues for all engines MemoryRange sdmaRange_; //!< SDMA memory range for write access std::vector wrtBackImageBuffer_; //!< Array of images for write back + + void* hostcallBuffer_; //!< Hostcall buffer }; inline void VirtualGPU::logVmMemory(const std::string name, const Memory* memory) { diff --git a/projects/clr/rocclr/device/rocm/rocsignal.cpp b/projects/clr/rocclr/device/rocm/rocsignal.cpp index c252075e43..013f4f3e55 100644 --- a/projects/clr/rocclr/device/rocm/rocsignal.cpp +++ b/projects/clr/rocclr/device/rocm/rocsignal.cpp @@ -26,7 +26,7 @@ Signal::~Signal() { hsa_signal_destroy(signal_); } -bool Signal::Init(uint64_t init, device::Signal::WaitState ws) { +bool Signal::Init(const amd::Device& dev, 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; diff --git a/projects/clr/rocclr/device/rocm/rocsignal.hpp b/projects/clr/rocclr/device/rocm/rocsignal.hpp index 1c54fce40b..2350c14e2f 100644 --- a/projects/clr/rocclr/device/rocm/rocsignal.hpp +++ b/projects/clr/rocclr/device/rocm/rocsignal.hpp @@ -33,7 +33,7 @@ private: public: ~Signal() override; - bool Init(uint64_t init, device::Signal::WaitState ws) override; + bool Init(const amd::Device& dev, uint64_t init, device::Signal::WaitState ws) override; uint64_t Wait(uint64_t value, device::Signal::Condition c, uint64_t timeout) override;