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

@@ -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

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 <hsa.h>
-
 #include <assert.h>
 #include <string.h>
 #include <set>
@@ -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<HostcallBuffer>::value,
@@ -262,7 +259,7 @@ void HostcallBuffer::initialize(uint32_t num_packets) {
  */
 class HostcallListener {
   std::set<HostcallBuffer*> 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<HostcallBuffer*>(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;

rocclr/device/devhostcall.hpp

@@ -20,6 +20,11 @@
 
 #pragma once
 
+#include "top.hpp"
+#include "device/device.hpp"
+
+#include <cstddef>
+
 /** \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);

rocclr/device/device.hpp

@@ -36,6 +36,7 @@
 #include "amdocl/cl_profile_amd.h"
 #include "acl.h"
 #include "hwdebug.hpp"
+#include "devsignal.hpp"
 
 #include <cassert>
 #include <cstdint>
@@ -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.

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; }
+};
+
+};

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

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

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

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 <algorithm>
 #include <cstring>
@@ -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

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,

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<hsa_signal_condition_t>(c),
+    value,
+    timeout,
+    static_cast<hsa_wait_state_t>(ws_));
+}
+
+void Signal::Reset(uint64_t value) {
+  hsa_signal_store_screlease(signal_, value);
+}
+
+};

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 <hsa.h>
+
+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<void*>(signal_.handle);
+  }
+};
+
+};

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"