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;