[libomptarget][amdgpu] Implement dlopen of libhsa

[JonChesterfield]

Summary

AMDGPU plugin equivalent of D95155, build without HSA installed locally

Compiles a new file, plugins/amdgpu/dynamic_hsa/hsa.cpp, to an object file that
exposes the same symbols that the plugin presently uses from hsa. The object
file contains dlopen of hsa and cached dlsym calls. Also provides header files
corresponding to the subset that is used.

This is behind a feature flag, LIBOMPTARGET_FORCE_DLOPEN_LIBHSA, default off.
That allows developers to build against the dlopen/dlsym implementation, e.g.
while testing this mode.

Enabling by default will cause this plugin to build on a wider variety of
machines than it does at present so may break some CI builds. That risk can
be minimised by reviewing the header dependencies of the library and ensuring
it doesn't use any libraries that are not already used by libomptarget.

Separating the implementation from enabling by default in case the latter needs
to be rolled back after wider CI results.

[JonChesterfield]

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

91

drive by fix, this logic was inspired by the nvptx one, where CMAKE_DL_LIBS is used instead of dl

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h

2

This file was semi-automatically reduced from https://github.com/RadeonOpenCompute/ROCR-Runtime/tree/master/src/inc and preserves the header block as written there.

Dumped hsa symbols used by the amdgpu plugin then used grep to find them in the ROCR headers, manually cleaned up, then copied across the type definitions and enums that are currently in use.

The ROCR API is expected to be backwards compatible. Elided the comments on the basis that developers are better off reading the official header files.

[jdoerfert]

LG, it's a plugin.

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h

42

Please use the LLVM licence header.

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa_ext_amd.h

42

LLVM licence.

[JonChesterfield]

• Change header to llvm as requested in review

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

if(NOT CMAKE_SYSTEM_PROCESSOR MATCHES "(x86_64)|(ppc64le)|(aarch64)$" AND CMAKE_SYSTEM_NAME MATCHES "Linux")
   return()
 endif()
 
 if (NOT LIBOMPTARGET_LLVM_INCLUDE_DIRS)
   libomptarget_say("Not building AMDGPU plugin: Missing definition for LIBOMPTARGET_LLVM_INCLUDE_DIRS")
   return()
 endif()
 
-libomptarget_say("Building amdgpu offloading plugin")
-
 ################################################################################
 # Define the suffix for the runtime messaging dumps.
 add_definitions(-DTARGET_NAME=AMDGPU)
 if(CMAKE_SYSTEM_PROCESSOR MATCHES "(ppc64le)|(aarch64)$")
    add_definitions(-DLITTLEENDIAN_CPU=1)
 endif()
 
 if(CMAKE_BUILD_TYPE MATCHES Debug)
   add_definitions(-DDEBUG)
 endif()
 
 include_directories(
   ${CMAKE_CURRENT_SOURCE_DIR}/impl
   ${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
 )
 
+set(LIBOMPTARGET_DLOPEN_LIBHSA OFF)
+option(LIBOMPTARGET_FORCE_DLOPEN_LIBHSA "Build with dlopened libhsa" ${LIBOMPTARGET_DLOPEN_LIBHSA})
+
+if (NOT LIBOMPTARGET_FORCE_DLOPEN_LIBHSA)
+  libomptarget_say("Building AMDGPU plugin linked against libhsa")
+  set(LIBOMPTARGET_EXTRA_SOURCE)
+  set(LIBOMPTARGET_DEP_LIBRARIES hsa-runtime64::hsa-runtime64)
+else()
+  libomptarget_say("Building AMDGPU plugin for dlopened libhsa")
+  include_directories(dynamic_hsa)
+  set(LIBOMPTARGET_EXTRA_SOURCE dynamic_hsa/hsa.cpp)
+  set(LIBOMPTARGET_DEP_LIBRARIES)
+endif()
+
 add_library(omptarget.rtl.amdgpu SHARED
       impl/atmi.cpp
       impl/atmi_interop_hsa.cpp
       impl/data.cpp
       impl/get_elf_mach_gfx_name.cpp
       impl/machine.cpp
       impl/system.cpp
       impl/utils.cpp
       impl/msgpack.cpp
       src/rtl.cpp
+      ${LIBOMPTARGET_EXTRA_SOURCE}
       )
 
 # Install plugin under the lib destination folder.
 # When we build for debug, OPENMP_LIBDIR_SUFFIX get set to -debug
 install(TARGETS omptarget.rtl.amdgpu LIBRARY DESTINATION "lib${OPENMP_LIBDIR_SUFFIX}")
 
 set_property(TARGET omptarget.rtl.amdgpu PROPERTY INSTALL_RPATH "$ORIGIN")
 target_link_libraries(
   omptarget.rtl.amdgpu
   PRIVATE
   elf_common
-  hsa-runtime64::hsa-runtime64
-  dl
+  ${LIBOMPTARGET_DEP_LIBRARIES}
+  ${CMAKE_DL_LIBS}
   ${LIBOMPTARGET_DEP_LIBELF_LIBRARIES}
   ${OPENMP_PTHREAD_LIB}
   "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exports"
   "-Wl,-z,defs"
   )
 
 if (LLVM_BINARY_DIR)
   # for in-tree build using LLVM_ENABLE_RUNTIMES

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h

+//===--- amdgpu/dynamic_hsa/hsa.h --------------------------------- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// The parts of the hsa api that are presently in use by the amdgpu plugin
+//
+//===----------------------------------------------------------------------===//
+#ifndef HSA_RUNTIME_INC_HSA_H_
+#define HSA_RUNTIME_INC_HSA_H_
+
+#include <stddef.h>
+#include <stdint.h>
+
+// Detect and set large model builds.
+#undef HSA_LARGE_MODEL
+#if defined(__LP64__) || defined(_M_X64)
+#define HSA_LARGE_MODEL
+#endif
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef enum {
+  HSA_STATUS_SUCCESS = 0x0,
+  HSA_STATUS_INFO_BREAK = 0x1,
+  HSA_STATUS_ERROR = 0x1000,
+  HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
+  HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
+} hsa_status_t;
+
+hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);
+
+typedef struct hsa_dim3_s {
+  uint32_t x;
+  uint32_t y;
+  uint32_t z;
+} hsa_dim3_t;
+
+hsa_status_t hsa_init();
+
+hsa_status_t hsa_shut_down();
+
+typedef struct hsa_agent_s {
+  uint64_t handle;
+} hsa_agent_t;
+
+typedef enum {
+  HSA_DEVICE_TYPE_CPU = 0,
+  HSA_DEVICE_TYPE_GPU = 1,
+  HSA_DEVICE_TYPE_DSP = 2
+} hsa_device_type_t;
+
+typedef enum {
+  HSA_AGENT_INFO_NAME = 0,
+  HSA_AGENT_INFO_PROFILE = 4,
+  HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
+  HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
+  HSA_AGENT_INFO_GRID_MAX_DIM = 9,
+  HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
+  HSA_AGENT_INFO_DEVICE = 17,
+} hsa_agent_info_t;
+
+hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
+                                void *value);
+
+hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
+                                                         void *data),
+                                void *data);
+
+typedef struct hsa_signal_s {
+  uint64_t handle;
+} hsa_signal_t;
+
+#ifdef HSA_LARGE_MODEL
+typedef int64_t hsa_signal_value_t;
+#else
+typedef int32_t hsa_signal_value_t;
+#endif
+
+hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
+                               uint32_t num_consumers,
+                               const hsa_agent_t *consumers,
+                               hsa_signal_t *signal);
+
+hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
+
+void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
+
+typedef enum {
+  HSA_SIGNAL_CONDITION_EQ = 0,
+  HSA_SIGNAL_CONDITION_NE = 1,
+} hsa_signal_condition_t;
+
+typedef enum {
+  HSA_WAIT_STATE_BLOCKED = 0,
+  HSA_WAIT_STATE_ACTIVE = 1
+} hsa_wait_state_t;
+
+hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
+                                             hsa_signal_condition_t condition,
+                                             hsa_signal_value_t compare_value,
+                                             uint64_t timeout_hint,
+                                             hsa_wait_state_t wait_state_hint);
+
+typedef enum {
+  HSA_QUEUE_TYPE_MULTI = 0,
+  HSA_QUEUE_TYPE_SINGLE = 1,
+} hsa_queue_type_t;
+
+typedef uint32_t hsa_queue_type32_t;
+
+typedef struct hsa_queue_s {
+  hsa_queue_type32_t type;
+  uint32_t features;
+
+#ifdef HSA_LARGE_MODEL
+  void *base_address;
+#elif defined HSA_LITTLE_ENDIAN
+  void *base_address;
+  uint32_t reserved0;
+#else
+  uint32_t reserved0;
+  void *base_address;
+#endif
+  hsa_signal_t doorbell_signal;
+  uint32_t size;
+  uint32_t reserved1;
+  uint64_t id;
+} hsa_queue_t;
+
+hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size,
+                              hsa_queue_type32_t type,
+                              void (*callback)(hsa_status_t status,
+                                               hsa_queue_t *source, void *data),
+                              void *data, uint32_t private_segment_size,
+                              uint32_t group_segment_size, hsa_queue_t **queue);
+
+uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);
+
+uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
+                                           uint64_t value);
+
+typedef enum {
+  HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
+} hsa_packet_type_t;
+
+typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;
+
+typedef enum {
+  HSA_PACKET_HEADER_TYPE = 0,
+  HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
+  HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
+} hsa_packet_header_t;
+
+typedef enum {
+  HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
+} hsa_kernel_dispatch_packet_setup_t;
+
+typedef enum {
+  HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
+} hsa_kernel_dispatch_packet_setup_width_t;
+
+typedef struct hsa_kernel_dispatch_packet_s {
+  uint16_t header;
+  uint16_t setup;
+  uint16_t workgroup_size_x;
+  uint16_t workgroup_size_y;
+  uint16_t workgroup_size_z;
+  uint16_t reserved0;
+  uint32_t grid_size_x;
+  uint32_t grid_size_y;
+  uint32_t grid_size_z;
+  uint32_t private_segment_size;
+  uint32_t group_segment_size;
+  uint64_t kernel_object;
+#ifdef HSA_LARGE_MODEL
+  void *kernarg_address;
+#elif defined HSA_LITTLE_ENDIAN
+  void *kernarg_address;
+  uint32_t reserved1;
+#else
+  uint32_t reserved1;
+  void *kernarg_address;
+#endif
+  uint64_t reserved2;
+  hsa_signal_t completion_signal;
+} hsa_kernel_dispatch_packet_t;
+
+typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
+
+typedef enum {
+  HSA_EXECUTABLE_STATE_UNFROZEN = 0,
+  HSA_EXECUTABLE_STATE_FROZEN = 1
+} hsa_executable_state_t;
+
+typedef struct hsa_executable_s {
+  uint64_t handle;
+} hsa_executable_t;
+
+typedef struct hsa_executable_symbol_s {
+  uint64_t handle;
+} hsa_executable_symbol_t;
+
+typedef enum {
+  HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
+  HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
+  HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
+  HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
+  HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
+  HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
+  HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
+  HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
+  HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
+} hsa_executable_symbol_info_t;
+
+typedef struct hsa_code_object_s {
+  uint64_t handle;
+} hsa_code_object_t;
+
+typedef enum {
+  HSA_SYMBOL_KIND_VARIABLE = 0,
+  HSA_SYMBOL_KIND_KERNEL = 1,
+  HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
+} hsa_symbol_kind_t;
+
+hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);
+
+hsa_status_t hsa_executable_create(hsa_profile_t profile,
+                                   hsa_executable_state_t executable_state,
+                                   const char *options,
+                                   hsa_executable_t *executable);
+
+hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
+
+hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
+                                   const char *options);
+
+hsa_status_t
+hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
+                               hsa_executable_symbol_info_t attribute,
+                               void *value);
+
+hsa_status_t hsa_executable_iterate_symbols(
+    hsa_executable_t executable,
+    hsa_status_t (*callback)(hsa_executable_t exec,
+                             hsa_executable_symbol_t symbol, void *data),
+    void *data);
+
+hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
+                                         size_t serialized_code_object_size,
+                                         const char *options,
+                                         hsa_code_object_t *code_object);
+
+hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
+                                             hsa_agent_t agent,
+                                             hsa_code_object_t code_object,
+                                             const char *options);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp

+//===--- amdgpu/dynamic_hsa/hsa.cpp ------------------------------- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Implement subset of hsa api by calling into hsa library via dlopen
+// Does the dlopen/dlsym calls as part of the call to hsa_init
+//
+//===----------------------------------------------------------------------===//
+#include "hsa.h"
+#include "Debug.h"
+#include "dlwrap.h"
+#include "hsa_ext_amd.h"
+
+#include <dlfcn.h>
+
+DLWRAP_INTERNAL(hsa_init, 0);
+
+DLWRAP(hsa_status_string, 2);
+DLWRAP(hsa_shut_down, 0);
+DLWRAP(hsa_agent_get_info, 3);
+DLWRAP(hsa_iterate_agents, 2);
+DLWRAP(hsa_signal_create, 4);
+DLWRAP(hsa_signal_destroy, 1);
+DLWRAP(hsa_signal_store_relaxed, 2);
+DLWRAP(hsa_signal_store_screlease, 2);
+DLWRAP(hsa_signal_wait_scacquire, 5);
+DLWRAP(hsa_queue_create, 8);
+DLWRAP(hsa_queue_load_read_index_scacquire, 1);
+DLWRAP(hsa_queue_add_write_index_relaxed, 2);
+DLWRAP(hsa_memory_copy, 3);
+DLWRAP(hsa_executable_create, 4);
+DLWRAP(hsa_executable_destroy, 1);
+DLWRAP(hsa_executable_freeze, 2);
+DLWRAP(hsa_executable_symbol_get_info, 3);
+DLWRAP(hsa_executable_iterate_symbols, 3);
+DLWRAP(hsa_code_object_deserialize, 4);
+DLWRAP(hsa_executable_load_code_object, 4);
+DLWRAP(hsa_amd_agent_memory_pool_get_info, 4);
+DLWRAP(hsa_amd_agent_iterate_memory_pools, 3);
+DLWRAP(hsa_amd_memory_pool_allocate, 4);
+DLWRAP(hsa_amd_memory_pool_free, 1);
+DLWRAP(hsa_amd_memory_async_copy, 8);
+DLWRAP(hsa_amd_memory_pool_get_info, 3);
+DLWRAP(hsa_amd_agents_allow_access, 4);
+DLWRAP(hsa_amd_memory_fill, 3);
+DLWRAP(hsa_amd_register_system_event_handler, 2);
+
+DLWRAP_FINALIZE();
+
+#ifndef DYNAMIC_HSA_PATH
+#define DYNAMIC_HSA_PATH "libhsa-runtime64.so"
+#endif
+
+#ifndef TARGET_NAME
+#define TARGET_NAME AMDHSA
+#endif
+#define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
+
+static bool checkForHSA() {
+  // return true if dlopen succeeded and all functions found
+
+  const char *HsaLib = DYNAMIC_HSA_PATH;
+  void *DynlibHandle = dlopen(HsaLib, RTLD_NOW);
+  if (!DynlibHandle) {
+    DP("Unable to load library '%s': %s!\n", HsaLib, dlerror());
+    return false;
+  }
+
+  for (size_t I = 0; I < dlwrap::size(); I++) {
+    const char *Sym = dlwrap::symbol(I);
+
+    void *P = dlsym(DynlibHandle, Sym);
+    if (P == nullptr) {
+      DP("Unable to find '%s' in '%s'!\n", Sym, HsaLib);
+      return false;
+    }
+    DP("Implementing %s with dlsym(%s) -> %p\n", Sym, Sym, P);
+
+    *dlwrap::pointer(I) = P;
+  }
+
+  return true;
+}
+
+hsa_status_t hsa_init() {
+  if (!checkForHSA()) {
+    return HSA_STATUS_ERROR;
+  }
+  return dlwrap_hsa_init();
+}

openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa_ext_amd.h

+//===--- amdgpu/dynamic_hsa/hsa_ext_amd.h ------------------------- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// The parts of the hsa api that are presently in use by the amdgpu plugin
+//
+//===----------------------------------------------------------------------===//
+#ifndef HSA_RUNTIME_EXT_AMD_H_
+#define HSA_RUNTIME_EXT_AMD_H_
+
+#include "hsa.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef struct hsa_amd_memory_pool_s {
+  uint64_t handle;
+} hsa_amd_memory_pool_t;
+
+typedef enum hsa_amd_memory_pool_global_flag_s {
+  HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT = 1,
+  HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED = 2,
+  HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED = 4
+} hsa_amd_memory_pool_global_flag_t;
+
+typedef enum {
+  HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS = 1,
+  HSA_AMD_MEMORY_POOL_INFO_SIZE = 2,
+  HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED = 5,
+} hsa_amd_memory_pool_info_t;
+
+typedef enum {
+  HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS = 0,
+} hsa_amd_agent_memory_pool_info_t;
+
+typedef enum {
+  HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED = 0,
+} hsa_amd_memory_pool_access_t;
+
+typedef enum hsa_amd_agent_info_s {
+  HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT = 0xA002,
+} hsa_amd_agent_info_t;
+
+hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
+                                          hsa_amd_memory_pool_info_t attribute,
+                                          void *value);
+
+hsa_status_t hsa_amd_agent_iterate_memory_pools(
+    hsa_agent_t agent,
+    hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void *data),
+    void *data);
+
+hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool,
+                                          size_t size, uint32_t flags,
+                                          void **ptr);
+
+hsa_status_t hsa_amd_memory_pool_free(void *ptr);
+
+hsa_status_t hsa_amd_memory_async_copy(void *dst, hsa_agent_t dst_agent,
+                                       const void *src, hsa_agent_t src_agent,
+                                       size_t size, uint32_t num_dep_signals,
+                                       const hsa_signal_t *dep_signals,
+                                       hsa_signal_t completion_signal);
+
+hsa_status_t hsa_amd_agent_memory_pool_get_info(
+    hsa_agent_t agent, hsa_amd_memory_pool_t memory_pool,
+    hsa_amd_agent_memory_pool_info_t attribute, void *value);
+
+hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents,
+                                         const hsa_agent_t *agents,
+                                         const uint32_t *flags,
+                                         const void *ptr);
+
+hsa_status_t hsa_amd_memory_fill(void *ptr, uint32_t value, size_t count);
+
+typedef enum hsa_amd_event_type_s {
+  HSA_AMD_GPU_MEMORY_FAULT_EVENT = 0,
+} hsa_amd_event_type_t;
+
+typedef struct hsa_amd_gpu_memory_fault_info_s {
+  hsa_agent_t agent;
+  uint64_t virtual_address;
+  uint32_t fault_reason_mask;
+} hsa_amd_gpu_memory_fault_info_t;
+
+typedef struct hsa_amd_event_s {
+  hsa_amd_event_type_t event_type;
+  union {
+    hsa_amd_gpu_memory_fault_info_t memory_fault;
+  };
+} hsa_amd_event_t;
+
+typedef hsa_status_t (*hsa_amd_system_event_callback_t)(
+    const hsa_amd_event_t *event, void *data);
+
+hsa_status_t
+hsa_amd_register_system_event_handler(hsa_amd_system_event_callback_t callback,
+                                      void *data);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif