[AMDGPU] Remove uses of deprecreated HSA executable functions (llvm#117241)

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.

Author: jhuber6

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;