399 lines · c
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-exception6//7//===----------------------------------------------------------------------===//8//9// The parts of the hsa api that are presently in use by the amdgpu plugin10//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_MODEL20#if defined(__LP64__) || defined(_M_X64)21#define HSA_LARGE_MODEL22#endif23 24#ifdef __cplusplus25extern "C" {26#endif27 28typedef 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_INVALID_SYMBOL_NAME = 0x1013,34 HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,35 HSA_STATUS_ERROR_EXCEPTION = 0x1016,36} hsa_status_t;37 38hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);39 40typedef struct hsa_dim3_s {41 uint32_t x;42 uint32_t y;43 uint32_t z;44} hsa_dim3_t;45 46hsa_status_t hsa_init();47 48hsa_status_t hsa_shut_down();49 50typedef struct hsa_agent_s {51 uint64_t handle;52} hsa_agent_t;53 54typedef struct hsa_loaded_code_object_s {55 uint64_t handle;56} hsa_loaded_code_object_t;57 58typedef struct hsa_code_object_reader_s {59 uint64_t handle;60} hsa_code_object_reader_t;61 62typedef enum {63 HSA_DEVICE_TYPE_CPU = 0,64 HSA_DEVICE_TYPE_GPU = 1,65 HSA_DEVICE_TYPE_DSP = 266} hsa_device_type_t;67 68typedef enum {69 HSA_ISA_INFO_NAME_LENGTH = 0,70 HSA_ISA_INFO_NAME = 171} hsa_isa_info_t;72 73typedef enum {74 HSA_MACHINE_MODEL_SMALL = 0,75 HSA_MACHINE_MODEL_LARGE = 176} hsa_machine_model_t;77 78typedef enum {79 HSA_AGENT_INFO_NAME = 0,80 HSA_AGENT_INFO_VENDOR_NAME = 1,81 HSA_AGENT_INFO_FEATURE = 2,82 HSA_AGENT_INFO_MACHINE_MODEL = 3,83 HSA_AGENT_INFO_PROFILE = 4,84 HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,85 HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,86 HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,87 HSA_AGENT_INFO_GRID_MAX_DIM = 9,88 HSA_AGENT_INFO_GRID_MAX_SIZE = 10,89 HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,90 HSA_AGENT_INFO_QUEUES_MAX = 12,91 HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,92 HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,93 HSA_AGENT_INFO_NODE = 16,94 HSA_AGENT_INFO_DEVICE = 17,95 HSA_AGENT_INFO_CACHE_SIZE = 18,96 HSA_AGENT_INFO_FAST_F16_OPERATION = 24,97} hsa_agent_info_t;98 99typedef enum {100 HSA_SYSTEM_INFO_VERSION_MAJOR = 0,101 HSA_SYSTEM_INFO_VERSION_MINOR = 1,102} hsa_system_info_t;103 104typedef enum {105 HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,106 HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,107} hsa_agent_feature_t;108 109typedef struct hsa_region_s {110 uint64_t handle;111} hsa_region_t;112 113typedef struct hsa_isa_s {114 uint64_t handle;115} hsa_isa_t;116 117hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);118 119hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,120 void *value);121 122hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,123 void *value);124 125hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,126 void *data),127 void *data);128 129hsa_status_t hsa_agent_iterate_isas(hsa_agent_t agent,130 hsa_status_t (*callback)(hsa_isa_t isa,131 void *data),132 void *data);133 134typedef struct hsa_signal_s {135 uint64_t handle;136} hsa_signal_t;137 138#ifdef HSA_LARGE_MODEL139typedef int64_t hsa_signal_value_t;140#else141typedef int32_t hsa_signal_value_t;142#endif143 144hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,145 uint32_t num_consumers,146 const hsa_agent_t *consumers,147 hsa_signal_t *signal);148 149hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,150 uint32_t num_consumers,151 const hsa_agent_t *consumers,152 uint64_t attributes, hsa_signal_t *signal);153 154hsa_status_t hsa_signal_destroy(hsa_signal_t signal);155 156void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);157 158void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);159 160hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);161 162void hsa_signal_subtract_screlease(hsa_signal_t signal,163 hsa_signal_value_t value);164 165typedef enum {166 HSA_SIGNAL_CONDITION_EQ = 0,167 HSA_SIGNAL_CONDITION_NE = 1,168} hsa_signal_condition_t;169 170typedef enum {171 HSA_WAIT_STATE_BLOCKED = 0,172 HSA_WAIT_STATE_ACTIVE = 1173} hsa_wait_state_t;174 175hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,176 hsa_signal_condition_t condition,177 hsa_signal_value_t compare_value,178 uint64_t timeout_hint,179 hsa_wait_state_t wait_state_hint);180 181typedef enum {182 HSA_QUEUE_TYPE_MULTI = 0,183 HSA_QUEUE_TYPE_SINGLE = 1,184} hsa_queue_type_t;185 186typedef enum {187 HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,188 HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2189} hsa_queue_feature_t;190 191typedef uint32_t hsa_queue_type32_t;192 193typedef struct hsa_queue_s {194 hsa_queue_type32_t type;195 uint32_t features;196 197#ifdef HSA_LARGE_MODEL198 void *base_address;199#elif defined HSA_LITTLE_ENDIAN200 void *base_address;201 uint32_t reserved0;202#else203 uint32_t reserved0;204 void *base_address;205#endif206 hsa_signal_t doorbell_signal;207 uint32_t size;208 uint32_t reserved1;209 uint64_t id;210} hsa_queue_t;211 212hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size,213 hsa_queue_type32_t type,214 void (*callback)(hsa_status_t status,215 hsa_queue_t *source, void *data),216 void *data, uint32_t private_segment_size,217 uint32_t group_segment_size, hsa_queue_t **queue);218 219hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);220 221uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);222 223uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,224 uint64_t value);225 226typedef enum {227 HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,228 HSA_PACKET_TYPE_BARRIER_AND = 3,229} hsa_packet_type_t;230 231typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;232 233typedef enum {234 HSA_PACKET_HEADER_TYPE = 0,235 HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,236 HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11237} hsa_packet_header_t;238 239typedef enum {240 HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0241} hsa_kernel_dispatch_packet_setup_t;242 243typedef enum {244 HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2245} hsa_kernel_dispatch_packet_setup_width_t;246 247typedef struct hsa_kernel_dispatch_packet_s {248 uint16_t header;249 uint16_t setup;250 uint16_t workgroup_size_x;251 uint16_t workgroup_size_y;252 uint16_t workgroup_size_z;253 uint16_t reserved0;254 uint32_t grid_size_x;255 uint32_t grid_size_y;256 uint32_t grid_size_z;257 uint32_t private_segment_size;258 uint32_t group_segment_size;259 uint64_t kernel_object;260#ifdef HSA_LARGE_MODEL261 void *kernarg_address;262#elif defined HSA_LITTLE_ENDIAN263 void *kernarg_address;264 uint32_t reserved1;265#else266 uint32_t reserved1;267 void *kernarg_address;268#endif269 uint64_t reserved2;270 hsa_signal_t completion_signal;271} hsa_kernel_dispatch_packet_t;272 273typedef struct hsa_barrier_and_packet_s {274 uint16_t header;275 uint16_t reserved0;276 uint32_t reserved1;277 hsa_signal_t dep_signal[5];278 uint64_t reserved2;279 hsa_signal_t completion_signal;280} hsa_barrier_and_packet_t;281 282typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;283 284typedef enum {285 HSA_EXECUTABLE_STATE_UNFROZEN = 0,286 HSA_EXECUTABLE_STATE_FROZEN = 1287} hsa_executable_state_t;288 289typedef struct hsa_executable_s {290 uint64_t handle;291} hsa_executable_t;292 293typedef struct hsa_executable_symbol_s {294 uint64_t handle;295} hsa_executable_symbol_t;296 297typedef enum {298 HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,299 HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,300 HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,301 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,302 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,303 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,304 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,305 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,306 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,307 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,308} hsa_executable_symbol_info_t;309 310typedef struct hsa_code_object_s {311 uint64_t handle;312} hsa_code_object_t;313 314typedef enum {315 HSA_SYMBOL_KIND_VARIABLE = 0,316 HSA_SYMBOL_KIND_KERNEL = 1,317 HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2318} hsa_symbol_kind_t;319 320typedef enum {321 HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,322 HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,323 HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,324} hsa_default_float_rounding_mode_t;325 326hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);327 328hsa_status_t hsa_executable_create(hsa_profile_t profile,329 hsa_executable_state_t executable_state,330 const char *options,331 hsa_executable_t *executable);332 333hsa_status_t hsa_executable_create_alt(334 hsa_profile_t profile,335 hsa_default_float_rounding_mode_t default_float_rounding_mode,336 const char *options, hsa_executable_t *executable);337 338hsa_status_t hsa_executable_destroy(hsa_executable_t executable);339 340hsa_status_t hsa_executable_freeze(hsa_executable_t executable,341 const char *options);342 343hsa_status_t hsa_executable_validate(hsa_executable_t executable,344 uint32_t *result);345 346hsa_status_t347hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,348 hsa_executable_symbol_info_t attribute,349 void *value);350 351hsa_status_t hsa_executable_iterate_symbols(352 hsa_executable_t executable,353 hsa_status_t (*callback)(hsa_executable_t exec,354 hsa_executable_symbol_t symbol, void *data),355 void *data);356 357hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,358 const char *symbol_name,359 const hsa_agent_t *agent,360 hsa_executable_symbol_t *symbol);361 362hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,363 size_t serialized_code_object_size,364 const char *options,365 hsa_code_object_t *code_object);366 367hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,368 hsa_agent_t agent,369 hsa_code_object_t code_object,370 const char *options);371 372hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);373 374typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);375 376hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,377 hsa_signal_condition_t cond,378 hsa_signal_value_t value,379 hsa_amd_signal_handler handler,380 void *arg);381 382hsa_status_t hsa_code_object_reader_create_from_memory(383 const void *code_object, size_t size,384 hsa_code_object_reader_t *code_object_reader);385 386hsa_status_t387hsa_code_object_reader_destroy(hsa_code_object_reader_t code_object_reader);388 389hsa_status_t hsa_executable_load_agent_code_object(390 hsa_executable_t executable, hsa_agent_t agent,391 hsa_code_object_reader_t code_object_reader, const char *options,392 hsa_loaded_code_object_t *loaded_code_object);393 394#ifdef __cplusplus395}396#endif397 398#endif399