From 9bc38e2ee6c6d5e5eb870e692c4342dc657e1970 Mon Sep 17 00:00:00 2001 From: Honglei Huang Date: Fri, 27 Jun 2025 13:29:03 +0800 Subject: [PATCH] rocr/driver: add support for getting GPU tile configuration - Implemented GetTileConfig in KfdDriver to retrieve tile configuration for a specific node. - Added a stub implementation of GetTileConfig in XdnaDriver. - Updated driver.h to include a virtual GetTileConfig method. - Extended hsa_internal.h with a new hsa_get_tile_config function. - Integrated hsa_get_tile_config into hsa.cpp to call the driver-specific implementation. - Updated driver headers to declare the new GetTileConfig method. Signed-off-by: Honglei Huang --- .../hsa-runtime/core/driver/kfd/amd_kfd_driver.cpp | 10 ++++++++++ .../core/driver/xdna/amd_xdna_driver.cpp | 5 +++++ runtime/hsa-runtime/core/inc/amd_kfd_driver.h | 1 + runtime/hsa-runtime/core/inc/amd_xdna_driver.h | 1 + runtime/hsa-runtime/core/inc/driver.h | 7 +++++++ runtime/hsa-runtime/core/inc/hsa_internal.h | 1 + runtime/hsa-runtime/core/runtime/hsa.cpp | 13 +++++++++++++ runtime/hsa-runtime/image/image_manager_kv.cpp | 4 ++-- 8 files changed, 40 insertions(+), 2 deletions(-) diff --git a/runtime/hsa-runtime/core/driver/kfd/amd_kfd_driver.cpp b/runtime/hsa-runtime/core/driver/kfd/amd_kfd_driver.cpp index f47d4ac4ee..eafdbe76ec 100644 --- a/runtime/hsa-runtime/core/driver/kfd/amd_kfd_driver.cpp +++ b/runtime/hsa-runtime/core/driver/kfd/amd_kfd_driver.cpp @@ -591,6 +591,16 @@ hsa_status_t KfdDriver::GetClockCounters(uint32_t node_id, HsaClockCounters* clo return HSA_STATUS_SUCCESS; } +hsa_status_t KfdDriver::GetTileConfig(uint32_t node_id, HsaGpuTileConfig* config) const { + assert(config); + + if (HSAKMT_CALL(hsaKmtGetTileConfig(node_id, config)) != HSAKMT_STATUS_SUCCESS) { + return HSA_STATUS_ERROR; + } + + return HSA_STATUS_SUCCESS; +} + hsa_status_t KfdDriver::IsModelEnabled(bool* enable) const { // AIE does not support streaming performance monitor. HSAKMT_STATUS status = HSAKMT_STATUS_ERROR; diff --git a/runtime/hsa-runtime/core/driver/xdna/amd_xdna_driver.cpp b/runtime/hsa-runtime/core/driver/xdna/amd_xdna_driver.cpp index 375364ca55..e2a0f60e82 100644 --- a/runtime/hsa-runtime/core/driver/xdna/amd_xdna_driver.cpp +++ b/runtime/hsa-runtime/core/driver/xdna/amd_xdna_driver.cpp @@ -885,5 +885,10 @@ hsa_status_t XdnaDriver::GetClockCounters(uint32_t node_id, HsaClockCounters* cl return HSA_STATUS_ERROR; } + +hsa_status_t XdnaDriver::GetTileConfig(uint32_t node_id, HsaGpuTileConfig* config) const { + return HSA_STATUS_ERROR; +} + } // namespace AMD } // namespace rocr diff --git a/runtime/hsa-runtime/core/inc/amd_kfd_driver.h b/runtime/hsa-runtime/core/inc/amd_kfd_driver.h index f20619d20a..f673221731 100644 --- a/runtime/hsa-runtime/core/inc/amd_kfd_driver.h +++ b/runtime/hsa-runtime/core/inc/amd_kfd_driver.h @@ -126,6 +126,7 @@ public: const void* buffer_base, uint64_t buffer_base_size) const override; hsa_status_t GetDeviceHandle(uint32_t node_id, void** device_handle) const override; hsa_status_t GetClockCounters(uint32_t node_id, HsaClockCounters* clock_counter) const override; + hsa_status_t GetTileConfig(uint32_t node_id, HsaGpuTileConfig* config) const override; hsa_status_t OpenSMI(uint32_t node_id, int* fd) const override; diff --git a/runtime/hsa-runtime/core/inc/amd_xdna_driver.h b/runtime/hsa-runtime/core/inc/amd_xdna_driver.h index 6857c0e973..f3f0ac2c9b 100644 --- a/runtime/hsa-runtime/core/inc/amd_xdna_driver.h +++ b/runtime/hsa-runtime/core/inc/amd_xdna_driver.h @@ -240,6 +240,7 @@ public: const void* buffer_base, uint64_t buffer_base_size) const override; hsa_status_t GetDeviceHandle(uint32_t node_id, void** device_handle) const override; hsa_status_t GetClockCounters(uint32_t node_id, HsaClockCounters* clock_counter) const override; + hsa_status_t GetTileConfig(uint32_t node_id, HsaGpuTileConfig* config) const override; hsa_status_t IsModelEnabled(bool* enable) const override; diff --git a/runtime/hsa-runtime/core/inc/driver.h b/runtime/hsa-runtime/core/inc/driver.h index 5ae5278597..70ea3b1601 100644 --- a/runtime/hsa-runtime/core/inc/driver.h +++ b/runtime/hsa-runtime/core/inc/driver.h @@ -283,6 +283,13 @@ public: virtual hsa_status_t GetClockCounters(uint32_t node_id, HsaClockCounters* clock_counter) const = 0; + /// @brief Get the tile configuration for a specific node. + /// + /// @param[in] node_id Node ID of the agent + /// @param[out] config Pointer to tile configuration + /// @return HSA_STATUS_SUCCESS if the driver successfully returns the tile configuration. + virtual hsa_status_t GetTileConfig(uint32_t node_id, HsaGpuTileConfig* config) const = 0; + /// @brief Check if the HSA KMT Model is enabled /// @param[out] enable True if the model is enabled, false otherwise virtual hsa_status_t IsModelEnabled(bool* enable) const = 0; diff --git a/runtime/hsa-runtime/core/inc/hsa_internal.h b/runtime/hsa-runtime/core/inc/hsa_internal.h index fb86eb9578..b40c049976 100644 --- a/runtime/hsa-runtime/core/inc/hsa_internal.h +++ b/runtime/hsa-runtime/core/inc/hsa_internal.h @@ -406,6 +406,7 @@ namespace HSA { hsa_executable_symbol_t symbol, void *data), void *data); + hsa_status_t hsa_get_tile_config(hsa_agent_t agent_handle, void* config); //===--- Runtime Notifications ------------------------------------------===// diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index 9152f070db..147ce1ba6d 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -2590,6 +2590,19 @@ hsa_status_t hsa_executable_iterate_program_symbols( CATCH; } +hsa_status_t hsa_get_tile_config(hsa_agent_t agent_handle, void* config) { + TRY; + IS_OPEN(); + IS_BAD_PTR(config); + + const core::Agent* agent_object = core::Agent::Convert(agent_handle); + IS_VALID(agent_object); + + return agent_object->driver().GetTileConfig(agent_object->node_id(), + static_cast(config)); + CATCH; +} + //===--- Runtime Notifications --------------------------------------------===// hsa_status_t hsa_status_string( diff --git a/runtime/hsa-runtime/image/image_manager_kv.cpp b/runtime/hsa-runtime/image/image_manager_kv.cpp index 1cd243d848..b7e142fe3a 100644 --- a/runtime/hsa-runtime/image/image_manager_kv.cpp +++ b/runtime/hsa-runtime/image/image_manager_kv.cpp @@ -110,8 +110,8 @@ hsa_status_t ImageManagerKv::Initialize(hsa_agent_t agent_handle) { status = HSA::hsa_agent_get_info( agent_, static_cast(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID), &node_id); assert(status == HSA_STATUS_SUCCESS); - HSAKMT_STATUS stat = HSAKMT_CALL(hsaKmtGetTileConfig(node_id, &tileConfig)); - assert(stat == HSAKMT_STATUS_SUCCESS); + hsa_status_t stat = HSA::hsa_get_tile_config(agent_handle, &tileConfig); + assert(stat == HSA_STATUS_SUCCESS); // Initialize address library. // TODO(bwicakso) hard coded based on UGL parameters.