[AMDGPU] Remove uses of deprecreated HSA executable functions

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);
 

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()
 

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

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<offloading::amdgpu::AMDGPUKernelMetaData> 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;