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 <Honglei1.Huang@amd.com>

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;

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

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;
 

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;
 

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;

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 ------------------------------------------===//
 

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<HsaGpuTileConfig*>(config));
+  CATCH;
+}
+
 //===--- Runtime Notifications --------------------------------------------===//
 
 hsa_status_t hsa_status_string(

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_agent_info_t>(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.