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: 99e8ac55cd]

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<HostcallBuffer*>(bfr);
   buffer->initialize(numPackets);
 

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);

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

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

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.

projects/clr/rocclr/device/pal/palkernel.cpp

@@ -349,6 +349,18 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const
           WriteAqlArgAt(const_cast<address>(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<address>(parameters), &buffer, it.size_, it.offset_);
+        }
+        break;
       case amd::KernelParameterDescriptor::HiddenDefaultQueue:
         if (vmDefQueue != 0) {
           WriteAqlArgAt(const_cast<address>(parameters), &vmDefQueue, it.size_, it.offset_);

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 <functional>
+
+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<const pal::Device*>(&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<bool(uint64_t, uint64_t)> {
+    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;
+}
+
+};

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 <amd_hsa_signal.h>
+
+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_;
+  }
+};
+
+};

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

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<Image*> wrtBackImageBuffer_;  //!< Array of images for write back
+
+  void* hostcallBuffer_;  //!< Hostcall buffer
 };
 
 inline void VirtualGPU::logVmMemory(const std::string name, const Memory* memory) {

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;

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;