diff --git a/runtime/hsa-runtime/CMakeLists.txt b/runtime/hsa-runtime/CMakeLists.txt index 7bf93832c0..93796b3f7d 100644 --- a/runtime/hsa-runtime/CMakeLists.txt +++ b/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() diff --git a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp index 30167db973..58b35f4028 100644 --- a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp +++ b/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() { diff --git a/runtime/hsa-runtime/core/inc/hsa_amd_tool_int.hpp b/runtime/hsa-runtime/core/inc/hsa_amd_tool_int.hpp new file mode 100644 index 0000000000..d0a0248376 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h b/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h index 91f8cc3f15..f270efb72f 100644 --- a/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h +++ b/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(); diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index e8b8f3e80f..fd2891531f 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/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; } diff --git a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index 044b357526..d26b1c99a5 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/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); } diff --git a/runtime/hsa-runtime/hsacore.so.def b/runtime/hsa-runtime/hsacore.so.def index c439643828..4d9c92186f 100644 --- a/runtime/hsa-runtime/hsacore.so.def +++ b/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: *; diff --git a/runtime/hsa-runtime/inc/hsa_amd_tool.h b/runtime/hsa-runtime/inc/hsa_amd_tool.h new file mode 100644 index 0000000000..fa9cac804a --- /dev/null +++ b/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 +#include +#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 \ No newline at end of file diff --git a/runtime/hsa-runtime/inc/hsa_api_trace.h b/runtime/hsa-runtime/inc/hsa_api_trace.h index 51abf6a3ef..e46c777af1 100644 --- a/runtime/hsa-runtime/inc/hsa_api_trace.h +++ b/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 @@ -59,12 +61,13 @@ #include // 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