From 2b79f8fd20e62d5b3eb57d9ba15af192ac1cb5e5 Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Thu, 21 Nov 2024 15:08:49 -0600 Subject: [PATCH] [AMDGPU] Remove uses of deprecreated HSA executable functions Summary: These functions were deprecated in ROCR 1.3 which was released quite some time ago. The main functionality that was lost was modifying and inspecting the code object indepedently of the executable, however we do all of that custom through our ELF API. This should be within the versions of other functions we use. --- .../utils/gpu/loader/amdgpu/amdhsa-loader.cpp | 16 +++++----- .../amdgpu/dynamic_hsa/hsa.cpp | 3 ++ .../plugins-nextgen/amdgpu/dynamic_hsa/hsa.h | 20 +++++++++++++ offload/plugins-nextgen/amdgpu/src/rtl.cpp | 30 ++++++++++--------- 4 files changed, 48 insertions(+), 21 deletions(-) diff --git a/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp b/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp index cb81a866622f9..d825a6299263a 100644 --- a/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp +++ b/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp @@ -371,8 +371,9 @@ int load(int argc, const char **argv, const char **envp, void *image, handle_error(err); // Load the code object's ISA information and executable data segments. - hsa_code_object_t object; - if (hsa_status_t err = hsa_code_object_deserialize(image, size, "", &object)) + hsa_code_object_reader_t reader; + if (hsa_status_t err = + hsa_code_object_reader_create_from_memory(image, size, &reader)) handle_error(err); hsa_executable_t executable; @@ -381,8 +382,9 @@ int load(int argc, const char **argv, const char **envp, void *image, &executable)) handle_error(err); - if (hsa_status_t err = - hsa_executable_load_code_object(executable, dev_agent, object, "")) + hsa_loaded_code_object_t object; + if (hsa_status_t err = hsa_executable_load_agent_code_object( + executable, dev_agent, reader, "", &object)) handle_error(err); // No modifications to the executable are allowed after this point. @@ -397,6 +399,9 @@ int load(int argc, const char **argv, const char **envp, void *image, if (result) handle_error(HSA_STATUS_ERROR); + if (hsa_status_t err = hsa_code_object_reader_destroy(reader)) + handle_error(err); + // Obtain memory pools to exchange data between the host and the device. The // fine-grained pool acts as pinned memory on the host for DMA transfers to // the device, the coarse-grained pool is for allocations directly on the @@ -617,9 +622,6 @@ int load(int argc, const char **argv, const char **envp, void *image, if (hsa_status_t err = hsa_executable_destroy(executable)) handle_error(err); - if (hsa_status_t err = hsa_code_object_destroy(object)) - handle_error(err); - if (hsa_status_t err = hsa_shut_down()) handle_error(err); diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp index cb82180b9a27e..bc92f4a46a5c0 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp @@ -68,6 +68,9 @@ DLWRAP(hsa_amd_register_system_event_handler, 2) DLWRAP(hsa_amd_signal_create, 5) DLWRAP(hsa_amd_signal_async_handler, 5) DLWRAP(hsa_amd_pointer_info, 5) +DLWRAP(hsa_code_object_reader_create_from_memory, 3) +DLWRAP(hsa_code_object_reader_destroy, 1) +DLWRAP(hsa_executable_load_agent_code_object, 5) DLWRAP_FINALIZE() diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h index 5d9fb5d7dc7cd..27e4541481301 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h @@ -50,6 +50,14 @@ typedef struct hsa_agent_s { uint64_t handle; } hsa_agent_t; +typedef struct hsa_loaded_code_object_s { + uint64_t handle; +} hsa_loaded_code_object_t; + +typedef struct hsa_code_object_reader_s { + uint64_t handle; +} hsa_code_object_reader_t; + typedef enum { HSA_DEVICE_TYPE_CPU = 0, HSA_DEVICE_TYPE_GPU = 1, @@ -364,6 +372,18 @@ hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal, hsa_amd_signal_handler handler, void *arg); +hsa_status_t hsa_code_object_reader_create_from_memory( + const void *code_object, size_t size, + hsa_code_object_reader_t *code_object_reader); + +hsa_status_t +hsa_code_object_reader_destroy(hsa_code_object_reader_t code_object_reader); + +hsa_status_t hsa_executable_load_agent_code_object( + hsa_executable_t executable, hsa_agent_t agent, + hsa_code_object_reader_t code_object_reader, const char *options, + hsa_loaded_code_object_t *loaded_code_object); + #ifdef __cplusplus } #endif diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp index bdb33d4f4ab27..6356fa0554a9c 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -468,11 +468,7 @@ struct AMDGPUDeviceImageTy : public DeviceImageTy { /// Unload the executable. Error unloadExecutable() { hsa_status_t Status = hsa_executable_destroy(Executable); - if (auto Err = Plugin::check(Status, "Error in hsa_executable_destroy: %s")) - return Err; - - Status = hsa_code_object_destroy(CodeObject); - return Plugin::check(Status, "Error in hsa_code_object_destroy: %s"); + return Plugin::check(Status, "Error in hsa_executable_destroy: %s"); } /// Get the executable. @@ -499,7 +495,6 @@ struct AMDGPUDeviceImageTy : public DeviceImageTy { private: /// The exectuable loaded on the agent. hsa_executable_t Executable; - hsa_code_object_t CodeObject; StringMap KernelInfoMap; uint16_t ELFABIVersion; }; @@ -2954,10 +2949,11 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { }; Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) { - hsa_status_t Status; - Status = hsa_code_object_deserialize(getStart(), getSize(), "", &CodeObject); - if (auto Err = - Plugin::check(Status, "Error in hsa_code_object_deserialize: %s")) + hsa_code_object_reader_t Reader; + hsa_status_t Status = + hsa_code_object_reader_create_from_memory(getStart(), getSize(), &Reader); + if (auto Err = Plugin::check( + Status, "Error in hsa_code_object_reader_create_from_memory: %s")) return Err; Status = hsa_executable_create_alt( @@ -2966,10 +2962,11 @@ Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) { Plugin::check(Status, "Error in hsa_executable_create_alt: %s")) return Err; - Status = hsa_executable_load_code_object(Executable, Device.getAgent(), - CodeObject, ""); - if (auto Err = - Plugin::check(Status, "Error in hsa_executable_load_code_object: %s")) + hsa_loaded_code_object_t Object; + Status = hsa_executable_load_agent_code_object(Executable, Device.getAgent(), + Reader, "", &Object); + if (auto Err = Plugin::check( + Status, "Error in hsa_executable_load_agent_code_object: %s")) return Err; Status = hsa_executable_freeze(Executable, ""); @@ -2984,6 +2981,11 @@ Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) { if (Result) return Plugin::error("Loaded HSA executable does not validate"); + Status = hsa_code_object_reader_destroy(Reader); + if (auto Err = + Plugin::check(Status, "Error in hsa_code_object_reader_destroy: %s")) + return Err; + if (auto Err = hsa_utils::readAMDGPUMetaDataFromImage( getMemoryBuffer(), KernelInfoMap, ELFABIVersion)) return Err;