[libc-commits] [libc] [llvm] [AMDGPU] Remove uses of deprecreated HSA executable functions (PR #117241)
via libc-commits
libc-commits at lists.llvm.org
Thu Nov 21 13:12:16 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-libc
Author: Joseph Huber (jhuber6)
<details>
<summary>Changes</summary>
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.
---
Full diff: https://github.com/llvm/llvm-project/pull/117241.diff
4 Files Affected:
- (modified) libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp (+9-7)
- (modified) offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp (+3)
- (modified) offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h (+20)
- (modified) offload/plugins-nextgen/amdgpu/src/rtl.cpp (+16-14)
``````````diff
diff --git a/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp b/libc/utils/gpu/loader/amdgpu/amdhsa-loader.cpp
index cb81a866622f93..d825a6299263ae 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 cb82180b9a27e8..bc92f4a46a5c0e 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 5d9fb5d7dc7cdd..27e45414813011 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 bdb33d4f4ab27c..6356fa0554a9c1 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<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;
``````````
</details>
https://github.com/llvm/llvm-project/pull/117241
More information about the libc-commits
mailing list