Initial support for scratch allocation tracking

Add new tools table and functions to notify in case of an event

Change-Id: I47f0c2f3c8e02d7bcb74d649903eb4f86721c154

runtime/hsa-runtime/CMakeLists.txt

@@ -85,7 +85,7 @@ if (ROCM_CCACHE_BUILD)
 endif() # if (ROCM_CCACHE_BUILD)
 
 ## Get version strings
-get_version ( "1.12.0" )
+get_version ( "1.13.0" )
 if ( ${ROCM_PATCH_VERSION} )
   set ( VERSION_PATCH ${ROCM_PATCH_VERSION})
 endif()

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

@@ -46,11 +46,13 @@
 static const HsaApiTable* hsaApiTable;
 static const CoreApiTable* coreApiTable;
 static const AmdExtTable* amdExtTable;
+static const ToolsApiTable* toolsApiTable;
 
 void hsa_table_interface_init(const HsaApiTable* apiTable) {
     hsaApiTable = apiTable;
     coreApiTable = apiTable->core_;
     amdExtTable = apiTable->amd_ext_;
+    toolsApiTable = apiTable->tools_;
 }
 
 const HsaApiTable* hsa_table_interface_get_table() {

runtime/hsa-runtime/core/inc/hsa_amd_tool_int.hpp

@@ -0,0 +1,141 @@
+#ifndef HSA_RUNTIME_INC_HSA_TOOL_HOOK_IMPL_H
+#define HSA_RUNTIME_INC_HSA_TOOL_HOOK_IMPL_H
+
+#include "inc/hsa_amd_tool.h"
+#include "runtime.h"
+
+namespace rocr::AMD::tool {
+
+using scratch_alloc_flag = hsa_amd_event_scratch_alloc_flag_t;
+
+__forceinline void notify_event_scratch_alloc_start(const hsa_queue_t* queue,
+                                                    scratch_alloc_flag flag, uint64_t dispatch_id);
+
+__forceinline void notify_event_scratch_alloc_end(const hsa_queue_t* queue, scratch_alloc_flag flag,
+                                                  uint64_t dispatch_id, size_t size,
+                                                  size_t num_slots);
+
+__forceinline void notify_event_scratch_free_start(const hsa_queue_t* queue,
+                                                   scratch_alloc_flag flag);
+
+__forceinline void notify_event_scratch_free_end(const hsa_queue_t* queue, scratch_alloc_flag flag);
+
+__forceinline void notify_event_scratch_async_reclaim_start(const hsa_queue_t* queue,
+                                                            scratch_alloc_flag flag);
+
+__forceinline void notify_event_scratch_async_reclaim_end(const hsa_queue_t* queue,
+                                                          scratch_alloc_flag flag);
+
+
+// Impl
+
+__forceinline void notify_event_scratch_alloc_start(const hsa_queue_t* queue,
+                                                    scratch_alloc_flag flags,
+                                                    uint64_t dispatch_id) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_alloc_start_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_alloc_start_t{.kind = HSA_AMD_TOOL_EVENT_SCRATCH_ALLOC_START,
+                                                   .queue = queue,
+                                                   .flags = flags,
+                                                   .dispatch_id = dispatch_id};
+
+  tool_table.hsa_amd_tool_scratch_event_alloc_start_fn(
+      hsa_amd_tool_event_t{.scratch_alloc_start = &event});
+}
+
+__forceinline void notify_event_scratch_alloc_end(const hsa_queue_t* queue,
+                                                  scratch_alloc_flag flags, uint64_t dispatch_id,
+                                                  size_t size, size_t num_slots) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_alloc_end_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_alloc_end_t{
+      .kind = HSA_AMD_TOOL_EVENT_SCRATCH_ALLOC_END,
+      .queue = queue,
+      .flags = flags,
+      .dispatch_id = dispatch_id,
+      .size = size,
+      .num_slots = num_slots,
+  };
+
+  tool_table.hsa_amd_tool_scratch_event_alloc_end_fn(
+      hsa_amd_tool_event_t{.scratch_alloc_end = &event});
+}
+
+__forceinline void notify_event_scratch_free_start(const hsa_queue_t* queue,
+                                                   scratch_alloc_flag flags) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_free_start_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_free_start_t{
+      .kind = HSA_AMD_TOOL_EVENT_SCRATCH_FREE_START,
+      .queue = queue,
+      .flags = flags,
+  };
+
+  tool_table.hsa_amd_tool_scratch_event_free_start_fn(
+      hsa_amd_tool_event_t{.scratch_free_start = &event});
+}
+
+__forceinline void notify_event_scratch_free_end(const hsa_queue_t* queue,
+                                                 scratch_alloc_flag flags) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_free_end_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_free_end_t{
+      .kind = HSA_AMD_TOOL_EVENT_SCRATCH_FREE_END,
+      .queue = queue,
+      .flags = flags,
+  };
+
+  tool_table.hsa_amd_tool_scratch_event_free_end_fn(
+      hsa_amd_tool_event_t{.scratch_free_end = &event});
+}
+
+__forceinline void notify_event_scratch_async_reclaim_start(const hsa_queue_t* queue,
+                                                            scratch_alloc_flag flags) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_async_reclaim_start_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_async_reclaim_start_t{
+      .kind = HSA_AMD_TOOL_EVENT_SCRATCH_ASYNC_RECLAIM_START,
+      .queue = queue,
+      .flags = flags,
+  };
+
+  tool_table.hsa_amd_tool_scratch_event_async_reclaim_start_fn(
+      hsa_amd_tool_event_t{.scratch_async_reclaim_start = &event});
+}
+
+__forceinline void notify_event_scratch_async_reclaim_end(const hsa_queue_t* queue,
+                                                          scratch_alloc_flag flags) {
+  const auto& tool_table = core::hsa_api_table_.tools_api;
+  if (!tool_table.hsa_amd_tool_scratch_event_async_reclaim_end_fn) {
+    return;
+  }
+
+  auto event = hsa_amd_event_scratch_async_reclaim_end_t{
+      .kind = HSA_AMD_TOOL_EVENT_SCRATCH_ASYNC_RECLAIM_END,
+      .queue = queue,
+      .flags = flags,
+  };
+
+  tool_table.hsa_amd_tool_scratch_event_async_reclaim_end_fn(
+      hsa_amd_tool_event_t{.scratch_async_reclaim_end = &event});
+}
+
+
+}  // namespace rocr::AMD::tool
+
+#endif

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

@@ -59,11 +59,13 @@ namespace core {
     ::AmdExtTable amd_ext_api;
     ::FinalizerExtTable finalizer_api;
     ::ImageExtTable image_api;
+    ::ToolsApiTable tools_api;
 
     HsaApiTable();
     void Init();
     void UpdateCore();
     void UpdateAmdExts();
+    void UpdateTools();
     void CloneExts(void* ptr, uint32_t table_id);
     void LinkExts(void* ptr, uint32_t table_id);
     void Reset();

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

@@ -67,6 +67,7 @@
 #include "core/inc/default_signal.h"
 #include "core/inc/hsa_ext_amd_impl.h"
 #include "core/inc/amd_gpu_pm4.h"
+#include "core/inc/hsa_amd_tool_int.hpp"
 
 namespace rocr {
 namespace AMD {
@@ -827,6 +828,9 @@ void AqlQueue::AsyncReclaimMainScratch() {
 
   // Notify CP that we are trying to reclaim scratch. CP will assume scratch is reclaimed on next
   // dispatch
+  tool::notify_event_scratch_async_reclaim_start(public_handle(),
+                                                 HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_NONE);
+
   amd_queue_.scratch_wave64_lane_byte_size = 0;
   uint64_t last_used =
       atomic::Exchange(&amd_queue_.scratch_last_used_index, UINT64_MAX, std::memory_order_relaxed);
@@ -837,6 +841,8 @@ void AqlQueue::AsyncReclaimMainScratch() {
 
     if (std::min(last, last_used) < amd_queue_.read_dispatch_id) {
       FreeMainScratchSpace();
+      tool::notify_event_scratch_async_reclaim_end(public_handle(),
+                                                   HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_NONE);
       return;
     }
   }
@@ -859,6 +865,9 @@ void AqlQueue::AsyncReclaimAltScratch() {
 
   // Notify CP that we are trying to reclaim scratch. CP will assume scratch is reclaimed on next
   // dispatch
+  tool::notify_event_scratch_async_reclaim_start(public_handle(),
+                                                 HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_ALT);
+
   amd_queue_.alt_scratch_wave64_lane_byte_size = 0;
   uint64_t last_used = atomic::Exchange(
       &amd_queue_.alt_scratch_last_used_index, UINT64_MAX,
@@ -869,6 +878,8 @@ void AqlQueue::AsyncReclaimAltScratch() {
 
     if (std::min(last, last_used) < amd_queue_.read_dispatch_id) {
       FreeAltScratchSpace();
+      tool::notify_event_scratch_async_reclaim_end(public_handle(),
+                                                   HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_ALT);
       return;
     }
   }
@@ -912,6 +923,10 @@ void AqlQueue::HandleInsufficientScratch(hsa_signal_value_t& error_code,
    *
    *******************************************************************************************/
 
+  const auto& dispatch_id = amd_queue_.read_dispatch_id;
+  tool::notify_event_scratch_alloc_start(public_handle(), HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_NONE,
+                                         dispatch_id);
+
   auto calc_dispatch_waves_per_group = [&](core::AqlPacket& pkt) {
     const uint64_t lanes_per_group =
         (uint64_t(pkt.dispatch.workgroup_size_x) * pkt.dispatch.workgroup_size_y) *
@@ -970,7 +985,7 @@ void AqlQueue::HandleInsufficientScratch(hsa_signal_value_t& error_code,
 
   scratch.cooperative = (amd_queue_.hsa_queue.type == HSA_QUEUE_TYPE_COOPERATIVE);
 
-  uint64_t pkt_slot_idx = amd_queue_.read_dispatch_id & (amd_queue_.hsa_queue.size - 1);
+  uint64_t pkt_slot_idx = dispatch_id & (amd_queue_.hsa_queue.size - 1);
 
   core::AqlPacket& pkt = ((core::AqlPacket*)amd_queue_.hsa_queue.base_address)[pkt_slot_idx];
 
@@ -1010,7 +1025,8 @@ void AqlQueue::HandleInsufficientScratch(hsa_signal_value_t& error_code,
       InitScratchSRD();
       // Restart the queue.
       HSA::hsa_signal_store_screlease(amd_queue_.queue_inactive_signal, 0);
-
+      tool::notify_event_scratch_alloc_end(public_handle(), HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_ALT,
+                                           dispatch_id, scratch.alt_size, dispatch_slots);
       return;
     }
     // Could not allocate enough memory for alternate scratch fallback to primary scratch
@@ -1036,6 +1052,8 @@ void AqlQueue::HandleInsufficientScratch(hsa_signal_value_t& error_code,
     waitVal = error_code;
   } else if (scratch.main_queue_base == nullptr) {
     // We could not allocate memory to fit even 1 wave
+    tool::notify_event_scratch_alloc_end(public_handle(), HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_USE_ONCE,
+                                         dispatch_id, scratch.main_size, dispatch_slots);
     return;
   }
 
@@ -1058,6 +1076,12 @@ void AqlQueue::HandleInsufficientScratch(hsa_signal_value_t& error_code,
   InitScratchSRD();
   // Restart the queue.
   HSA::hsa_signal_store_screlease(amd_queue_.queue_inactive_signal, 0);
+
+  auto alloc_flag = (scratch.large) ? HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_USE_ONCE
+                                    : HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_NONE;
+  tool::notify_event_scratch_alloc_end(public_handle(), alloc_flag, dispatch_id, scratch.main_size,
+                                       dispatch_slots);
+
   return;
 }
 
@@ -1080,6 +1104,9 @@ bool AqlQueue::DynamicQueueEventsHandler(hsa_signal_value_t error_code, void* ar
   // Process errors only if queue is not terminating.
   if ((queue->dynamicScratchState & ERROR_HANDLER_TERMINATE) != ERROR_HANDLER_TERMINATE) {
     if (error_code == 512) {  // Large scratch reclaim
+      tool::notify_event_scratch_free_start(queue->public_handle(),
+                                            HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_USE_ONCE);
+
       auto& scratch = queue->queue_scratch_;
       queue->agent_->ReleaseQueueMainScratch(scratch);
       scratch.main_queue_base = nullptr;
@@ -1094,6 +1121,8 @@ bool AqlQueue::DynamicQueueEventsHandler(hsa_signal_value_t error_code, void* ar
                     queue->amd_queue_.queue_properties & (~AMD_QUEUE_PROPERTIES_USE_SCRATCH_ONCE),
                     std::memory_order_release);
       atomic::Fence(std::memory_order_release);
+      tool::notify_event_scratch_free_end(queue->public_handle(),
+                                          HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_USE_ONCE);
       return true;
     }
 

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

@@ -83,6 +83,7 @@ void HsaApiTable::Init() {
   constexpr size_t expected_amd_ext_table_size = 560;
   constexpr size_t expected_image_ext_table_size = 120;
   constexpr size_t expected_finalizer_ext_table_size = 64;
+  constexpr size_t expected_tools_table_size = 64;
 
   static_assert(sizeof(CoreApiTable) == expected_core_api_table_size,
                 "HSA core API table size changed, bump HSA_CORE_API_TABLE_STEP_VERSION and set "
@@ -97,6 +98,9 @@ void HsaApiTable::Init() {
   static_assert(sizeof(FinalizerExtTable) == expected_finalizer_ext_table_size,
                 "HSA finalizer ext table size changed, bump HSA_FINALIZER_API_TABLE_STEP_VERSION "
                 "and set expected_finalizer_ext_table_size to the new size of the struct");
+  static_assert(sizeof(ToolsApiTable) == expected_tools_table_size,
+                "HSA tools table size changed, bump HSA_TOOLS_API_TABLE_STEP_VERSION "
+                "and set expected_tools_table_size to the new size of the struct");
 
   // Initialize Version of Api Table
   hsa_api.version.major_id = HSA_API_TABLE_MAJOR_VERSION;
@@ -116,6 +120,8 @@ void HsaApiTable::Init() {
   // of Hsa Runtime initialization, including their major ids
   hsa_api.finalizer_ext_ = NULL;
   hsa_api.image_ext_ = NULL;
+
+  hsa_api.tools_ = &tools_api;
 }
 
 void HsaApiTable::Reset() {
@@ -440,6 +446,15 @@ void HsaApiTable::UpdateAmdExts() {
   amd_ext_api.hsa_amd_agent_set_async_scratch_limit_fn = AMD::hsa_amd_agent_set_async_scratch_limit;
 }
 
+void HsaApiTable::UpdateTools() {
+  tools_api.hsa_amd_tool_scratch_event_alloc_start_fn = nullptr;
+  tools_api.hsa_amd_tool_scratch_event_alloc_end_fn = nullptr;
+  tools_api.hsa_amd_tool_scratch_event_free_start_fn = nullptr;
+  tools_api.hsa_amd_tool_scratch_event_free_end_fn = nullptr;
+  tools_api.hsa_amd_tool_scratch_event_async_reclaim_start_fn = nullptr;
+  tools_api.hsa_amd_tool_scratch_event_async_reclaim_end_fn = nullptr;
+}
+
 void LoadInitialHsaApiTable() {
   hsa_table_interface_init(&hsa_api_table_.hsa_api);
 }

runtime/hsa-runtime/hsacore.so.def

@@ -246,6 +246,12 @@ global:
 	hsa_amd_vmem_retain_alloc_handle;
 	hsa_amd_vmem_get_alloc_properties_from_handle;
 	hsa_amd_agent_set_async_scratch_limit;
+	hsa_tools_scratch_event_alloc_start;
+	hsa_tools_scratch_event_alloc_end;
+	hsa_tools_scratch_event_free_start;
+	hsa_tools_scratch_event_free_end;
+	hsa_tools_scratch_event_async_reclaim_start;
+	hsa_tools_scratch_event_async_reclaim_end;
 
 local:
     *;

runtime/hsa-runtime/inc/hsa_amd_tool.h

@@ -0,0 +1,91 @@
+#ifndef HSA_RUNTIME_AMD_TOOL_EVENTS_H_
+#define HSA_RUNTIME_AMD_TOOL_EVENTS_H_
+
+// Insert license header
+
+#include <stddef.h>
+#include <stdint.h>
+#include "hsa.h"
+
+
+typedef enum {
+  HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_NONE = 0,
+  HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_USE_ONCE =
+      (1 << 0),  // This scratch allocation is only valid for 1 dispatch.
+  HSA_AMD_EVENT_SCRATCH_ALLOC_FLAG_ALT =
+      (1 << 1),  // Used alternate scratch instead of main scratch
+} hsa_amd_event_scratch_alloc_flag_t;
+
+typedef enum {
+  HSA_AMD_TOOL_EVENT_MIN = 0,
+
+  // Scratch memory tracking
+  HSA_AMD_TOOL_EVENT_SCRATCH_ALLOC_START,
+  HSA_AMD_TOOL_EVENT_SCRATCH_ALLOC_END,
+  HSA_AMD_TOOL_EVENT_SCRATCH_FREE_START,
+  HSA_AMD_TOOL_EVENT_SCRATCH_FREE_END,
+  HSA_AMD_TOOL_EVENT_SCRATCH_ASYNC_RECLAIM_START,
+  HSA_AMD_TOOL_EVENT_SCRATCH_ASYNC_RECLAIM_END,
+
+  // Add new events above ^
+  HSA_AMD_TOOL_EVENT_MAX
+} hsa_amd_tool_event_kind_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+} hsa_amd_tool_event_none_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+  uint64_t dispatch_id;  // Dispatch ID of the AQL packet that needs more scratch memory
+} hsa_amd_event_scratch_alloc_start_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+  uint64_t dispatch_id;  // Dispatch ID of the AQL packet that needs more scratch memory
+  size_t size;           // Amount of scratch allocated - in bytes
+  size_t num_slots;      // limit of number of waves
+} hsa_amd_event_scratch_alloc_end_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+} hsa_amd_event_scratch_free_start_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+} hsa_amd_event_scratch_free_end_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+} hsa_amd_event_scratch_async_reclaim_start_t;
+
+typedef struct {
+  hsa_amd_tool_event_kind_t kind;
+  const hsa_queue_t* queue;
+  hsa_amd_event_scratch_alloc_flag_t flags;
+} hsa_amd_event_scratch_async_reclaim_end_t;
+
+typedef union {
+  const hsa_amd_tool_event_none_t* none;
+  const hsa_amd_event_scratch_alloc_start_t* scratch_alloc_start;
+  const hsa_amd_event_scratch_alloc_end_t* scratch_alloc_end;
+  const hsa_amd_event_scratch_free_start_t* scratch_free_start;
+  const hsa_amd_event_scratch_free_end_t* scratch_free_end;
+  const hsa_amd_event_scratch_async_reclaim_start_t* scratch_async_reclaim_start;
+  const hsa_amd_event_scratch_async_reclaim_end_t* scratch_async_reclaim_end;
+} hsa_amd_tool_event_t;
+
+typedef hsa_status_t (*hsa_amd_tool_event)(hsa_amd_tool_event_t);
+
+
+#endif

runtime/hsa-runtime/inc/hsa_api_trace.h

@@ -48,10 +48,12 @@
 #include "hsa_ext_image.h"
 #include "hsa_ext_amd.h"
 #include "hsa_ext_finalize.h"
+#include "hsa_amd_tool.h"
 #else
 #include "inc/hsa_ext_image.h"
 #include "inc/hsa_ext_amd.h"
 #include "inc/hsa_ext_finalize.h"
+#include "inc/hsa_amd_tool.h"
 #endif
 
 #include <string.h>
@@ -59,12 +61,13 @@
 #include <stddef.h>
 
 // Major Ids of the Api tables exported by Hsa Core Runtime
-#define HSA_API_TABLE_MAJOR_VERSION               0x02
+#define HSA_API_TABLE_MAJOR_VERSION               0x03
 #define HSA_CORE_API_TABLE_MAJOR_VERSION          0x02
 #define HSA_AMD_EXT_API_TABLE_MAJOR_VERSION       0x02
 #define HSA_FINALIZER_API_TABLE_MAJOR_VERSION     0x02
 #define HSA_IMAGE_API_TABLE_MAJOR_VERSION         0x02
 #define HSA_AQLPROFILE_API_TABLE_MAJOR_VERSION    0x01
+#define HSA_TOOLS_API_TABLE_MAJOR_VERSION         0x01
 
 // Step Ids of the Api tables exported by Hsa Core Runtime
 #define HSA_API_TABLE_STEP_VERSION                0x00
@@ -73,6 +76,7 @@
 #define HSA_FINALIZER_API_TABLE_STEP_VERSION      0x00
 #define HSA_IMAGE_API_TABLE_STEP_VERSION          0x00
 #define HSA_AQLPROFILE_API_TABLE_STEP_VERSION     0x00
+#define HSA_TOOLS_API_TABLE_STEP_VERSION          0x00
 
 // Min function used to copy Api Tables
 static inline uint32_t Min(const uint32_t a, const uint32_t b) {
@@ -147,6 +151,17 @@ struct ApiTableVersion {
   uint32_t reserved;
 };
 
+struct ToolsApiTable {
+  ApiTableVersion version;
+
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_alloc_start_fn;
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_alloc_end_fn;
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_free_start_fn;
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_free_end_fn;
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_async_reclaim_start_fn;
+  hsa_amd_tool_event hsa_amd_tool_scratch_event_async_reclaim_end_fn;
+};
+
 // Table to export HSA Finalizer Extension Apis
 struct FinalizerExtTable {
   ApiTableVersion version;
@@ -446,6 +461,9 @@ struct HsaApiTable {
 
   // Table of function pointers to HSA Image Extension
 	ImageExtTable* image_ext_;
+
+  // Table of function pointers for tools to use
+  ToolsApiTable* tools_;
 };
 
 // Structure containing instances of different api tables
@@ -455,6 +473,7 @@ struct HsaApiTableContainer {
 	AmdExtTable amd_ext;
 	FinalizerExtTable finalizer_ext;
 	ImageExtTable image_ext;
+	ToolsApiTable tools;
 
   // Default initialization of a container instance
   HsaApiTableContainer() {
@@ -475,12 +494,17 @@ struct HsaApiTableContainer {
     finalizer_ext.version.major_id = HSA_FINALIZER_API_TABLE_MAJOR_VERSION;
     finalizer_ext.version.minor_id = sizeof(FinalizerExtTable);
     finalizer_ext.version.step_id = HSA_FINALIZER_API_TABLE_STEP_VERSION;
-    root.finalizer_ext_ = & finalizer_ext;
+    root.finalizer_ext_ = &finalizer_ext;
 
     image_ext.version.major_id = HSA_IMAGE_API_TABLE_MAJOR_VERSION;
     image_ext.version.minor_id = sizeof(ImageExtTable);
     image_ext.version.step_id = HSA_IMAGE_API_TABLE_STEP_VERSION;
     root.image_ext_ = &image_ext;
+
+    tools.version.major_id = HSA_TOOLS_API_TABLE_MAJOR_VERSION;
+    tools.version.minor_id = sizeof(ToolsApiTable);
+    tools.version.step_id = HSA_TOOLS_API_TABLE_STEP_VERSION;
+    root.tools_ = &tools;
   }
 };
 
@@ -536,5 +560,7 @@ static void inline copyTables(const HsaApiTable* src, HsaApiTable* dest) {
     copyElement(&dest->finalizer_ext_->version, &src->finalizer_ext_->version);
   if ((offsetof(HsaApiTable, image_ext_) < dest->version.minor_id))
     copyElement(&dest->image_ext_->version, &src->image_ext_->version);
+  if ((offsetof(HsaApiTable, tools_) < dest->version.minor_id))
+    copyElement(&dest->tools_->version, &src->tools_->version);
 }
 #endif