Initial IPC signal support.

Added an API for creating signals with attributes.
Added two APIs for IPC operations on signals.
Initial use of exceptions for error handling.

Add ref counting to signals.
Removed spin loops from signal destructors.
Signals are no longer to be destroyed with delete, use DeleteSignal instead.
Added delete safety to doorbells.
Added secondary hsa_signal_t -> Signal* translation path for IPC enabled signals.

Change-Id: Id59065d002f0c2566b0a9425694da2ed27cb7d7f

runtime/hsa-runtime/CMakeLists.txt

@@ -133,6 +133,7 @@ set ( SRCS "core/util/lnx/os_linux.cpp"
             "core/runtime/hsa_ext_amd.cpp"
             "core/runtime/hsa_ext_interface.cpp"
             "core/runtime/interrupt_signal.cpp"
+            "core/runtime/ipc_signal.cpp"
             "core/runtime/isa.cpp"
             "core/runtime/runtime.cpp"
             "core/runtime/signal.cpp"

runtime/hsa-runtime/core/common/hsa_table_interface.cpp

@@ -1075,3 +1075,22 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t l
 hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr) {
   return amdExtTable->hsa_amd_ipc_memory_detach_fn(mapped_ptr);
 }
+
+// Mirrors Amd Extension Apis
+hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
+                                   const hsa_agent_t* consumers, uint64_t attributes,
+                                   hsa_signal_t* signal) {
+  return amdExtTable->hsa_amd_signal_create_fn(initial_value, num_consumers, consumers, attributes,
+                                               signal);
+}
+
+// Mirrors Amd Extension Apis
+hsa_status_t HSA_API hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle) {
+  return amdExtTable->hsa_amd_ipc_signal_create_fn(signal, handle);
+}
+
+// Mirrors Amd Extension Apis
+hsa_status_t HSA_API hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle,
+                                               hsa_signal_t* signal) {
+  return amdExtTable->hsa_amd_ipc_signal_attach_fn(handle, signal);
+}

runtime/hsa-runtime/core/inc/amd_aql_queue.h

@@ -339,6 +339,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
  protected:
   bool _IsA(rtti_t id) const override { return id == &rtti_id_; }
 
+  /// @brief Disallow destroying doorbell apart from its queue.
+  void doDestroySignal() override { assert(false); }
+
  private:
   uint32_t ComputeRingBufferMinPkts();
   uint32_t ComputeRingBufferMaxPkts();

runtime/hsa-runtime/core/inc/default_signal.h

@@ -51,21 +51,18 @@
 
 namespace core {
 
-/// @brief Simple pure memory based signal.
+/// @brief Operations for a simple pure memory based signal.
 /// @brief See base class Signal.
-class DefaultSignal : private LocalSignal, public Signal {
+class BusyWaitSignal : public Signal {
  public:
-  /// @brief Determines if a Signal* can be safely converted to DefaultSignal*
+  /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
   /// via static_cast.
   static __forceinline bool IsType(Signal* ptr) {
     return ptr->IsType(&rtti_id_);
   }
 
   /// @brief See base class Signal.
-  explicit DefaultSignal(hsa_signal_value_t initial_value);
-
-  /// @brief See base class Signal.
-  ~DefaultSignal();
+  explicit BusyWaitSignal(SharedSignal* abi_block, bool enableIPC);
 
   // Below are various methods corresponding to the APIs, which load/store the
   // signal value or modify the existing signal value automically and with
@@ -158,6 +155,29 @@ class DefaultSignal : private LocalSignal, public Signal {
  protected:
   bool _IsA(rtti_t id) const { return id == &rtti_id_; }
 
+ private:
+  static int rtti_id_;
+
+  DISALLOW_COPY_AND_ASSIGN(BusyWaitSignal);
+};
+
+/// @brief Simple memory only signal using a new ABI block.
+class DefaultSignal : private LocalSignal, public BusyWaitSignal {
+ public:
+  /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
+  /// via static_cast.
+  static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); }
+
+  /// @brief See base class Signal.
+  explicit DefaultSignal(hsa_signal_value_t initial_value, bool enableIPC = false)
+      : LocalSignal(initial_value), BusyWaitSignal(signal(), enableIPC) {}
+
+ protected:
+  bool _IsA(rtti_t id) const {
+    if (id == &rtti_id_) return true;
+    return BusyWaitSignal::_IsA(id);
+  }
+
  private:
   static int rtti_id_;
 

runtime/hsa-runtime/core/inc/exceptions.h

@@ -0,0 +1,67 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// The University of Illinois/NCSA
+// Open Source License (NCSA)
+//
+// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
+//
+// Developed by:
+//
+//                 AMD Research and AMD HSA Software Development
+//
+//                 Advanced Micro Devices, Inc.
+//
+//                 www.amd.com
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to
+// deal with 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:
+//
+//  - Redistributions of source code must retain the above copyright notice,
+//    this list of conditions and the following disclaimers.
+//  - Redistributions in binary form must reproduce the above copyright
+//    notice, this list of conditions and the following disclaimers in
+//    the documentation and/or other materials provided with the distribution.
+//  - Neither the names of Advanced Micro Devices, Inc,
+//    nor the names of its contributors may be used to endorse or promote
+//    products derived from this Software without specific prior written
+//    permission.
+//
+// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+//
+////////////////////////////////////////////////////////////////////////////////
+
+#ifndef HSA_RUNTIME_CORE_INC_EXCEPTIONS_H
+#define HSA_RUNTIME_CORE_INC_EXCEPTIONS_H
+
+#include <exception>
+#include <string>
+
+#include "core/inc/hsa_internal.h"
+
+namespace AMD {
+
+/// @brief Exception type which carries an error code to return to the user.
+class hsa_exception : public std::exception {
+ public:
+  hsa_exception(hsa_status_t error, const char* description) : err_(error), desc_(description) {}
+  hsa_status_t error_code() const noexcept { return err_; }
+  const char* what() const noexcept override { return desc_.c_str(); }
+
+ private:
+  hsa_status_t err_;
+  std::string desc_;
+};
+
+}  // namespace AMD
+
+#endif  // HSA_RUNTIME_CORE_INC_EXCEPTIONS_H

runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h

@@ -94,6 +94,11 @@ hsa_status_t HSA_API
 hsa_status_t HSA_API
     hsa_amd_async_function(void (*callback)(void* arg), void* arg);
 
+// Mirrors Amd Extension Apis
+hsa_status_t HSA_API hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
+                                           const hsa_agent_t* consumers, uint64_t attributes,
+                                           hsa_signal_t* signal);
+
 // Mirrors Amd Extension Apis
 uint32_t HSA_API
     hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* signals,
@@ -199,6 +204,12 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* handle, size_
 // Mirrors Amd Extension Apis
 hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr);
 
+// Mirrors Amd Extension Apis
+hsa_status_t hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle);
+
+// Mirrors Amd Extension Apis
+hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, hsa_signal_t* signal);
+
 }  // end of AMD namespace
 
 #endif  // header guard

runtime/hsa-runtime/core/inc/ipc_signal.h

@@ -0,0 +1,112 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// The University of Illinois/NCSA
+// Open Source License (NCSA)
+//
+// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
+//
+// Developed by:
+//
+//                 AMD Research and AMD HSA Software Development
+//
+//                 Advanced Micro Devices, Inc.
+//
+//                 www.amd.com
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to
+// deal with 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:
+//
+//  - Redistributions of source code must retain the above copyright notice,
+//    this list of conditions and the following disclaimers.
+//  - Redistributions in binary form must reproduce the above copyright
+//    notice, this list of conditions and the following disclaimers in
+//    the documentation and/or other materials provided with the distribution.
+//  - Neither the names of Advanced Micro Devices, Inc,
+//    nor the names of its contributors may be used to endorse or promote
+//    products derived from this Software without specific prior written
+//    permission.
+//
+// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+//
+////////////////////////////////////////////////////////////////////////////////
+
+#ifndef HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_
+#define HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_
+
+#include <atomic>
+#include <utility>
+
+#include "core/inc/signal.h"
+#include "core/inc/default_signal.h"
+#include "core/util/locks.h"
+
+namespace core {
+
+/// @brief Container for ipc shared memory.
+class SharedMemory {
+ public:
+  SharedMemory(const hsa_amd_ipc_memory_t* handle, size_t len);
+  ~SharedMemory();
+  SharedMemory(SharedMemory&&);
+
+  void* ptr() const { return ptr_; }
+
+ private:
+  void* ptr_;
+};
+
+/// @brief Container for ipc signal abi block.
+class SharedMemorySignal {
+ public:
+  explicit SharedMemorySignal(const hsa_amd_ipc_memory_t* handle) : signal_(handle, 4096) {
+    if (!signal()->IsValid())
+      throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "IPC Signal handle is invalid.");
+  }
+  SharedSignal* signal() const { return reinterpret_cast<SharedSignal*>(signal_.ptr()); }
+
+ private:
+  SharedMemory signal_;
+};
+
+/// @brief Memory only signal using a shared memory ABI block.
+class IPCSignal : private SharedMemorySignal, public BusyWaitSignal {
+ public:
+  /// @brief Creates a sharable handle for an IPC enabled signal.
+  static void CreateHandle(Signal* signal, hsa_amd_ipc_signal_t* ipc_handle);
+
+  /// @brief Opens an IPC signal from its IPC handle.
+  static Signal* Attach(const hsa_amd_ipc_signal_t* ipc_handle);
+
+  /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
+  /// via static_cast.
+  static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); }
+
+ protected:
+  bool _IsA(rtti_t id) const {
+    if (id == &rtti_id_) return true;
+    return BusyWaitSignal::_IsA(id);
+  }
+
+ private:
+  static int rtti_id_;
+  static KernelMutex lock_;
+
+  explicit IPCSignal(SharedMemorySignal&& abi_block)
+      : SharedMemorySignal(std::move(abi_block)), BusyWaitSignal(signal(), true) {}
+
+  DISALLOW_COPY_AND_ASSIGN(IPCSignal);
+};
+
+}  // namespace core
+
+#endif  // HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_

runtime/hsa-runtime/core/inc/signal.h

@@ -45,12 +45,15 @@
 #ifndef HSA_RUNTME_CORE_INC_SIGNAL_H_
 #define HSA_RUNTME_CORE_INC_SIGNAL_H_
 
+#include <map>
+
 #include "hsakmt.h"
 
 #include "core/common/shared.h"
 
 #include "core/inc/runtime.h"
 #include "core/inc/checked.h"
+#include "core/inc/exceptions.h"
 
 #include "core/util/utils.h"
 #include "core/util/locks.h"
@@ -67,20 +70,28 @@ struct SharedSignal {
   Signal* core_signal;
   Check<0x71FCCA6A3D5D5276, true> id;
 
-  bool IsValid() const { return id.IsValid(); }
-
   SharedSignal() {
+    memset(&amd_signal, 0, sizeof(amd_signal));
     amd_signal.kind = AMD_SIGNAL_KIND_INVALID;
     core_signal = nullptr;
   }
 
+  bool IsValid() const { return (Convert(this).handle != 0) && id.IsValid(); }
+
+  bool IsIPC() const { return core_signal == nullptr; }
+
   static __forceinline SharedSignal* Convert(hsa_signal_t signal) {
-    if (signal.handle == 0) throw std::bad_cast();
     SharedSignal* ret = reinterpret_cast<SharedSignal*>(static_cast<uintptr_t>(signal.handle) -
                                                         offsetof(SharedSignal, amd_signal));
-    if (!ret->IsValid()) throw std::bad_cast();
     return ret;
   }
+
+  static __forceinline hsa_signal_t Convert(const SharedSignal* signal) {
+    assert(signal != nullptr && "Conversion on null Signal object.");
+    const uint64_t handle = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(&signal->amd_signal));
+    const hsa_signal_t signal_handle = {handle};
+    return signal_handle;
+  }
 };
 static_assert(std::is_standard_layout<SharedSignal>::value,
               "SharedSignal must remain standard layout for IPC use.");
@@ -107,49 +118,86 @@ class LocalSignal {
 class Signal {
  public:
   /// @brief Constructor Links and publishes the signal interface object.
-  explicit Signal(SharedSignal* abi_block)
-      : signal_(abi_block->amd_signal), async_copy_agent_(NULL) {
-    if (abi_block == nullptr) throw std::bad_alloc();
+  explicit Signal(SharedSignal* abi_block, bool enableIPC = false)
+      : signal_(abi_block->amd_signal), async_copy_agent_(NULL), refcount_(1) {
+    assert(abi_block != nullptr && "Signal abi_block must not be NULL");
 
-    invalid_ = false;
     waiting_ = 0;
-    retained_ = 0;
+    retained_ = 1;
 
-    abi_block->core_signal = this;
+    if (enableIPC) {
+      abi_block->core_signal = nullptr;
+      registerIpc();
+    } else {
+      abi_block->core_signal = this;
+    }
   }
 
-  virtual ~Signal() { invalid_ = true; }
+  /// @brief Interface to discard a signal handle (hsa_signal_t)
+  /// Decrements signal ref count and invokes doDestroySignal() when
+  /// Signal is no longer in use.
+  void DestroySignal() {
+    // If handle is now invalid wake any retained sleepers.
+    if (--refcount_ == 0) CasRelaxed(0, 0);
+    // Release signal, last release will destroy the object.
+    Release();
+  }
 
-  bool IsValid() const { return !invalid_; }
-
-  /// @brief Converts from this implementation class to the public
+  /// @brief Converts from this interface class to the public
   /// hsa_signal_t type - an opaque handle.
   static __forceinline hsa_signal_t Convert(Signal* signal) {
-    assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object.");
+    assert(signal != nullptr && "Conversion on null Signal object.");
     const uint64_t handle = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(&signal->signal_));
     const hsa_signal_t signal_handle = {handle};
     return signal_handle;
   }
 
-  /// @brief Converts from this implementation class to the public
+  /// @brief Converts from this interface class to the public
   /// hsa_signal_t type - an opaque handle.
   static __forceinline const hsa_signal_t Convert(const Signal* signal) {
-    assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object.");
+    assert(signal != nullptr && "Conversion on null Signal object.");
     const uint64_t handle = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(&signal->signal_));
     const hsa_signal_t signal_handle = {handle};
     return signal_handle;
   }
 
   /// @brief Converts from public hsa_signal_t type (an opaque handle) to
-  /// this implementation class object.
+  /// this interface class object.
   static __forceinline Signal* Convert(hsa_signal_t signal) {
+    if (signal.handle == 0)
+      throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "Signal handle is invalid.");
     SharedSignal* shared = SharedSignal::Convert(signal);
-    if (shared->core_signal != nullptr)
+    if (!shared->IsValid())
+      throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_SIGNAL, "Signal handle is invalid.");
+    if (shared->IsIPC()) {
+      Signal* ret = lookupIpc(signal);
+      if (ret == nullptr)
+        throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_SIGNAL, "Signal handle is invalid.");
+      return ret;
+    } else {
       return shared->core_signal;
-    else
-      throw std::bad_cast();
+    }
   }
 
+  static Signal* DuplicateHandle(hsa_signal_t signal) {
+    if (signal.handle == 0) return nullptr;
+    SharedSignal* shared = SharedSignal::Convert(signal);
+
+    if (!shared->IsIPC()) {
+      if (!shared->IsValid()) return nullptr;
+      shared->core_signal->refcount_++;
+      shared->core_signal->Retain();
+      return shared->core_signal;
+    }
+
+    // IPC signals may only be duplicated while holding the ipcMap lock.
+    return duplicateIpc(signal);
+  }
+
+  bool IsValid() const { return refcount_ != 0; }
+
+  bool __forceinline isIPC() const { return SharedSignal::Convert(Convert(this))->IsIPC(); }
+
   // Below are various methods corresponding to the APIs, which load/store the
   // signal value or modify the existing signal value automically and with
   // specified memory ordering semantics.
@@ -230,13 +278,9 @@ class Signal {
 
   __forceinline bool IsType(rtti_t id) { return _IsA(id); }
 
-  /// @brief Allows special case interaction with signal destruction cleanup.
-  void Retain() { atomic::Increment(&retained_); }
-  void Release() { atomic::Decrement(&retained_); }
-
-  /// @brief Checks if signal is currently in use such that it should not be
-  /// deleted.
-  bool InUse() const { return (retained_ != 0) || (waiting_ != 0); }
+  /// @brief Prevents the signal from being destroyed until the matching Release().
+  void Retain() { retained_++; }
+  void Release();
 
   /// @brief Checks if signal is currently in use by a wait API.
   bool InWaiting() const { return waiting_ != 0; }
@@ -253,25 +297,40 @@ class Signal {
   amd_signal_t& signal_;
 
  protected:
+  virtual ~Signal();
+
+  /// @brief Overrideable deletion function
+  virtual void doDestroySignal() { delete this; }
+
   /// @brief Simple RTTI type checking helper
   /// Returns true if the object can be converted to the query type via
   /// static_cast.
   /// Do not use directly.  Use IsType in the desired derived type instead.
   virtual bool _IsA(rtti_t id) const = 0;
 
-  /// @variable  Indicates if signal is valid or not.
-  volatile bool invalid_;
-
   /// @variable Indicates number of runtime threads waiting on this signal.
   /// Value of zero means no waits.
-  volatile uint32_t waiting_;
-
-  volatile uint32_t retained_;
+  std::atomic<uint32_t> waiting_;
 
   /// @variable Pointer to agent used to perform an async copy.
   core::Agent* async_copy_agent_;
 
  private:
+  static KernelMutex ipcLock_;
+  static std::map<decltype(hsa_signal_t::handle), Signal*> ipcMap_;
+
+  static Signal* lookupIpc(hsa_signal_t signal);
+  static Signal* duplicateIpc(hsa_signal_t signal);
+
+  /// @variable Ref count of this signal's handle (see IPC APIs)
+  std::atomic<uint32_t> refcount_;
+
+  /// @variable Count of handle references and Retain() calls for this handle (see IPC APIs)
+  std::atomic<uint32_t> retained_;
+
+  void registerIpc();
+  bool deregisterIpc();
+
   DISALLOW_COPY_AND_ASSIGN(Signal);
 };
 

runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp

@@ -582,6 +582,7 @@ void GpuAgent::InitDma() {
 
       if (blits_[BlitDevToHost] == NULL) {
         // Share utility queue with device-to-host blits.
+        if (queues_[QueueUtility] == nullptr) queues_[QueueUtility] = CreateInterceptibleQueue();
         blits_[BlitDevToHost] = CreateBlitKernel(queues_[QueueUtility]);
         assert(blits_[BlitDevToHost] != NULL && "Blit creation failed");
       }
@@ -597,14 +598,15 @@ hsa_status_t GpuAgent::PostToolsInit() {
   BindTrapHandler();
 
   // Defer utility queue creation to allow tools to intercept.
-  queues_[QueueUtility] = CreateInterceptibleQueue();
+  if (queues_[QueueUtility] == nullptr) queues_[QueueUtility] = CreateInterceptibleQueue();
 
   if (queues_[QueueUtility] == NULL) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
   }
 
   // Share utility queue with device-to-device blits.
-  blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]);
+  if (blits_[BlitDevToDev] == nullptr)
+    blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]);
 
   if (blits_[BlitDevToDev] == NULL) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;

runtime/hsa-runtime/core/runtime/default_signal.cpp

@@ -46,49 +46,45 @@
 namespace core {
 
 int DefaultSignal::rtti_id_ = 0;
+int BusyWaitSignal::rtti_id_ = 0;
 
-DefaultSignal::DefaultSignal(hsa_signal_value_t initial_value)
-    : LocalSignal(initial_value), Signal(signal()) {
+BusyWaitSignal::BusyWaitSignal(SharedSignal* abi_block, bool enableIPC)
+    : Signal(abi_block, enableIPC) {
   signal_.kind = AMD_SIGNAL_KIND_USER;
   signal_.event_mailbox_ptr = NULL;
-  HSA::hsa_memory_register(this, sizeof(DefaultSignal));
 }
 
-DefaultSignal::~DefaultSignal() {
-  invalid_ = true;
-  while (InUse())
-    ;
-  HSA::hsa_memory_deregister(this, sizeof(DefaultSignal));
-}
-
-hsa_signal_value_t DefaultSignal::LoadRelaxed() {
+hsa_signal_value_t BusyWaitSignal::LoadRelaxed() {
   return hsa_signal_value_t(
       atomic::Load(&signal_.value, std::memory_order_relaxed));
 }
 
-hsa_signal_value_t DefaultSignal::LoadAcquire() {
+hsa_signal_value_t BusyWaitSignal::LoadAcquire() {
   return hsa_signal_value_t(
       atomic::Load(&signal_.value, std::memory_order_acquire));
 }
 
-void DefaultSignal::StoreRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::StoreRelaxed(hsa_signal_value_t value) {
   atomic::Store(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::StoreRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::StoreRelease(hsa_signal_value_t value) {
   atomic::Store(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition,
-                                              hsa_signal_value_t compare_value,
-                                              uint64_t timeout,
-                                              hsa_wait_state_t wait_hint) {
-  atomic::Increment(&waiting_);
-  MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); });
+hsa_signal_value_t BusyWaitSignal::WaitRelaxed(hsa_signal_condition_t condition,
+                                               hsa_signal_value_t compare_value, uint64_t timeout,
+                                               hsa_wait_state_t wait_hint) {
+  Retain();
+  MAKE_SCOPE_GUARD([&]() { Release(); });
+
+  waiting_++;
+  MAKE_SCOPE_GUARD([&]() { waiting_--; });
   bool condition_met = false;
   int64_t value;
 
-  debug_warning(!g_use_interrupt_wait && "Use of non-host signal in host signal wait API.");
+  debug_warning((!g_use_interrupt_wait || isIPC()) &&
+                "Use of non-host signal in host signal wait API.");
 
   timer::fast_clock::time_point start_time, time;
   start_time = timer::fast_clock::now();
@@ -100,7 +96,7 @@ hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition,
           double(timeout) / double(hsa_freq));
 
   while (true) {
-    if (invalid_) return 0;
+    if (!IsValid()) return 0;
 
     value = atomic::Load(&signal_.value, std::memory_order_relaxed);
 
@@ -135,139 +131,138 @@ hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition,
   }
 }
 
-hsa_signal_value_t DefaultSignal::WaitAcquire(hsa_signal_condition_t condition,
-                                              hsa_signal_value_t compare_value,
-                                              uint64_t timeout,
-                                              hsa_wait_state_t wait_hint) {
+hsa_signal_value_t BusyWaitSignal::WaitAcquire(hsa_signal_condition_t condition,
+                                               hsa_signal_value_t compare_value, uint64_t timeout,
+                                               hsa_wait_state_t wait_hint) {
   hsa_signal_value_t ret =
       WaitRelaxed(condition, compare_value, timeout, wait_hint);
   std::atomic_thread_fence(std::memory_order_acquire);
   return ret;
 }
 
-void DefaultSignal::AndRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::AndRelaxed(hsa_signal_value_t value) {
   atomic::And(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::AndAcquire(hsa_signal_value_t value) {
+void BusyWaitSignal::AndAcquire(hsa_signal_value_t value) {
   atomic::And(&signal_.value, int64_t(value), std::memory_order_acquire);
 }
 
-void DefaultSignal::AndRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::AndRelease(hsa_signal_value_t value) {
   atomic::And(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-void DefaultSignal::AndAcqRel(hsa_signal_value_t value) {
+void BusyWaitSignal::AndAcqRel(hsa_signal_value_t value) {
   atomic::And(&signal_.value, int64_t(value), std::memory_order_acq_rel);
 }
 
-void DefaultSignal::OrRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::OrRelaxed(hsa_signal_value_t value) {
   atomic::Or(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::OrAcquire(hsa_signal_value_t value) {
+void BusyWaitSignal::OrAcquire(hsa_signal_value_t value) {
   atomic::Or(&signal_.value, int64_t(value), std::memory_order_acquire);
 }
 
-void DefaultSignal::OrRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::OrRelease(hsa_signal_value_t value) {
   atomic::Or(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-void DefaultSignal::OrAcqRel(hsa_signal_value_t value) {
+void BusyWaitSignal::OrAcqRel(hsa_signal_value_t value) {
   atomic::Or(&signal_.value, int64_t(value), std::memory_order_acq_rel);
 }
 
-void DefaultSignal::XorRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::XorRelaxed(hsa_signal_value_t value) {
   atomic::Xor(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::XorAcquire(hsa_signal_value_t value) {
+void BusyWaitSignal::XorAcquire(hsa_signal_value_t value) {
   atomic::Xor(&signal_.value, int64_t(value), std::memory_order_acquire);
 }
 
-void DefaultSignal::XorRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::XorRelease(hsa_signal_value_t value) {
   atomic::Xor(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-void DefaultSignal::XorAcqRel(hsa_signal_value_t value) {
+void BusyWaitSignal::XorAcqRel(hsa_signal_value_t value) {
   atomic::Xor(&signal_.value, int64_t(value), std::memory_order_acq_rel);
 }
 
-void DefaultSignal::AddRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::AddRelaxed(hsa_signal_value_t value) {
   atomic::Add(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::AddAcquire(hsa_signal_value_t value) {
+void BusyWaitSignal::AddAcquire(hsa_signal_value_t value) {
   atomic::Add(&signal_.value, int64_t(value), std::memory_order_acquire);
 }
 
-void DefaultSignal::AddRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::AddRelease(hsa_signal_value_t value) {
   atomic::Add(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-void DefaultSignal::AddAcqRel(hsa_signal_value_t value) {
+void BusyWaitSignal::AddAcqRel(hsa_signal_value_t value) {
   atomic::Add(&signal_.value, int64_t(value), std::memory_order_acq_rel);
 }
 
-void DefaultSignal::SubRelaxed(hsa_signal_value_t value) {
+void BusyWaitSignal::SubRelaxed(hsa_signal_value_t value) {
   atomic::Sub(&signal_.value, int64_t(value), std::memory_order_relaxed);
 }
 
-void DefaultSignal::SubAcquire(hsa_signal_value_t value) {
+void BusyWaitSignal::SubAcquire(hsa_signal_value_t value) {
   atomic::Sub(&signal_.value, int64_t(value), std::memory_order_acquire);
 }
 
-void DefaultSignal::SubRelease(hsa_signal_value_t value) {
+void BusyWaitSignal::SubRelease(hsa_signal_value_t value) {
   atomic::Sub(&signal_.value, int64_t(value), std::memory_order_release);
 }
 
-void DefaultSignal::SubAcqRel(hsa_signal_value_t value) {
+void BusyWaitSignal::SubAcqRel(hsa_signal_value_t value) {
   atomic::Sub(&signal_.value, int64_t(value), std::memory_order_acq_rel);
 }
 
-hsa_signal_value_t DefaultSignal::ExchRelaxed(hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::ExchRelaxed(hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value),
                                              std::memory_order_relaxed));
 }
 
-hsa_signal_value_t DefaultSignal::ExchAcquire(hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::ExchAcquire(hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value),
                                              std::memory_order_acquire));
 }
 
-hsa_signal_value_t DefaultSignal::ExchRelease(hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::ExchRelease(hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value),
                                              std::memory_order_release));
 }
 
-hsa_signal_value_t DefaultSignal::ExchAcqRel(hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::ExchAcqRel(hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value),
                                              std::memory_order_acq_rel));
 }
 
-hsa_signal_value_t DefaultSignal::CasRelaxed(hsa_signal_value_t expected,
-                                             hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::CasRelaxed(hsa_signal_value_t expected,
+                                              hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value),
                                         int64_t(expected),
                                         std::memory_order_relaxed));
 }
 
-hsa_signal_value_t DefaultSignal::CasAcquire(hsa_signal_value_t expected,
-                                             hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::CasAcquire(hsa_signal_value_t expected,
+                                              hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value),
                                         int64_t(expected),
                                         std::memory_order_acquire));
 }
 
-hsa_signal_value_t DefaultSignal::CasRelease(hsa_signal_value_t expected,
-                                             hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::CasRelease(hsa_signal_value_t expected,
+                                              hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value),
                                         int64_t(expected),
                                         std::memory_order_release));
 }
 
-hsa_signal_value_t DefaultSignal::CasAcqRel(hsa_signal_value_t expected,
-                                            hsa_signal_value_t value) {
+hsa_signal_value_t BusyWaitSignal::CasAcqRel(hsa_signal_value_t expected,
+                                             hsa_signal_value_t value) {
   return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value),
                                         int64_t(expected),
                                         std::memory_order_acq_rel));

runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp

@@ -378,6 +378,9 @@ void HsaApiTable::UpdateAmdExts() {
   amd_ext_api.hsa_amd_ipc_memory_create_fn = AMD::hsa_amd_ipc_memory_create;
   amd_ext_api.hsa_amd_ipc_memory_attach_fn = AMD::hsa_amd_ipc_memory_attach;
   amd_ext_api.hsa_amd_ipc_memory_detach_fn = AMD::hsa_amd_ipc_memory_detach;
+  amd_ext_api.hsa_amd_signal_create_fn = AMD::hsa_amd_signal_create;
+  amd_ext_api.hsa_amd_ipc_signal_create_fn = AMD::hsa_amd_ipc_signal_create;
+  amd_ext_api.hsa_amd_ipc_signal_attach_fn = AMD::hsa_amd_ipc_signal_attach;
 }
 
 class Init {

runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp

@@ -43,6 +43,7 @@
 #include <new>
 #include <typeinfo>
 #include <exception>
+#include <set>
 
 #include "hsakmt.h"
 
@@ -52,7 +53,10 @@
 #include "core/inc/amd_gpu_agent.h"
 #include "core/inc/amd_memory_region.h"
 #include "core/inc/signal.h"
+#include "core/inc/default_signal.h"
 #include "core/inc/interrupt_signal.h"
+#include "core/inc/ipc_signal.h"
+#include "core/inc/exceptions.h"
 
 template <class T>
 struct ValidityError;
@@ -124,11 +128,13 @@ hsa_status_t handleException() {
     throw;
   } catch (const std::bad_alloc& e) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-  } catch (const std::bad_cast& e) {
-    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
-  } catch (std::nested_exception& e) {  // Rethrow exceptions from callbacks after unwinding HSA.
-    e.rethrow_nested();
-    return HSA_STATUS_ERROR;
+  } catch (const hsa_exception& e) {
+    return e.error_code();
+    // Enable when callback exception support is added.
+    // } catch (std::nested_exception& e) {  // Rethrow exceptions from callbacks after unwinding
+    // HSA.
+    //  e.rethrow_nested();
+    //  return HSA_STATUS_ERROR;
   } catch (...) {
     assert(false && "Unhandled exception.");
     abort();
@@ -143,7 +149,7 @@ template <class T> static __forceinline T handleExceptionT() {
 }
 
 hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_coherency_type_t* type) {
-  TRY
+  TRY;
   IS_OPEN();
 
   const core::Agent* agent = core::Agent::Convert(agent_handle);
@@ -162,12 +168,12 @@ hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_cohere
   *type = gpu_agent->current_coherency_type();
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle,
                                         hsa_amd_coherency_type_t type) {
-  TRY
+  TRY;
   IS_OPEN();
 
   core::Agent* agent = core::Agent::Convert(agent_handle);
@@ -190,11 +196,11 @@ hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle,
   }
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (ptr == NULL) {
@@ -206,14 +212,14 @@ hsa_status_t hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) {
   }
 
   return core::Runtime::runtime_singleton_->FillMemory(ptr, value, count);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent_handle, const void* src,
                                        hsa_agent_t src_agent_handle, size_t size,
                                        uint32_t num_dep_signals, const hsa_signal_t* dep_signals,
                                        hsa_signal_t completion_signal) {
-  TRY
+  TRY;
   if (dst == NULL || src == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; }
 
   if ((num_dep_signals == 0 && dep_signals != NULL) ||
@@ -246,11 +252,11 @@ hsa_status_t hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent_handle,
   }
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) {
-  TRY
+  TRY;
   IS_OPEN();
 
   core::Queue* cmd_queue = core::Queue::Convert(queue);
@@ -261,11 +267,11 @@ hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enab
                    AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, (enable != 0));
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) {
-  TRY
+  TRY;
   IS_OPEN();
 
   return core::Runtime::runtime_singleton_->IterateAgent(
@@ -274,13 +280,13 @@ hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) {
         return core::Agent::Convert(agent_handle)->profiling_enabled(enable);
       },
       reinterpret_cast<void*>(&enable));
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_profiling_get_dispatch_time(
     hsa_agent_t agent_handle, hsa_signal_t hsa_signal,
     hsa_amd_profiling_dispatch_time_t* time) {
-  TRY
+  TRY;
   IS_OPEN();
 
   IS_BAD_PTR(time);
@@ -303,12 +309,12 @@ hsa_status_t hsa_amd_profiling_get_dispatch_time(
   gpu_agent->TranslateTime(signal, *time);
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_profiling_get_async_copy_time(
     hsa_signal_t hsa_signal, hsa_amd_profiling_async_copy_time_t* time) {
-  TRY
+  TRY;
   IS_OPEN();
 
   IS_BAD_PTR(time);
@@ -333,13 +339,13 @@ hsa_status_t hsa_amd_profiling_get_async_copy_time(
   time->start = signal->signal_.start_ts;
   time->end = signal->signal_.end_ts;
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_handle,
                                                              uint64_t agent_tick,
                                                              uint64_t* system_tick) {
-  TRY
+  TRY;
   IS_OPEN();
 
   IS_BAD_PTR(system_tick);
@@ -357,42 +363,89 @@ hsa_status_t hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_h
   *system_tick = gpu_agent->TranslateTime(agent_tick);
 
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
+}
+
+hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
+                                   const hsa_agent_t* consumers, uint64_t attributes,
+                                   hsa_signal_t* hsa_signal) {
+  struct AgentHandleCompare {
+    bool operator()(const hsa_agent_t& lhs, const hsa_agent_t& rhs) const {
+      return lhs.handle < rhs.handle;
+    }
+  };
+
+  TRY;
+  IS_OPEN();
+  IS_BAD_PTR(hsa_signal);
+
+  core::Signal* ret;
+
+  bool enable_ipc = attributes & HSA_AMD_SIGNAL_IPC;
+  bool use_default =
+      enable_ipc || (attributes & HSA_AMD_SIGNAL_AMD_GPU_ONLY) || (!core::g_use_interrupt_wait);
+
+  if ((!use_default) && (num_consumers != 0)) {
+    IS_BAD_PTR(consumers);
+
+    // Check for duplicates in consumers.
+    std::set<hsa_agent_t, AgentHandleCompare> consumer_set(consumers, consumers + num_consumers);
+    if (consumer_set.size() != num_consumers) {
+      return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+    }
+
+    use_default = true;
+    for (const core::Agent* cpu_agent : core::Runtime::runtime_singleton_->cpu_agents()) {
+      use_default &= (consumer_set.find(cpu_agent->public_handle()) == consumer_set.end());
+    }
+  }
+
+  if (use_default) {
+    ret = new core::DefaultSignal(initial_value, enable_ipc);
+  } else {
+    ret = new core::InterruptSignal(initial_value);
+  }
+
+  *hsa_signal = core::Signal::Convert(ret);
+  return HSA_STATUS_SUCCESS;
+  CATCH;
 }
 
 uint32_t hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* hsa_signals,
                                  hsa_signal_condition_t* conds, hsa_signal_value_t* values,
                                  uint64_t timeout_hint, hsa_wait_state_t wait_hint,
                                  hsa_signal_value_t* satisfying_value) {
-  TRY
+  TRY;
+  IS_OPEN();
   // Do not check for signal invalidation.  Invalidation may occur during async
   // signal handler loop and is not an error.
   for (uint i = 0; i < signal_count; i++)
-    assert(core::SharedSignal::Convert(hsa_signals[i])->IsValid() && "Invalid signal.");
+    assert(hsa_signals[i].handle != 0 && core::SharedSignal::Convert(hsa_signals[i])->IsValid() &&
+           "Invalid signal.");
 
   return core::Signal::WaitAny(signal_count, hsa_signals, conds, values,
                                timeout_hint, wait_hint, satisfying_value);
-  CATCHRET(uint32_t)
+  CATCHRET(uint32_t);
 }
 
 hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t hsa_signal, hsa_signal_condition_t cond,
                                           hsa_signal_value_t value, hsa_amd_signal_handler handler,
                                           void* arg) {
-  TRY
+  TRY;
   IS_OPEN();
+  IS_BAD_PTR(handler);
 
   core::Signal* signal = core::Signal::Convert(hsa_signal);
   IS_VALID(signal);
-  IS_BAD_PTR(handler);
   if (core::g_use_interrupt_wait && (!core::InterruptSignal::IsType(signal)))
     return HSA_STATUS_ERROR_INVALID_SIGNAL;
   return core::Runtime::runtime_singleton_->SetAsyncSignalHandler(
       hsa_signal, cond, value, handler, arg);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_async_function(void (*callback)(void* arg), void* arg) {
-  TRY
+  TRY;
   IS_OPEN();
 
   IS_BAD_PTR(callback);
@@ -400,26 +453,26 @@ hsa_status_t hsa_amd_async_function(void (*callback)(void* arg), void* arg) {
   return core::Runtime::runtime_singleton_->SetAsyncSignalHandler(
       null_signal, HSA_SIGNAL_CONDITION_EQ, 0, (hsa_amd_signal_handler)callback,
       arg);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
                                                uint32_t num_cu_mask_count,
                                                const uint32_t* cu_mask) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(cu_mask);
 
   core::Queue* cmd_queue = core::Queue::Convert(queue);
   IS_VALID(cmd_queue);
   return cmd_queue->SetCUMasking(num_cu_mask_count, cu_mask);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size,
                                  hsa_agent_t* agents, int num_agent,
                                  void** agent_ptr) {
-  TRY
+  TRY;
   IS_OPEN();
   *agent_ptr = NULL;
 
@@ -437,11 +490,11 @@ hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size,
           core::Runtime::runtime_singleton_->system_regions_fine()[0]);
 
   return system_region->Lock(num_agent, agents, host_ptr, size, agent_ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_unlock(void* host_ptr) {
-  TRY
+  TRY;
   IS_OPEN();
 
   const amd::MemoryRegion* system_region =
@@ -449,12 +502,12 @@ hsa_status_t hsa_amd_memory_unlock(void* host_ptr) {
           core::Runtime::runtime_singleton_->system_regions_fine()[0]);
 
   return system_region->Unlock(host_ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
                                           hsa_amd_memory_pool_info_t attribute, void* value) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(value);
 
@@ -465,14 +518,14 @@ hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
   }
 
   return mem_region->GetPoolInfo(attribute, value);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_agent_iterate_memory_pools(
     hsa_agent_t agent_handle,
     hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data),
     void* data) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(callback);
   const core::Agent* agent = core::Agent::Convert(agent_handle);
@@ -490,12 +543,12 @@ hsa_status_t hsa_amd_agent_iterate_memory_pools(
       reinterpret_cast<hsa_status_t (*)(hsa_region_t memory_pool, void* data)>(
           callback),
       data);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size,
                                           uint32_t flags, void** ptr) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (size == 0 || ptr == NULL) {
@@ -511,7 +564,7 @@ hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, siz
 
   return core::Runtime::runtime_singleton_->AllocateMemory(
       mem_region, size, core::MemoryRegion::AllocateRestrict, ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_pool_free(void* ptr) {
@@ -520,7 +573,7 @@ hsa_status_t hsa_amd_memory_pool_free(void* ptr) {
 
 hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents,
                                          const uint32_t* flags, const void* ptr) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (num_agents == 0 || agents == NULL || flags != NULL || ptr == NULL) {
@@ -529,12 +582,12 @@ hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t*
 
   return core::Runtime::runtime_singleton_->AllowAccess(num_agents, agents,
                                                         ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_pool,
                                              hsa_amd_memory_pool_t dst_memory_pool, bool* result) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (result == NULL) {
@@ -558,13 +611,13 @@ hsa_status_t hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_po
   }
 
   return src_mem_region->CanMigrate(*dst_mem_region, *result);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_memory_migrate(const void* ptr,
                                     hsa_amd_memory_pool_t memory_pool,
                                     uint32_t flags) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (ptr == NULL || flags != 0) {
@@ -580,13 +633,13 @@ hsa_status_t hsa_amd_memory_migrate(const void* ptr,
   }
 
   return dst_mem_region->Migrate(flags, ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_agent_memory_pool_get_info(
     hsa_agent_t agent_handle, hsa_amd_memory_pool_t memory_pool,
     hsa_amd_agent_memory_pool_info_t attribute, void* value) {
-  TRY
+  TRY;
   IS_OPEN();
 
   if (value == NULL) {
@@ -605,7 +658,7 @@ hsa_status_t hsa_amd_agent_memory_pool_get_info(
   }
 
   return mem_region->GetAgentPoolInfo(*agent, attribute, value);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents,
@@ -614,7 +667,7 @@ hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents,
                                         void** ptr, size_t* metadata_size,
                                         const void** metadata) {
   static const int tinyArraySize=8;
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(agents);
   IS_BAD_PTR(size);
@@ -641,49 +694,49 @@ hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents,
 
   if (num_agents > tinyArraySize) delete[] core_agents;
   return ret;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_interop_unmap_buffer(void* ptr) {
-  TRY
+  TRY;
   IS_OPEN();
   if (ptr != NULL) core::Runtime::runtime_singleton_->InteropUnmap(ptr);
   return HSA_STATUS_SUCCESS;
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_pointer_info(void* ptr, hsa_amd_pointer_info_t* info, void* (*alloc)(size_t),
                                   uint32_t* num_accessible, hsa_agent_t** accessible) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(ptr);
   IS_BAD_PTR(info);
   return core::Runtime::runtime_singleton_->PtrInfo(ptr, info, alloc, num_accessible, accessible);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_pointer_info_set_userdata(void* ptr, void* userdata) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(ptr);
   return core::Runtime::runtime_singleton_->SetPtrInfoData(ptr, userdata);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_ipc_memory_create(void* ptr, size_t len, hsa_amd_ipc_memory_t* handle) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(ptr);
   IS_BAD_PTR(handle);
   return core::Runtime::runtime_singleton_->IPCCreate(ptr, len, handle);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t len,
                                        uint32_t num_agents, const hsa_agent_t* mapping_agents,
                                        void** mapped_ptr) {
   static const int tinyArraySize = 8;
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(mapped_ptr);
   if (num_agents != 0) IS_BAD_PTR(mapping_agents);
@@ -706,15 +759,38 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t l
 
   return core::Runtime::runtime_singleton_->IPCAttach(ipc, len, num_agents, core_agents,
                                                       mapped_ptr);
-  CATCH
+  CATCH;
 }
 
 hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr) {
-  TRY
+  TRY;
   IS_OPEN();
   IS_BAD_PTR(mapped_ptr);
   return core::Runtime::runtime_singleton_->IPCDetach(mapped_ptr);
-  CATCH
+  CATCH;
+}
+
+hsa_status_t hsa_amd_ipc_signal_create(hsa_signal_t hsa_signal, hsa_amd_ipc_signal_t* handle) {
+  TRY;
+  IS_OPEN();
+  IS_BAD_PTR(handle);
+  core::Signal* signal = core::Signal::Convert(hsa_signal);
+  IS_VALID(signal);
+  core::IPCSignal::CreateHandle(signal, handle);
+  return HSA_STATUS_SUCCESS;
+  CATCH;
+}
+
+hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle,
+                                       hsa_signal_t* hsa_signal) {
+  TRY;
+  IS_OPEN();
+  IS_BAD_PTR(handle);
+  IS_BAD_PTR(hsa_signal);
+  core::Signal* signal = core::IPCSignal::Attach(handle);
+  *hsa_signal = core::Signal::Convert(signal);
+  return HSA_STATUS_SUCCESS;
+  CATCH;
 }
 
 } // end of AMD namespace

runtime/hsa-runtime/core/runtime/interrupt_signal.cpp

@@ -89,10 +89,6 @@ InterruptSignal::InterruptSignal(hsa_signal_value_t initial_value, HsaEvent* use
 }
 
 InterruptSignal::~InterruptSignal() {
-  invalid_ = true;
-  SetEvent();
-  while (InUse())
-    ;
   if (free_event_) hsaKmtDestroyEvent(event_);
 }
 
@@ -119,14 +115,14 @@ void InterruptSignal::StoreRelease(hsa_signal_value_t value) {
 hsa_signal_value_t InterruptSignal::WaitRelaxed(
     hsa_signal_condition_t condition, hsa_signal_value_t compare_value,
     uint64_t timeout, hsa_wait_state_t wait_hint) {
-  uint32_t prior = atomic::Increment(&waiting_);
+  Retain();
+  MAKE_SCOPE_GUARD([&]() { Release(); });
 
-  // assert(prior == 0 && "Multiple waiters on interrupt signal!");
+  uint32_t prior = waiting_++;
+  MAKE_SCOPE_GUARD([&]() { waiting_--; });
   // Allow only the first waiter to sleep (temporary, known to be bad).
   if (prior != 0) wait_hint = HSA_WAIT_STATE_ACTIVE;
 
-  MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); });
-
   int64_t value;
 
   timer::fast_clock::time_point start_time = timer::fast_clock::now();
@@ -143,7 +139,7 @@ hsa_signal_value_t InterruptSignal::WaitRelaxed(
 
   bool condition_met = false;
   while (true) {
-    if (invalid_) return 0;
+    if (!IsValid()) return 0;
 
     value = atomic::Load(&signal_.value, std::memory_order_relaxed);
 

runtime/hsa-runtime/core/runtime/ipc_signal.cpp

@@ -0,0 +1,94 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// The University of Illinois/NCSA
+// Open Source License (NCSA)
+//
+// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
+//
+// Developed by:
+//
+//                 AMD Research and AMD HSA Software Development
+//
+//                 Advanced Micro Devices, Inc.
+//
+//                 www.amd.com
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to
+// deal with 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:
+//
+//  - Redistributions of source code must retain the above copyright notice,
+//    this list of conditions and the following disclaimers.
+//  - Redistributions in binary form must reproduce the above copyright
+//    notice, this list of conditions and the following disclaimers in
+//    the documentation and/or other materials provided with the distribution.
+//  - Neither the names of Advanced Micro Devices, Inc,
+//    nor the names of its contributors may be used to endorse or promote
+//    products derived from this Software without specific prior written
+//    permission.
+//
+// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+//
+////////////////////////////////////////////////////////////////////////////////
+
+#include "core/inc/ipc_signal.h"
+
+#include <utility>
+
+#include "core/inc/runtime.h"
+#include "core/inc/exceptions.h"
+
+namespace core {
+
+int IPCSignal::rtti_id_ = 0;
+KernelMutex IPCSignal::lock_;
+
+SharedMemory::SharedMemory(const hsa_amd_ipc_memory_t* handle, size_t len) {
+  hsa_status_t err = Runtime::runtime_singleton_->IPCAttach(handle, len, 0, NULL, &ptr_);
+  if (err != HSA_STATUS_SUCCESS) throw AMD::hsa_exception(err, "IPC memory attach failed.");
+}
+
+SharedMemory::SharedMemory(SharedMemory&& rhs) {
+  ptr_ = rhs.ptr_;
+  rhs.ptr_ = nullptr;
+}
+
+SharedMemory::~SharedMemory() {
+  if (ptr_ == nullptr) return;
+  auto err = Runtime::runtime_singleton_->IPCDetach(ptr_);
+  assert(err == HSA_STATUS_SUCCESS && "IPC detach failed.");
+}
+
+void IPCSignal::CreateHandle(Signal* signal, hsa_amd_ipc_signal_t* ipc_handle) {
+  if (!signal->isIPC())
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "Signal must be IPC enabled.");
+  SharedSignal* shared = SharedSignal::Convert(Convert(signal));
+  hsa_status_t err = Runtime::runtime_singleton_->IPCCreate(shared, 4096, ipc_handle);
+  if (err != HSA_STATUS_SUCCESS) throw AMD::hsa_exception(err, "IPC memory create failed.");
+}
+
+Signal* IPCSignal::Attach(const hsa_amd_ipc_signal_t* ipc_signal_handle) {
+  SharedMemorySignal shared(ipc_signal_handle);
+
+  if (!(shared.signal()->IsIPC()))
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT,
+                             "IPC memory does not contain an IPC signal abi block.");
+
+  hsa_signal_t handle = SharedSignal::Convert(shared.signal());
+
+  ScopedAcquire<KernelMutex> lock(&lock_);
+  Signal* ret = core::Signal::DuplicateHandle(handle);
+  if (ret == nullptr) ret = new IPCSignal(std::move(shared));
+  return ret;
+}
+
+}  // namespace core

runtime/hsa-runtime/core/runtime/runtime.cpp

@@ -1026,7 +1026,7 @@ void Runtime::Unload() {
 
   async_events_control_.Shutdown();
 
-  delete vm_fault_signal_;
+  vm_fault_signal_->DestroySignal();
   core::InterruptSignal::DestroyEvent(vm_fault_event_);
 
   DestroyAgents();

runtime/hsa-runtime/core/runtime/signal.cpp

@@ -44,24 +44,83 @@
 #define HSA_RUNTME_CORE_SIGNAL_CPP_
 
 #include "core/inc/signal.h"
-#include "core/util/timer.h"
+
 #include <algorithm>
+#include "core/util/timer.h"
 
 namespace core {
 
+KernelMutex Signal::ipcLock_;
+std::map<decltype(hsa_signal_t::handle), Signal*> Signal::ipcMap_;
+
+void Signal::registerIpc() {
+  ScopedAcquire<KernelMutex> lock(&ipcLock_);
+  auto handle = Convert(this);
+  assert(ipcMap_.find(handle.handle) == ipcMap_.end() &&
+         "Can't register the same IPC signal twice.");
+  ipcMap_[handle.handle] = this;
+}
+
+bool Signal::deregisterIpc() {
+  ScopedAcquire<KernelMutex> lock(&ipcLock_);
+  if (refcount_ != 0) return false;
+  auto handle = Convert(this);
+  const auto& it = ipcMap_.find(handle.handle);
+  assert(it != ipcMap_.end() && "Deregister on non-IPC signal.");
+  ipcMap_.erase(it);
+  return true;
+}
+
+Signal* Signal::lookupIpc(hsa_signal_t signal) {
+  ScopedAcquire<KernelMutex> lock(&ipcLock_);
+  const auto& it = ipcMap_.find(signal.handle);
+  if (it == ipcMap_.end()) return nullptr;
+  return it->second;
+}
+
+Signal* Signal::duplicateIpc(hsa_signal_t signal) {
+  ScopedAcquire<KernelMutex> lock(&ipcLock_);
+  const auto& it = ipcMap_.find(signal.handle);
+  if (it == ipcMap_.end()) return nullptr;
+  it->second->refcount_++;
+  it->second->Retain();
+  return it->second;
+}
+
+void Signal::Release() {
+  if (--retained_ != 0) return;
+  if (!isIPC())
+    doDestroySignal();
+  else if (deregisterIpc())
+    doDestroySignal();
+}
+
+Signal::~Signal() {
+  signal_.kind = AMD_SIGNAL_KIND_INVALID;
+  if (refcount_ == 1 && isIPC()) {
+    refcount_ = 0;
+    deregisterIpc();
+  }
+}
+
 uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
                          const hsa_signal_condition_t* conds, const hsa_signal_value_t* values,
                          uint64_t timeout, hsa_wait_state_t wait_hint,
                          hsa_signal_value_t* satisfying_value) {
   hsa_signal_handle* signals =
       reinterpret_cast<hsa_signal_handle*>(const_cast<hsa_signal_t*>(hsa_signals));
-  uint32_t prior = 0;
-  for (uint32_t i = 0; i < signal_count; i++)
-    prior = Max(prior, atomic::Increment(&signals[i]->waiting_));
+
+  for (uint32_t i = 0; i < signal_count; i++) signals[i]->Retain();
 
   MAKE_SCOPE_GUARD([&]() {
-    for (uint32_t i = 0; i < signal_count; i++)
-      atomic::Decrement(&signals[i]->waiting_);
+    for (uint32_t i = 0; i < signal_count; i++) signals[i]->Release();
+  });
+
+  uint32_t prior = 0;
+  for (uint32_t i = 0; i < signal_count; i++) prior = Max(prior, signals[i]->waiting_++);
+
+  MAKE_SCOPE_GUARD([&]() {
+    for (uint32_t i = 0; i < signal_count; i++) signals[i]->waiting_--;
   });
 
   // Allow only the first waiter to sleep (temporary, known to be bad).
@@ -113,7 +172,7 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
   bool condition_met = false;
   while (true) {
     for (uint32_t i = 0; i < signal_count; i++) {
-      if (signals[i]->invalid_) return uint32_t(-1);
+      if (!signals[i]->IsValid()) return uint32_t(-1);
 
       // Handling special event.
       if (signals[i]->EopEvent() != NULL) {

runtime/hsa-runtime/hsacore.so.def

@@ -173,6 +173,7 @@ global:
 	hsa_amd_profiling_async_copy_enable;
 	hsa_amd_profiling_get_async_copy_time;
 	hsa_amd_profiling_convert_tick_to_system_domain;
+	hsa_amd_signal_create;
 	hsa_amd_signal_wait_any;
 	hsa_amd_signal_async_handler;
 	hsa_amd_async_function;
@@ -203,14 +204,16 @@ global:
 	hsa_ext_image_destroy;
 	hsa_ext_sampler_create;
 	hsa_ext_sampler_destroy;
-        hsa_ext_image_get_capability_with_layout;
-        hsa_ext_image_data_get_info_with_layout;
-        hsa_ext_image_create_with_layout;
+    hsa_ext_image_get_capability_with_layout;
+    hsa_ext_image_data_get_info_with_layout;
+    hsa_ext_image_create_with_layout;
 	hsa_amd_pointer_info;
 	hsa_amd_pointer_info_set_userdata;
 	hsa_amd_ipc_memory_create;
 	hsa_amd_ipc_memory_attach;
 	hsa_amd_ipc_memory_detach;
+	hsa_amd_ipc_signal_create;
+	hsa_amd_ipc_signal_attach;
 
 local:
     *;

runtime/hsa-runtime/inc/hsa_api_trace.h

@@ -166,6 +166,9 @@ struct AmdExtTable {
   decltype(hsa_amd_ipc_memory_create)* hsa_amd_ipc_memory_create_fn;
   decltype(hsa_amd_ipc_memory_attach)* hsa_amd_ipc_memory_attach_fn;
   decltype(hsa_amd_ipc_memory_detach)* hsa_amd_ipc_memory_detach_fn;
+  decltype(hsa_amd_signal_create)* hsa_amd_signal_create_fn;
+  decltype(hsa_amd_ipc_signal_create)* hsa_amd_ipc_signal_create_fn;
+  decltype(hsa_amd_ipc_signal_attach)* hsa_amd_ipc_signal_attach_fn;
 };
 
 // Table to export HSA Core Runtime Apis

runtime/hsa-runtime/inc/hsa_ext_amd.h

@@ -377,6 +377,62 @@ hsa_status_t HSA_API
                                                     uint64_t agent_tick,
                                                     uint64_t* system_tick);
 
+/**
+ * @brief Signal attribute flags.
+ */
+typedef enum {
+  /**
+   * Signal will only be consumed by AMD GPUs.  Limits signal consumption to
+   * AMD GPU agents only.  Ignored if @p num_consumers is not zero (all agents).
+   */
+  HSA_AMD_SIGNAL_AMD_GPU_ONLY = 1,
+  /**
+   * Signal may be used for interprocess communication.
+   * This signal may not be used with profiling APIs.  Errors or inaccurate
+   * timing data may result from such use.
+   */
+  HSA_AMD_SIGNAL_IPC = 2,
+} hsa_amd_signal_attribute_t;
+
+/**
+ * @brief Create a signal with specific attributes.
+ *
+ * @param[in] initial_value Initial value of the signal.
+ *
+ * @param[in] num_consumers Size of @p consumers. A value of 0 indicates that
+ * any agent might wait on the signal.
+ *
+ * @param[in] consumers List of agents that might consume (wait on) the
+ * signal. If @p num_consumers is 0, this argument is ignored; otherwise, the
+ * HSA runtime might use the list to optimize the handling of the signal
+ * object. If an agent not listed in @p consumers waits on the returned
+ * signal, the behavior is undefined. The memory associated with @p consumers
+ * can be reused or freed after the function returns.
+ *
+ * @param[in] attributes Requested signal attributes.  Multiple signal attributes
+ * may be requested by combining them with bitwise OR.  Requesting no attributes
+ * (@p attributes == 0) results in the same signal as would have been obtained
+ * via hsa_signal_create.
+ *
+ * @param[out] signal Pointer to a memory location where the HSA runtime will
+ * store the newly created signal handle. Must not be NULL.
+ *
+ * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
+ *
+ * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
+ * initialized.
+ *
+ * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
+ * the required resources.
+ *
+ * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p signal is NULL, @p
+ * num_consumers is greater than 0 but @p consumers is NULL, or @p consumers
+ * contains duplicates.
+ */
+hsa_status_t HSA_API hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
+                                           const hsa_agent_t* consumers, uint64_t attributes,
+                                           hsa_signal_t* signal);
+
 /**
  * @brief Asyncronous signal handler function type.
  *
@@ -1511,6 +1567,60 @@ hsa_status_t HSA_API hsa_amd_ipc_memory_attach(
  */
 hsa_status_t HSA_API hsa_amd_ipc_memory_detach(void* mapped_ptr);
 
+/**
+ * @brief 256-bit process independent identifier for a ROCr IPC signal.
+ */
+typedef hsa_amd_ipc_memory_t hsa_amd_ipc_signal_t;
+
+/**
+ * @brief Obtains an interprocess sharing handle for a signal.  The handle is
+ * valid while the signal it references remains valid in any process.  In
+ * general applications should confirm that the signal has been attached (via
+ * hsa_amd_ipc_signal_attach) in the remote process prior to destroying that
+ * signal in the local process.
+ * Repeated calls for the same signal may, but are not required to, return
+ * unique handles.
+ *
+ * @param[in] signal Signal created with attribute HSA_AMD_SIGNAL_IPC.
+ *
+ * @param[out] handle Process independent identifier referencing the shared
+ * signal.
+ *
+ * @retval HSA_STATUS_SUCCESS @p handle is ready to use for interprocess sharing.
+ *
+ * @retval HSA_STATUS_ERROR_NOT_INITIALIZED if HSA is not initialized
+ *
+ * @retval HSA_STATUS_ERROR_OUT_OF_RESOURCES if there is a failure in allocating
+ * necessary resources
+ *
+ * @retval HSA_STATUS_ERROR_INVALID_ARGUMENT @p signal is not a valid signal
+ * created with attribute HSA_AMD_SIGNAL_IPC or handle is NULL.
+ */
+hsa_status_t HSA_API hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle);
+
+/**
+ * @brief Imports an IPC capable signal into the local process.  If an IPC
+ * signal handle is attached multiple times in a process each attach may return
+ * a different signal handle.  Each returned signal handle is refcounted and
+ * requires a matching number of calls to hsa_signal_destroy to release the
+ * shared signal.
+ *
+ * @param[in] handle Pointer to the identifier for the shared signal.
+ *
+ * @param[out] signal Recieves a process local signal handle to the shared signal.
+ *
+ * @retval HSA_STATUS_SUCCESS if the signal is successfully imported.
+ *
+ * @retval HSA_STATUS_ERROR_NOT_INITIALIZED if HSA is not initialized
+ *
+ * @retval HSA_STATUS_ERROR_OUT_OF_RESOURCES if there is a failure in allocating
+ * necessary resources
+ *
+ * @retval HSA_STATUS_ERROR_INVALID_ARGUMENT @p handle is not valid.
+ */
+hsa_status_t HSA_API hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle,
+                                               hsa_signal_t* signal);
+
 #ifdef __cplusplus
 }  // end extern "C" block
 #endif