Add queue intercept support to the runtime.

Queue intercept is exposed as two tools-only APIs via the API
intercept table.

Change-Id: Iac9602ed3143974d85c3569e9092295ad18037f8

runtime/hsa-runtime/CMakeLists.txt

@@ -135,6 +135,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/intercept_queue.cpp"
             "core/runtime/ipc_signal.cpp"
             "core/runtime/isa.cpp"
             "core/runtime/runtime.cpp"

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

@@ -1102,3 +1102,19 @@ hsa_status_t HSA_API hsa_amd_register_system_event_handler(
     void* data) {
   return amdExtTable->hsa_amd_register_system_event_handler_fn(type, callback, data);
 }
+
+// Mirrors Amd Extension Apis
+hsa_status_t hsa_amd_queue_intercept_create(
+    hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type,
+    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
+    uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) {
+  return amdExtTable->hsa_amd_queue_intercept_create_fn(
+      agent_handle, size, type, callback, data, private_segment_size, group_segment_size, queue);
+}
+
+// Mirrors Amd Extension Apis
+hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue,
+                                              hsa_amd_queue_intercept_handler callback,
+                                              void* user_data) {
+  return amdExtTable->hsa_amd_queue_intercept_register_fn(queue, callback, user_data);
+}

runtime/hsa-runtime/core/common/shared.h

@@ -43,11 +43,12 @@
 #ifndef HSA_RUNTME_CORE_INC_SHARED_H_
 #define HSA_RUNTME_CORE_INC_SHARED_H_
 
-#include "core/util/utils.h"
 #include <assert.h>
-
 #include <cstring>
 #include <functional>
+#include <memory>
+
+#include "core/util/utils.h"
 
 namespace core {
 /// @brief Base class encapsulating the allocator and deallocator for

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

@@ -59,6 +59,8 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
     return signal->IsType(&rtti_id_);
   }
 
+  static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
+
   // Acquires/releases queue resources and requests HW schedule/deschedule.
   AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id,
            ScratchInfo& scratch, core::HsaEventCallback callback,
@@ -337,7 +339,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
   }
 
  protected:
-  bool _IsA(rtti_t id) const override { return id == &rtti_id_; }
+  bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; }
 
   /// @brief Disallow destroying doorbell apart from its queue.
   void doDestroySignal() override { assert(false); }

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

@@ -51,6 +51,8 @@
 namespace core {
 class HostQueue : public Queue {
  public:
+  static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
+
   HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type,
             uint32_t features, hsa_signal_t doorbell_signal);
 
@@ -158,7 +160,11 @@ class HostQueue : public Queue {
 
   void operator delete(void*, void*) {}
 
+ protected:
+  bool _IsA(Queue::rtti_t id) const { return id == &rtti_id_; }
+
  private:
+  static int rtti_id_;
   static const size_t kRingAlignment = 256;
   const uint32_t size_;
   bool active_;

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

@@ -0,0 +1,413 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// 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_INTERCEPT_QUEUE_H_
+#define HSA_RUNTIME_CORE_INC_INTERCEPT_QUEUE_H_
+
+#include <vector>
+#include <memory>
+#include <utility>
+
+#include "core/inc/runtime.h"
+#include "core/inc/queue.h"
+#include "core/inc/signal.h"
+#include "core/inc/interrupt_signal.h"
+#include "core/inc/exceptions.h"
+#include "core/util/locks.h"
+
+namespace core {
+
+// @brief Generic container to forward Queue interfaces into Queue* member.
+// Class only has utility as a base type customized Queue wrappers.
+class QueueWrapper : public Queue {
+ public:
+  std::unique_ptr<Queue> wrapped;
+
+  explicit QueueWrapper(std::unique_ptr<Queue> queue) : Queue(), wrapped(std::move(queue)) {
+    if (Queue::Shared::shared_object() == NULL) {
+      return;
+    }
+
+    memcpy(&amd_queue_, &wrapped->amd_queue_, sizeof(amd_queue_t));
+    wrapped->set_public_handle(wrapped.get(), public_handle_);
+  }
+
+  ~QueueWrapper() {}
+
+  hsa_status_t Inactivate() override { return wrapped->Inactivate(); }
+  uint64_t LoadReadIndexAcquire() override { return wrapped->LoadReadIndexAcquire(); }
+  uint64_t LoadReadIndexRelaxed() override { return wrapped->LoadReadIndexRelaxed(); }
+  uint64_t LoadWriteIndexRelaxed() override { return wrapped->LoadWriteIndexRelaxed(); }
+  uint64_t LoadWriteIndexAcquire() override { return wrapped->LoadWriteIndexAcquire(); }
+  void StoreReadIndexRelaxed(uint64_t value) override {
+    return wrapped->StoreReadIndexRelaxed(value);
+  }
+  void StoreReadIndexRelease(uint64_t value) override {
+    return wrapped->StoreReadIndexRelease(value);
+  }
+  void StoreWriteIndexRelaxed(uint64_t value) override {
+    return wrapped->StoreWriteIndexRelaxed(value);
+  }
+  void StoreWriteIndexRelease(uint64_t value) override {
+    return wrapped->StoreWriteIndexRelease(value);
+  }
+  uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) override {
+    return wrapped->CasWriteIndexAcqRel(expected, value);
+  }
+  uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) override {
+    return wrapped->CasWriteIndexAcquire(expected, value);
+  }
+  uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) override {
+    return wrapped->CasWriteIndexRelaxed(expected, value);
+  }
+  uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) override {
+    return wrapped->CasWriteIndexRelease(expected, value);
+  }
+  uint64_t AddWriteIndexAcqRel(uint64_t value) override {
+    return wrapped->AddWriteIndexAcqRel(value);
+  }
+  uint64_t AddWriteIndexAcquire(uint64_t value) override {
+    return wrapped->AddWriteIndexAcquire(value);
+  }
+  uint64_t AddWriteIndexRelaxed(uint64_t value) override {
+    return wrapped->AddWriteIndexRelaxed(value);
+  }
+  uint64_t AddWriteIndexRelease(uint64_t value) override {
+    return wrapped->AddWriteIndexRelease(value);
+  }
+  hsa_status_t SetCUMasking(const uint32_t num_cu_mask_count, const uint32_t* cu_mask) override {
+    return wrapped->SetCUMasking(num_cu_mask_count, cu_mask);
+  }
+  void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) override {
+    wrapped->ExecutePM4(cmd_data, cmd_size_b);
+  }
+
+ protected:
+  void do_set_public_handle(hsa_queue_t* handle) {
+    public_handle_ = handle;
+    wrapped->set_public_handle(wrapped.get(), handle);
+  }
+};
+
+// @brief Generic container for a proxy queue.
+// Presents an proxy packet buffer and doorbell signal for an underlying Queue.  Write index
+// operations act on the proxy buffer while all other operations pass through to the underlying
+// queue.
+class QueueProxy : public QueueWrapper {
+ public:
+  explicit QueueProxy(std::unique_ptr<Queue> queue) : QueueWrapper(std::move(queue)) {}
+
+  uint64_t LoadReadIndexAcquire() override {
+    return atomic::Load(&amd_queue_.read_dispatch_id, std::memory_order_acquire);
+  }
+  uint64_t LoadReadIndexRelaxed() override {
+    return atomic::Load(&amd_queue_.read_dispatch_id, std::memory_order_relaxed);
+  }
+  void StoreReadIndexRelaxed(uint64_t value) override { assert(false); }
+  void StoreReadIndexRelease(uint64_t value) override { assert(false); }
+
+  uint64_t LoadWriteIndexRelaxed() override {
+    return atomic::Load(&amd_queue_.write_dispatch_id, std::memory_order_relaxed);
+  }
+  uint64_t LoadWriteIndexAcquire() override {
+    return atomic::Load(&amd_queue_.write_dispatch_id, std::memory_order_acquire);
+  }
+  void StoreWriteIndexRelaxed(uint64_t value) override {
+    atomic::Store(&amd_queue_.write_dispatch_id, value, std::memory_order_relaxed);
+  }
+  void StoreWriteIndexRelease(uint64_t value) override {
+    atomic::Store(&amd_queue_.write_dispatch_id, value, std::memory_order_release);
+  }
+  uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) override {
+    return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_acq_rel);
+  }
+  uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) override {
+    return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_acquire);
+  }
+  uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) override {
+    return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_relaxed);
+  }
+  uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) override {
+    return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_release);
+  }
+  uint64_t AddWriteIndexAcqRel(uint64_t value) override {
+    return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_acq_rel);
+  }
+  uint64_t AddWriteIndexAcquire(uint64_t value) override {
+    return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_acquire);
+  }
+  uint64_t AddWriteIndexRelaxed(uint64_t value) override {
+    return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_relaxed);
+  }
+  uint64_t AddWriteIndexRelease(uint64_t value) override {
+    return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_release);
+  }
+};
+
+// @brief Provides packet intercept and rewrite capability for a queue.
+// Host-side dispatches are processed during doorbell ring.
+// Device-side dispatches are processed as an asynchronous signal event.
+class InterceptQueue : public QueueProxy, private LocalSignal, public Signal {
+ public:
+  // typedef void(*intercept_packet_writer)(const AqlPacket* pkts, uint64_t pkt_count);
+  // typedef void(*intercept_handler)(const AqlPacket* pkts, uint64_t pkt_count, uint64_t
+  // user_pkt_index, void* data, intercept_packet_writer writer);
+
+  std::vector<std::pair<hsa_amd_queue_intercept_handler, void*>> interceptors;
+
+  explicit InterceptQueue(std::unique_ptr<Queue> queue);
+  ~InterceptQueue();
+
+  void AddInterceptor(hsa_amd_queue_intercept_handler interceptor, void* data) {
+    interceptors.push_back(std::make_pair(interceptor, data));
+  }
+
+  hsa_status_t Inactivate() override {
+    active_ = false;
+    return wrapped->Inactivate();
+  }
+
+ private:
+  // Serialize packet interception processing.
+  KernelMutex lock_;
+
+  // Largest processed packet index.
+  uint64_t next_packet_;
+
+  // Post interception packet overflow buffer
+  std::vector<AqlPacket> overflow_;
+
+  // Index at which async intercept processing was scheduled.
+  uint64_t retry_index_;
+
+  // Event signal to use for async packet processing and control flag.
+  InterruptSignal* async_doorbell_;
+  std::atomic<bool> quit_;
+
+  // Indicates queue active/inactive state.
+  std::atomic<bool> active_;
+
+  static const hsa_signal_value_t DOORBELL_MAX = 0xFFFFFFFFFFFFFFFFull;
+
+  static bool HandleAsyncDoorbell(hsa_signal_value_t value, void* arg);
+  static void PacketWriter(const void* pkts, uint64_t pkt_count);
+
+  bool Submit(const AqlPacket* packets, uint64_t count);
+  static void Submit(const void* pkts, uint64_t pkt_count, uint64_t user_pkt_index, void* data,
+                     hsa_amd_queue_intercept_packet_writer writer);
+
+  /*
+   * Remaining Queue and Signal interface definitions.
+   */
+ public:
+  /// @brief Update signal value using Relaxed semantics
+  ///
+  /// @param value Value of signal to update with
+  void StoreRelaxed(hsa_signal_value_t value);
+
+  /// @brief Update signal value using Release semantics
+  ///
+  /// @param value Value of signal to update with
+  void StoreRelease(hsa_signal_value_t value) {
+    std::atomic_thread_fence(std::memory_order_release);
+    StoreRelaxed(value);
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t LoadRelaxed() {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t LoadAcquire() {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition, hsa_signal_value_t compare_value,
+                                 uint64_t timeout, hsa_wait_state_t wait_hint) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition, hsa_signal_value_t compare_value,
+                                 uint64_t timeout, hsa_wait_state_t wait_hint) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  void AndRelaxed(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AndAcquire(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AndRelease(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AndAcqRel(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void OrRelaxed(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void OrAcquire(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void OrRelease(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void OrAcqRel(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void XorRelaxed(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void XorAcquire(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void XorRelease(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void XorAcqRel(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AddRelaxed(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AddAcquire(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AddRelease(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void AddAcqRel(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void SubRelaxed(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void SubAcquire(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void SubRelease(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  void SubAcqRel(hsa_signal_value_t value) { assert(false); }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t ExchAcquire(hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t ExchRelease(hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected, hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t CasAcquire(hsa_signal_value_t expected, hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t CasRelease(hsa_signal_value_t expected, hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected, hsa_signal_value_t value) {
+    assert(false);
+    return 0;
+  }
+
+  /// @brief This operation is illegal
+  hsa_signal_value_t* ValueLocation() const {
+    assert(false);
+    return NULL;
+  }
+
+  /// @brief This operation is illegal
+  HsaEvent* EopEvent() {
+    assert(false);
+    return NULL;
+  }
+
+  static __forceinline bool IsType(core::Signal* signal) { return signal->IsType(&rtti_id_); }
+  static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
+
+ protected:
+  bool _IsA(Queue::rtti_t id) const { return id == &rtti_id_; }
+
+ private:
+  static int rtti_id_;
+};
+
+}  // namespace core
+
+#endif  // HSA_RUNTIME_CORE_INC_INTERCEPT_QUEUE_H_

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

@@ -306,6 +306,10 @@ class Queue : public Checked<0xFA3906A679F9DB49>,
 
   hsa_queue_t* public_handle() const { return public_handle_; }
 
+  typedef void* rtti_t;
+
+  bool IsType(rtti_t id) { return _IsA(id); }
+
  protected:
   static void set_public_handle(Queue* ptr, hsa_queue_t* handle) {
     ptr->do_set_public_handle(handle);
@@ -315,6 +319,8 @@ class Queue : public Checked<0xFA3906A679F9DB49>,
   }
   hsa_queue_t* public_handle_;
 
+  virtual bool _IsA(rtti_t id) const = 0;
+
  private:
   DISALLOW_COPY_AND_ASSIGN(Queue);
 };

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

@@ -50,6 +50,7 @@
 
 #include "core/inc/hsa_ext_interface.h"
 #include "core/inc/hsa_internal.h"
+#include "core/inc/hsa_ext_amd_impl.h"
 
 #include "core/inc/agent.h"
 #include "core/inc/memory_region.h"

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

@@ -964,10 +964,8 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {
 // Attempt to trim the maximum number of concurrent waves to allow scratch to fit.
 // This is somewhat dangerous as it limits the number of concurrent waves from future dispatches
 // on the queue if those waves use even small amounts of scratch.
-#ifndef NDEBUG
   if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
-    fprintf(stderr, "Failed to map requested scratch - reducing queue occupancy.\n");
-#endif
+    debug_print("Failed to map requested scratch - reducing queue occupancy.\n");
   uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU;
   uint64_t size_per_wave = AlignUp(scratch.size_per_thread * properties_.WaveFrontSize, 1024);
   uint64_t total_waves = scratch.size / size_per_wave;
@@ -994,10 +992,8 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {
 
   // Failed to allocate minimal scratch
   assert(scratch.queue_base == NULL && "bad scratch data");
-#ifndef NDEBUG
   if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
-    fprintf(stderr, "Could not allocate scratch for one wave per CU.\n");
-#endif
+    debug_print("Could not allocate scratch for one wave per CU.\n");
 }
 
 void GpuAgent::ReleaseQueueScratch(void* base) {

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

@@ -47,6 +47,7 @@
 
 namespace core {
 
+int HostQueue::rtti_id_ = 0;
 std::atomic<uint32_t> HostQueue::queue_count_(0x80000000);
 
 HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size,

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

@@ -47,6 +47,17 @@
 
 #include <iostream>
 
+// Tools only APIs.
+namespace AMD {
+hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue,
+                                              hsa_amd_queue_intercept_handler callback,
+                                              void* user_data);
+hsa_status_t hsa_amd_queue_intercept_create(
+    hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type,
+    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
+    uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue);
+}
+
 namespace core {
 
 HsaApiTable hsa_api_table_;
@@ -382,6 +393,8 @@ void HsaApiTable::UpdateAmdExts() {
   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;
   amd_ext_api.hsa_amd_register_system_event_handler_fn = AMD::hsa_amd_register_system_event_handler;
+  amd_ext_api.hsa_amd_queue_intercept_create_fn = AMD::hsa_amd_queue_intercept_create;
+  amd_ext_api.hsa_amd_queue_intercept_register_fn = AMD::hsa_amd_queue_intercept_register;
 }
 
 class Init {

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

@@ -44,8 +44,8 @@
 #include <typeinfo>
 #include <exception>
 #include <set>
-
-#include "hsakmt.h"
+#include <utility>
+#include <memory>
 
 #include "core/inc/runtime.h"
 #include "core/inc/agent.h"
@@ -56,6 +56,7 @@
 #include "core/inc/default_signal.h"
 #include "core/inc/interrupt_signal.h"
 #include "core/inc/ipc_signal.h"
+#include "core/inc/intercept_queue.h"
 #include "core/inc/exceptions.h"
 
 template <class T>
@@ -134,12 +135,10 @@ hsa_status_t handleException() {
     // Rethrow exceptions caught from callbacks.
     //  e.rethrow_nested();
     //  return HSA_STATUS_ERROR;
-#ifndef NDEBUG
   } catch (const std::exception& e) {
-    fprintf(stderr, "Unhandled exception: %s\n", e.what());
+    debug_print("Unhandled exception: %s\n", e.what());
     assert(false && "Unhandled exception.");
     return HSA_STATUS_ERROR;
-#endif
   } catch (...) {
     assert(false && "Unhandled exception.");
     abort();
@@ -798,6 +797,45 @@ hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle,
   CATCH;
 }
 
+// For use by tools only - not in library export table.
+hsa_status_t hsa_amd_queue_intercept_create(
+    hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type,
+    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
+    uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) {
+  TRY;
+  hsa_queue_t* lower_queue;
+  hsa_status_t err = HSA::hsa_queue_create(agent_handle, size, type, callback, data,
+                                           private_segment_size, group_segment_size, &lower_queue);
+  if (err != HSA_STATUS_SUCCESS) return err;
+  std::unique_ptr<core::Queue> lowerQueue(core::Queue::Convert(lower_queue));
+
+  std::unique_ptr<core::InterceptQueue> upperQueue(new core::InterceptQueue(std::move(lowerQueue)));
+
+  if (!upperQueue->IsSharedObjectAllocationValid()) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+
+  *queue = core::Queue::Convert(upperQueue.get());
+  upperQueue.release();
+
+  return HSA_STATUS_SUCCESS;
+  CATCH;
+}
+
+// For use by tools only - not in library export table.
+hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue,
+                                              hsa_amd_queue_intercept_handler callback,
+                                              void* user_data) {
+  TRY;
+  IS_OPEN();
+  IS_BAD_PTR(callback);
+  core::Queue* cmd_queue = core::Queue::Convert(queue);
+  IS_VALID(cmd_queue);
+  if (!core::InterceptQueue::IsType(cmd_queue)) return HSA_STATUS_ERROR_INVALID_QUEUE;
+  core::InterceptQueue* iQueue = static_cast<core::InterceptQueue*>(cmd_queue);
+  iQueue->AddInterceptor(callback, user_data);
+  return HSA_STATUS_SUCCESS;
+  CATCH;
+}
+
 hsa_status_t hsa_amd_register_system_event_handler(
     hsa_amd_event_t type,
     hsa_status_t (*callback)(const void* event_specific_data, void* data),

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

@@ -0,0 +1,260 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// 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/intercept_queue.h"
+#include "core/util/utils.h"
+
+namespace core {
+
+struct InterceptFrame {
+  InterceptQueue* queue;
+  uint64_t pkt_index;
+  size_t interceptor_index;
+};
+
+static thread_local InterceptFrame Cursor = {nullptr, 0, 0};
+
+static const uint16_t kInvalidHeader = (HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE) |
+    (1 << HSA_PACKET_HEADER_BARRIER) |
+    (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+    (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+
+static const uint16_t kBarrierHeader = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) |
+    (1 << HSA_PACKET_HEADER_BARRIER) |
+    (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+    (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+
+static const hsa_barrier_and_packet_t kBarrierPacket = {kInvalidHeader, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+
+int InterceptQueue::rtti_id_ = 0;
+
+InterceptQueue::InterceptQueue(std::unique_ptr<Queue> queue)
+    : QueueProxy(std::move(queue)),
+      LocalSignal(0),
+      Signal(signal()),
+      next_packet_(0),
+      retry_index_(0),
+      quit_(false),
+      active_(true) {
+  if (Queue::Shared::shared_object() == NULL) {
+    return;
+    // TODO skeely: cleanup queue constructor failure.
+    // throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Failed to allocate Queue ABI
+    // block.\n");
+  }
+
+  void* buffer = Runtime::runtime_singleton_->system_allocator()(
+      wrapped->amd_queue_.hsa_queue.size * sizeof(hsa_agent_dispatch_packet_t), 4096,
+      MemoryRegion::AllocateNoFlags);
+  if (buffer == nullptr) throw std::bad_alloc();
+  MAKE_NAMED_SCOPE_GUARD(buffGuard,
+                         [&]() { Runtime::runtime_singleton_->system_deallocator()(buffer); });
+  amd_queue_.hsa_queue.base_address = buffer;
+
+  // Match the queue's signal ABI block to async_doorbell_'s
+  // This allows us to use the queue's signal ABI block from devices to trigger async_doorbell while
+  // host side use jumps directly to the queue's signal implementation.
+  async_doorbell_ = new InterruptSignal(DOORBELL_MAX);
+  MAKE_NAMED_SCOPE_GUARD(sigGuard, [&]() { async_doorbell_->DestroySignal(); });
+  this->signal_ = async_doorbell_->signal_;
+  amd_queue_.hsa_queue.doorbell_signal = Signal::Convert(this);
+
+  // Install an async handler for device side dispatches.
+  auto err = Runtime::runtime_singleton_->SetAsyncSignalHandler(
+      core::Signal::Convert(async_doorbell_), HSA_SIGNAL_CONDITION_NE,
+      async_doorbell_->LoadRelaxed(), HandleAsyncDoorbell, this);
+  if (err != HSA_STATUS_SUCCESS)
+    throw AMD::hsa_exception(err, "Doorbell handler registration failed.\n");
+
+  // Install copy submission interceptor.
+  AddInterceptor(Submit, this);
+
+  sigGuard.Dismiss();
+  buffGuard.Dismiss();
+}
+
+InterceptQueue::~InterceptQueue() {
+  if (Queue::Shared::shared_object() == NULL) {
+    return;
+  }
+
+  active_ = false;
+
+  // Kill the async doorbell handler
+  // Doorbell may not be used during or after queue destroy, however an interrupt may be in flight.
+  // Ensure doorbell value is not 0, mark for exit, wake handler and wait for termination value.
+  async_doorbell_->StoreRelaxed(DOORBELL_MAX);
+  quit_ = true;
+  hsa_signal_value_t val = async_doorbell_->ExchRelaxed(1);
+  if (val != 0)
+    async_doorbell_->WaitRelaxed(HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_BLOCKED);
+  async_doorbell_->DestroySignal();
+
+  // Release buffer resources
+  Runtime::runtime_singleton_->system_deallocator()(amd_queue_.hsa_queue.base_address);
+}
+
+bool InterceptQueue::HandleAsyncDoorbell(hsa_signal_value_t value, void* arg) {
+  InterceptQueue* queue = reinterpret_cast<InterceptQueue*>(arg);
+  if (queue->quit_) {
+    queue->async_doorbell_->StoreRelaxed(0);
+    return false;
+  }
+  queue->async_doorbell_->StoreRelaxed(DOORBELL_MAX);
+  queue->StoreRelease(value);
+  return true;
+}
+
+void InterceptQueue::PacketWriter(const void* pkts, uint64_t pkt_count) {
+  Cursor.interceptor_index--;
+  auto& handler = Cursor.queue->interceptors[Cursor.interceptor_index];
+  handler.first(pkts, pkt_count, Cursor.pkt_index, handler.second, PacketWriter);
+}
+
+void InterceptQueue::Submit(const void* pkts, uint64_t pkt_count, uint64_t user_pkt_index,
+                            void* data, hsa_amd_queue_intercept_packet_writer writer) {
+  InterceptQueue* queue = reinterpret_cast<InterceptQueue*>(data);
+  const AqlPacket* packets = (const AqlPacket*)pkts;
+
+  // Submit final packet transform to hardware.
+  if (queue->Submit(packets, pkt_count)) return;
+
+  // Could not submit final packets, stash for later.
+  assert(queue->overflow_.empty() && "Packet intercept error: overflow buffer not empty.\n");
+  for (uint64_t i = 0; i < pkt_count; i++) queue->overflow_.push_back(packets[i]);
+}
+
+bool InterceptQueue::Submit(const AqlPacket* packets, uint64_t count) {
+  if (count == 0) return true;
+
+  AqlPacket* ring = reinterpret_cast<AqlPacket*>(wrapped->amd_queue_.hsa_queue.base_address);
+  uint64_t mask = wrapped->amd_queue_.hsa_queue.size - 1;
+
+  while (true) {
+    uint64_t write = wrapped->LoadWriteIndexRelaxed();
+    uint64_t read = wrapped->LoadReadIndexRelaxed();
+    uint64_t free_slots = wrapped->amd_queue_.hsa_queue.size - (write - read);
+
+    // If out of space defer packet insertion.
+    if (free_slots <= count) {
+      // If there is not already a pending retry point add one.
+      if (retry_index_ <= read) {
+        // Reserve and wait for one slot.
+        write = wrapped->AddWriteIndexRelaxed(1);
+        read = write - wrapped->amd_queue_.hsa_queue.size + 1;
+        while (wrapped->LoadReadIndexRelaxed() < read) os::YieldThread();
+
+        // Submit barrer which will wake async queue processing.
+        ring[write & mask].barrier_and = kBarrierPacket;
+        ring[write & mask].barrier_and.completion_signal = Signal::Convert(async_doorbell_);
+        atomic::Store(&ring[write & mask].barrier_and.header, kBarrierHeader,
+                      std::memory_order_release);
+        HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal, write);
+
+        // Record the retry point
+        retry_index_ = write;
+      }
+      return false;
+    }
+
+    // Attempt to reserve useable queue space
+    uint64_t new_write = wrapped->CasWriteIndexRelaxed(write, write + count);
+    if (new_write == write) {
+      AqlPacket first = packets[0];
+      uint16_t header = first.dispatch.header;
+      first.dispatch.header = kInvalidHeader;
+
+      ring[write & mask] = first;
+      for (uint64_t i = 1; i < count; i++) ring[(write + i) & mask] = packets[i];
+      atomic::Store(&ring[write & mask].dispatch.header, header, std::memory_order_release);
+      HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal,
+                                      write + count - 1);
+
+      return true;
+    }
+  }
+}
+
+void InterceptQueue::StoreRelaxed(hsa_signal_value_t value) {
+  if (!active_) return;
+
+  // If called recursively defer to async doorbell thread.
+  if (Cursor.queue != nullptr) {
+    debug_print("Likely incorrect queue use observed in an interceptor.\n");
+    async_doorbell_->StoreRelaxed(value);
+    return;
+  }
+
+  ScopedAcquire<KernelMutex> lock(&lock_);
+
+  // Submit overflow packets.
+  if (!overflow_.empty()) {
+    if (!Submit(&overflow_[0], overflow_.size())) return;
+    overflow_.clear();
+  }
+
+  Cursor.queue = this;
+
+  AqlPacket* ring = reinterpret_cast<AqlPacket*>(amd_queue_.hsa_queue.base_address);
+  uint64_t mask = wrapped->amd_queue_.hsa_queue.size - 1;
+
+  // Loop over valid packets and process.
+  uint64_t end = LoadWriteIndexAcquire();
+  uint64_t i;
+  for (i = next_packet_; i < end; i++) {
+    if (!ring[i & mask].IsValid()) break;
+
+    // Process callbacks.
+    Cursor.interceptor_index = interceptors.size() - 1;
+    Cursor.pkt_index = i;
+    auto& handler = interceptors[Cursor.interceptor_index];
+    handler.first(&ring[i & mask], 1, i, handler.second, PacketWriter);
+
+    // Invalidate consumed packet
+    atomic::Store(&ring[i & mask].dispatch.header, kInvalidHeader, std::memory_order_release);
+  }
+
+  next_packet_ = i;
+  Cursor.queue = nullptr;
+}
+
+}  // namespace core

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

@@ -1320,12 +1320,10 @@ void Runtime::LoadTools() {
         add = (tool_add_t)os::GetExportAddress(tool, "AddAgent");
         if (add) add(this);
       }
-#ifndef NDEBUG
       else {
         if (flag().report_tool_load_failures())
-          fprintf(stderr, "Tool lib \"%s\" failed to load.\n", names[i].c_str());
+          debug_print("Tool lib \"%s\" failed to load.\n", names[i].c_str());
       }
-#endif
     }
   }
 }

runtime/hsa-runtime/core/util/utils.h

@@ -78,12 +78,15 @@ static __forceinline void _aligned_free(void* ptr) { return free(ptr); }
 #elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64))
 #include "intrin.h"
 #define __ALIGNED__(x) __declspec(align(x))
-#if (_MSC_VER < 1800)
+#if (_MSC_VER < 1800)  // < VS 2013
 static __forceinline unsigned long long int strtoull(const char* str,
                                                      char** endptr, int base) {
   return static_cast<unsigned long long>(_strtoui64(str, endptr, base));
 }
 #endif
+#if (_MSC_VER < 1900)  // < VS 2015
+#define thread_local __declspec(thread)
+#endif
 #else
 #error "Compiler and/or processor not identified."
 #endif
@@ -102,7 +105,18 @@ static __forceinline unsigned long long int strtoull(const char* str,
     if (!(exp))                                                                                    \
       fprintf(stderr, "Warning: " STRING(exp) " in %s, " __FILE__ ":" STRING(__LINE__) "\n",       \
               __PRETTY_FUNCTION__);                                                                \
-  } while (false);
+  } while (false)
+#endif
+
+#ifdef NDEBUG
+#define debug_print(fmt, ...)                                                                      \
+  do {                                                                                             \
+  } while (false)
+#else
+#define debug_print(fmt, ...)                                                                      \
+  do {                                                                                             \
+    fprintf(stderr, fmt, ##__VA_ARGS__);                                                           \
+  } while (false)
 #endif
 
 // A macro to disallow the copy and move constructor and operator= functions

runtime/hsa-runtime/inc/hsa_api_trace.h

@@ -81,6 +81,19 @@ static inline uint32_t Min(const uint32_t a, const uint32_t b) {
   return (a > b) ? b : a;
 }
 
+// Declarations of APIs intended for use only by tools.
+typedef void (*hsa_amd_queue_intercept_packet_writer)(const void* pkts, uint64_t pkt_count);
+typedef void (*hsa_amd_queue_intercept_handler)(const void* pkts, uint64_t pkt_count,
+                                                uint64_t user_pkt_index, void* data,
+                                                hsa_amd_queue_intercept_packet_writer writer);
+hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue,
+                                              hsa_amd_queue_intercept_handler callback,
+                                              void* user_data);
+hsa_status_t hsa_amd_queue_intercept_create(
+    hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type,
+    void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
+    uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue);
+
 // Structure of Version used to identify an instance of Api table
 // Must be the first member (offsetof == 0) of all API tables.
 // This is the root of the table passing ABI.
@@ -170,6 +183,8 @@ struct AmdExtTable {
   decltype(hsa_amd_ipc_signal_create)* hsa_amd_ipc_signal_create_fn;
   decltype(hsa_amd_ipc_signal_attach)* hsa_amd_ipc_signal_attach_fn;
   decltype(hsa_amd_register_system_event_handler)* hsa_amd_register_system_event_handler_fn;
+  decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
+  decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
 };
 
 // Table to export HSA Core Runtime Apis