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