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

   1 //===--- amdgpu/dynamic_hsa/hsa.h --------------------------------- C++ -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 // The parts of the hsa api that are presently in use by the amdgpu plugin
  10 //
  11 //===----------------------------------------------------------------------===//
  12 #ifndef HSA_RUNTIME_INC_HSA_H_
  13 #define HSA_RUNTIME_INC_HSA_H_
  14
  15 #include <stddef.h>
  16 #include <stdint.h>
  17
  18 // Detect and set large model builds.
  19 #undef HSA_LARGE_MODEL
  20 #if defined(__LP64__) || defined(_M_X64)
  21 #define HSA_LARGE_MODEL
  22 #endif
  23
  24 #ifdef __cplusplus
  25 extern "C" {
  26 #endif
  27
  28 typedef enum {
  29   HSA_STATUS_SUCCESS = 0x0,
  30   HSA_STATUS_INFO_BREAK = 0x1,
  31   HSA_STATUS_ERROR = 0x1000,
  32   HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
  33   HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
  34 } hsa_status_t;
  35
  36 hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);
  37
  38 typedef struct hsa_dim3_s {
  39   uint32_t x;
  40   uint32_t y;
  41   uint32_t z;
  42 } hsa_dim3_t;
  43
  44 hsa_status_t hsa_init();
  45
  46 hsa_status_t hsa_shut_down();
  47
  48 typedef struct hsa_agent_s {
  49   uint64_t handle;
  50 } hsa_agent_t;
  51
  52 typedef enum {
  53   HSA_DEVICE_TYPE_CPU = 0,
  54   HSA_DEVICE_TYPE_GPU = 1,
  55   HSA_DEVICE_TYPE_DSP = 2
  56 } hsa_device_type_t;
  57
  58 typedef enum {
  59   HSA_ISA_INFO_NAME_LENGTH = 0,
  60   HSA_ISA_INFO_NAME = 1
  61 } hsa_isa_info_t;
  62
  63 typedef enum {
  64   HSA_AGENT_INFO_NAME = 0,
  65   HSA_AGENT_INFO_VENDOR_NAME = 1,
  66   HSA_AGENT_INFO_FEATURE = 2,
  67   HSA_AGENT_INFO_PROFILE = 4,
  68   HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
  69   HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
  70   HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
  71   HSA_AGENT_INFO_GRID_MAX_DIM = 9,
  72   HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
  73   HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
  74   HSA_AGENT_INFO_QUEUES_MAX = 12,
  75   HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
  76   HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
  77   HSA_AGENT_INFO_NODE = 16,
  78   HSA_AGENT_INFO_DEVICE = 17,
  79   HSA_AGENT_INFO_CACHE_SIZE = 18,
  80   HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
  81 } hsa_agent_info_t;
  82
  83 typedef enum {
  84   HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
  85   HSA_SYSTEM_INFO_VERSION_MINOR = 1,
  86 } hsa_system_info_t;
  87
  88 typedef enum {
  89   HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
  90   HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
  91 } hsa_agent_feature_t;
  92
  93 typedef struct hsa_region_s {
  94   uint64_t handle;
  95 } hsa_region_t;
  96
  97 typedef struct hsa_isa_s {
  98   uint64_t handle;
  99 } hsa_isa_t;
 100
 101 hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);
 102
 103 hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
 104                                 void *value);
 105
 106 hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,
 107                                   void *value);
 108
 109 hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
 110                                                          void *data),
 111                                 void *data);
 112
 113 hsa_status_t hsa_agent_iterate_isas(hsa_agent_t agent,
 114                                     hsa_status_t (*callback)(hsa_isa_t isa,
 115                                                              void *data),
 116                                     void *data);
 117
 118 typedef struct hsa_signal_s {
 119   uint64_t handle;
 120 } hsa_signal_t;
 121
 122 #ifdef HSA_LARGE_MODEL
 123 typedef int64_t hsa_signal_value_t;
 124 #else
 125 typedef int32_t hsa_signal_value_t;
 126 #endif
 127
 128 hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
 129                                uint32_t num_consumers,
 130                                const hsa_agent_t *consumers,
 131                                hsa_signal_t *signal);
 132
 133 hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
 134                                    uint32_t num_consumers,
 135                                    const hsa_agent_t *consumers,
 136                                    uint64_t attributes, hsa_signal_t *signal);
 137
 138 hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
 139
 140 void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
 141
 142 void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
 143
 144 hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);
 145
 146 void hsa_signal_subtract_screlease(hsa_signal_t signal,
 147                                    hsa_signal_value_t value);
 148
 149 typedef enum {
 150   HSA_SIGNAL_CONDITION_EQ = 0,
 151   HSA_SIGNAL_CONDITION_NE = 1,
 152 } hsa_signal_condition_t;
 153
 154 typedef enum {
 155   HSA_WAIT_STATE_BLOCKED = 0,
 156   HSA_WAIT_STATE_ACTIVE = 1
 157 } hsa_wait_state_t;
 158
 159 hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
 160                                              hsa_signal_condition_t condition,
 161                                              hsa_signal_value_t compare_value,
 162                                              uint64_t timeout_hint,
 163                                              hsa_wait_state_t wait_state_hint);
 164
 165 typedef enum {
 166   HSA_QUEUE_TYPE_MULTI = 0,
 167   HSA_QUEUE_TYPE_SINGLE = 1,
 168 } hsa_queue_type_t;
 169
 170 typedef enum {
 171   HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
 172   HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
 173 } hsa_queue_feature_t;
 174
 175 typedef uint32_t hsa_queue_type32_t;
 176
 177 typedef struct hsa_queue_s {
 178   hsa_queue_type32_t type;
 179   uint32_t features;
 180
 181 #ifdef HSA_LARGE_MODEL
 182   void *base_address;
 183 #elif defined HSA_LITTLE_ENDIAN
 184   void *base_address;
 185   uint32_t reserved0;
 186 #else
 187   uint32_t reserved0;
 188   void *base_address;
 189 #endif
 190   hsa_signal_t doorbell_signal;
 191   uint32_t size;
 192   uint32_t reserved1;
 193   uint64_t id;
 194 } hsa_queue_t;
 195
 196 hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size,
 197                               hsa_queue_type32_t type,
 198                               void (*callback)(hsa_status_t status,
 199                                                hsa_queue_t *source, void *data),
 200                               void *data, uint32_t private_segment_size,
 201                               uint32_t group_segment_size, hsa_queue_t **queue);
 202
 203 hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);
 204
 205 uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);
 206
 207 uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
 208                                            uint64_t value);
 209
 210 typedef enum {
 211   HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
 212   HSA_PACKET_TYPE_BARRIER_AND = 3,
 213 } hsa_packet_type_t;
 214
 215 typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;
 216
 217 typedef enum {
 218   HSA_PACKET_HEADER_TYPE = 0,
 219   HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
 220   HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
 221 } hsa_packet_header_t;
 222
 223 typedef enum {
 224   HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
 225 } hsa_kernel_dispatch_packet_setup_t;
 226
 227 typedef enum {
 228   HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
 229 } hsa_kernel_dispatch_packet_setup_width_t;
 230
 231 typedef struct hsa_kernel_dispatch_packet_s {
 232   uint16_t header;
 233   uint16_t setup;
 234   uint16_t workgroup_size_x;
 235   uint16_t workgroup_size_y;
 236   uint16_t workgroup_size_z;
 237   uint16_t reserved0;
 238   uint32_t grid_size_x;
 239   uint32_t grid_size_y;
 240   uint32_t grid_size_z;
 241   uint32_t private_segment_size;
 242   uint32_t group_segment_size;
 243   uint64_t kernel_object;
 244 #ifdef HSA_LARGE_MODEL
 245   void *kernarg_address;
 246 #elif defined HSA_LITTLE_ENDIAN
 247   void *kernarg_address;
 248   uint32_t reserved1;
 249 #else
 250   uint32_t reserved1;
 251   void *kernarg_address;
 252 #endif
 253   uint64_t reserved2;
 254   hsa_signal_t completion_signal;
 255 } hsa_kernel_dispatch_packet_t;
 256
 257 typedef struct hsa_barrier_and_packet_s {
 258   uint16_t header;
 259   uint16_t reserved0;
 260   uint32_t reserved1;
 261   hsa_signal_t dep_signal[5];
 262   uint64_t reserved2;
 263   hsa_signal_t completion_signal;
 264 } hsa_barrier_and_packet_t;
 265
 266 typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
 267
 268 typedef enum {
 269   HSA_EXECUTABLE_STATE_UNFROZEN = 0,
 270   HSA_EXECUTABLE_STATE_FROZEN = 1
 271 } hsa_executable_state_t;
 272
 273 typedef struct hsa_executable_s {
 274   uint64_t handle;
 275 } hsa_executable_t;
 276
 277 typedef struct hsa_executable_symbol_s {
 278   uint64_t handle;
 279 } hsa_executable_symbol_t;
 280
 281 typedef enum {
 282   HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
 283   HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
 284   HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
 285   HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
 286   HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
 287   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
 288   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
 289   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
 290   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
 291 } hsa_executable_symbol_info_t;
 292
 293 typedef struct hsa_code_object_s {
 294   uint64_t handle;
 295 } hsa_code_object_t;
 296
 297 typedef enum {
 298   HSA_SYMBOL_KIND_VARIABLE = 0,
 299   HSA_SYMBOL_KIND_KERNEL = 1,
 300   HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
 301 } hsa_symbol_kind_t;
 302
 303 typedef enum {
 304   HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
 305   HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
 306   HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,
 307 } hsa_default_float_rounding_mode_t;
 308
 309 hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);
 310
 311 hsa_status_t hsa_executable_create(hsa_profile_t profile,
 312                                    hsa_executable_state_t executable_state,
 313                                    const char *options,
 314                                    hsa_executable_t *executable);
 315
 316 hsa_status_t hsa_executable_create_alt(
 317     hsa_profile_t profile,
 318     hsa_default_float_rounding_mode_t default_float_rounding_mode,
 319     const char *options, hsa_executable_t *executable);
 320
 321 hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
 322
 323 hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
 324                                    const char *options);
 325
 326 hsa_status_t hsa_executable_validate(hsa_executable_t executable,
 327                                      uint32_t *result);
 328
 329 hsa_status_t
 330 hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
 331                                hsa_executable_symbol_info_t attribute,
 332                                void *value);
 333
 334 hsa_status_t hsa_executable_iterate_symbols(
 335     hsa_executable_t executable,
 336     hsa_status_t (*callback)(hsa_executable_t exec,
 337                              hsa_executable_symbol_t symbol, void *data),
 338     void *data);
 339
 340 hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,
 341                                                const char *symbol_name,
 342                                                const hsa_agent_t *agent,
 343                                                hsa_executable_symbol_t *symbol);
 344
 345 hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
 346                                          size_t serialized_code_object_size,
 347                                          const char *options,
 348                                          hsa_code_object_t *code_object);
 349
 350 hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
 351                                              hsa_agent_t agent,
 352                                              hsa_code_object_t code_object,
 353                                              const char *options);
 354
 355 hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);
 356
 357 typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);
 358
 359 hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,
 360                                           hsa_signal_condition_t cond,
 361                                           hsa_signal_value_t value,
 362                                           hsa_amd_signal_handler handler,
 363                                           void *arg);
 364
 365 #ifdef __cplusplus
 366 }
 367 #endif
 368
 369 #endif