[OpenMP][libomptarget] Add missing symbols in dynamic_hsa

[kevinsala]

Summary

This patch prepares for the new AMDGPU NextGen plugin.

[jhuber6]

If something has a simple uncontroversial fix, it's fine to just push it directly without review. Usually easier than reverting a whole patch.

[arsenm]

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

306

Is there a reason this is missing one?

[kevinsala]

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

306

That's used in https://reviews.llvm.org/D138389

[jhuber6]

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

306

typedef enum {                                                                                                                                                                                                    
  /**                                                                                                                                                                                                             
   * Use a default floating-point rounding mode specified elsewhere.                                                                                                                                              
   */                                                                                                                                                                                                             
  HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,    
  /**                                                                                                                                                                                                             
   * Operations that specify the default floating-point mode are rounded to zero                                                                                                                                  
   * by default.                                                                                                                                                                                                  
   */                                                                                                                                                                                                             
  HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,                                                                                                                                                                       
  /**                                                                                                                                                                                                             
   * Operations that specify the default floating-point mode are rounded to the                                                                                                                                   
   * nearest representable number and that ties should be broken by selecting                                                                                                                                     
   * the value with an even least significant bit.                                                                                                                                                                
   */                                                                                                                                                                                                             
  HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2                                                                                                                                                                        
} hsa_default_float_rounding_mode_t;

That's all I see in my installation, unless this was changed recently.

[arsenm]

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

306

There are 4 IEEE rounding modes and the hardware supports all of them so this is suspicious. Also "default" = near, so really this is missing +/- infinity

[jhuber6]

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

306

That's a good point, Clang defined the following, which matches what I just looked up at https://en.wikipedia.org/wiki/IEEE_754#Rounding_rules.

TowardZero        = 0,    ///< roundTowardZero.                                                                                                                                     
NearestTiesToEven = 1,    ///< roundTiesToEven.                                                                                                                                           
TowardPositive    = 2,    ///< roundTowardPositive.                                                                                                                      
TowardNegative    = 3,    ///< roundTowardNegative.                                                                                                          
NearestTiesToAway = 4,    ///< roundTiesToAway.

These three are all that's included in hsa.h for when we load the executable, not sure who would be the person to bring that up to, maybe @JonChesterfield has an idea. Maybe they're combining the +/- and nearest rounding modes somehow.

In any case this is unrelated to this patch as it just copied what hsa.h gives it.

[JonChesterfield]

This header file is a compromise between including enough of hsa.h to compile without the rocm source and not including all of it because it's large. Different people have taken different approaches to removing unused parts. There's a similar cuda.h near nvptx.

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

 typedef enum {
   HSA_ISA_INFO_NAME_LENGTH = 0,
   HSA_ISA_INFO_NAME = 1
 } hsa_isa_info_t;
 
 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,
   HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
   HSA_AGENT_INFO_GRID_MAX_DIM = 9,
   HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
   HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
   HSA_AGENT_INFO_QUEUES_MAX = 12,
   HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
   HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
   HSA_AGENT_INFO_DEVICE = 17,
   HSA_AGENT_INFO_CACHE_SIZE = 18,
   HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
 } hsa_agent_info_t;
 
 typedef enum {
   HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
   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;
 
 typedef struct hsa_isa_s {
   uint64_t handle;
 } hsa_isa_t;
 

 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_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,
 } 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 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 {
   hsa_queue_type32_t type;
   uint32_t features;
 
 #ifdef HSA_LARGE_MODEL
   void *base_address;

 
 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_BARRIER_AND = 3,
 } 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

 #else
   uint32_t reserved1;
   void *kernarg_address;
 #endif
   uint64_t reserved2;
   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 {
   HSA_EXECUTABLE_STATE_UNFROZEN = 0,
   HSA_EXECUTABLE_STATE_FROZEN = 1
 } hsa_executable_state_t;
 
 typedef struct hsa_executable_s {

 } 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;
 
+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,
                                    hsa_executable_state_t executable_state,
                                    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,
                                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_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,
                                          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);
 
+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
 
 #endif

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

 DLWRAP(hsa_isa_get_info_alt, 3)
 DLWRAP(hsa_iterate_agents, 2)
 DLWRAP(hsa_agent_iterate_isas, 3)
 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_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)
 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_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()
 
 #ifndef DYNAMIC_HSA_PATH
 #define DYNAMIC_HSA_PATH "libhsa-runtime64.so"
 #endif
 
 #ifndef TARGET_NAME