[Openmp-commits] [openmp] 7b97941 - [OpenMP][libomptarget] Add missing symbols in dynamic_hsa
Kevin Sala via Openmp-commits
openmp-commits at lists.llvm.org
Fri Dec 16 15:01:52 PST 2022
Author: Kevin Sala
Date: 2022-12-17T00:01:24+01:00
New Revision: 7b979417216f271fe36df50ed6a3480d5957a127
URL: https://github.com/llvm/llvm-project/commit/7b979417216f271fe36df50ed6a3480d5957a127
DIFF: https://github.com/llvm/llvm-project/commit/7b979417216f271fe36df50ed6a3480d5957a127.diff
LOG: [OpenMP][libomptarget] Add missing symbols in dynamic_hsa
This patch prepares for the new AMDGPU NextGen plugin.
Differential Revision: https://reviews.llvm.org/D140213
Added:
Modified:
openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
Removed:
################################################################################
diff --git a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
index 6b74b0a98e8ab..c43db066ae6ab 100644
--- a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
@@ -35,18 +35,24 @@ 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_signal_load_scacquire, 1)
+DLWRAP(hsa_signal_subtract_screlease, 2)
DLWRAP(hsa_queue_create, 8)
DLWRAP(hsa_queue_destroy, 1)
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_create_alt, 4)
DLWRAP(hsa_executable_destroy, 1)
DLWRAP(hsa_executable_freeze, 2)
+DLWRAP(hsa_executable_validate, 2)
DLWRAP(hsa_executable_symbol_get_info, 3)
+DLWRAP(hsa_executable_get_symbol_by_name, 4)
DLWRAP(hsa_executable_iterate_symbols, 3)
DLWRAP(hsa_code_object_deserialize, 4)
DLWRAP(hsa_executable_load_code_object, 4)
+DLWRAP(hsa_code_object_destroy, 1)
DLWRAP(hsa_amd_agent_memory_pool_get_info, 4)
DLWRAP(hsa_amd_agent_iterate_memory_pools, 3)
DLWRAP(hsa_amd_memory_pool_allocate, 4)
@@ -58,6 +64,8 @@ DLWRAP(hsa_amd_memory_lock, 5)
DLWRAP(hsa_amd_memory_unlock, 1)
DLWRAP(hsa_amd_memory_fill, 3)
DLWRAP(hsa_amd_register_system_event_handler, 2)
+DLWRAP(hsa_amd_signal_create, 5)
+DLWRAP(hsa_amd_signal_async_handler, 5)
DLWRAP_FINALIZE()
diff --git a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
index 8627860aef089..b57590820a55d 100644
--- a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
+++ b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
@@ -63,6 +63,7 @@ typedef enum {
typedef enum {
HSA_AGENT_INFO_NAME = 0,
HSA_AGENT_INFO_VENDOR_NAME = 1,
+ HSA_AGENT_INFO_FEATURE = 2,
HSA_AGENT_INFO_PROFILE = 4,
HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
@@ -83,6 +84,11 @@ typedef enum {
HSA_SYSTEM_INFO_VERSION_MINOR = 1,
} hsa_system_info_t;
+typedef enum {
+ HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
+ HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
+} hsa_agent_feature_t;
+
typedef struct hsa_region_s {
uint64_t handle;
} hsa_region_t;
@@ -123,12 +129,22 @@ hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
const hsa_agent_t *consumers,
hsa_signal_t *signal);
+hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
+ uint32_t num_consumers,
+ const hsa_agent_t *consumers,
+ uint64_t attributes, 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);
+hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);
+
+void hsa_signal_subtract_screlease(hsa_signal_t signal,
+ hsa_signal_value_t value);
+
typedef enum {
HSA_SIGNAL_CONDITION_EQ = 0,
HSA_SIGNAL_CONDITION_NE = 1,
@@ -150,6 +166,11 @@ typedef enum {
HSA_QUEUE_TYPE_SINGLE = 1,
} hsa_queue_type_t;
+typedef enum {
+ HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
+ HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
+} hsa_queue_feature_t;
+
typedef uint32_t hsa_queue_type32_t;
typedef struct hsa_queue_s {
@@ -187,6 +208,7 @@ uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
typedef enum {
HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
+ HSA_PACKET_TYPE_BARRIER_AND = 3,
} hsa_packet_type_t;
typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;
@@ -231,6 +253,15 @@ typedef struct hsa_kernel_dispatch_packet_s {
hsa_signal_t completion_signal;
} hsa_kernel_dispatch_packet_t;
+typedef struct hsa_barrier_and_packet_s {
+ uint16_t header;
+ uint16_t reserved0;
+ uint32_t reserved1;
+ hsa_signal_t dep_signal[5];
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_barrier_and_packet_t;
+
typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
typedef enum {
@@ -268,6 +299,12 @@ typedef enum {
HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
} hsa_symbol_kind_t;
+typedef enum {
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,
+} hsa_default_float_rounding_mode_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,
@@ -275,11 +312,19 @@ hsa_status_t hsa_executable_create(hsa_profile_t profile,
const char *options,
hsa_executable_t *executable);
+hsa_status_t hsa_executable_create_alt(
+ hsa_profile_t profile,
+ hsa_default_float_rounding_mode_t default_float_rounding_mode,
+ 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_validate(hsa_executable_t executable,
+ uint32_t *result);
+
hsa_status_t
hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
hsa_executable_symbol_info_t attribute,
@@ -291,6 +336,11 @@ hsa_status_t hsa_executable_iterate_symbols(
hsa_executable_symbol_t symbol, void *data),
void *data);
+hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,
+ const char *symbol_name,
+ const hsa_agent_t *agent,
+ hsa_executable_symbol_t *symbol);
+
hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
size_t serialized_code_object_size,
const char *options,
@@ -301,6 +351,16 @@ hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
hsa_code_object_t code_object,
const char *options);
+hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);
+
+typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);
+
+hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,
+ hsa_signal_condition_t cond,
+ hsa_signal_value_t value,
+ hsa_amd_signal_handler handler,
+ void *arg);
+
#ifdef __cplusplus
}
#endif
More information about the Openmp-commits
mailing list