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_AGENT_INFO_NAME = 0, 60 HSA_AGENT_INFO_PROFILE = 4, 61 HSA_AGENT_INFO_WAVEFRONT_SIZE = 6, 62 HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7, 63 HSA_AGENT_INFO_GRID_MAX_DIM = 9, 64 HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14, 65 HSA_AGENT_INFO_DEVICE = 17, 66 } hsa_agent_info_t; 67 68 hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute, 69 void *value); 70 71 hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, 72 void *data), 73 void *data); 74 75 typedef struct hsa_signal_s { 76 uint64_t handle; 77 } hsa_signal_t; 78 79 #ifdef HSA_LARGE_MODEL 80 typedef int64_t hsa_signal_value_t; 81 #else 82 typedef int32_t hsa_signal_value_t; 83 #endif 84 85 hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value, 86 uint32_t num_consumers, 87 const hsa_agent_t *consumers, 88 hsa_signal_t *signal); 89 90 hsa_status_t hsa_signal_destroy(hsa_signal_t signal); 91 92 void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value); 93 94 void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value); 95 96 typedef enum { 97 HSA_SIGNAL_CONDITION_EQ = 0, 98 HSA_SIGNAL_CONDITION_NE = 1, 99 } hsa_signal_condition_t; 100 101 typedef enum { 102 HSA_WAIT_STATE_BLOCKED = 0, 103 HSA_WAIT_STATE_ACTIVE = 1 104 } hsa_wait_state_t; 105 106 hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal, 107 hsa_signal_condition_t condition, 108 hsa_signal_value_t compare_value, 109 uint64_t timeout_hint, 110 hsa_wait_state_t wait_state_hint); 111 112 typedef enum { 113 HSA_QUEUE_TYPE_MULTI = 0, 114 HSA_QUEUE_TYPE_SINGLE = 1, 115 } hsa_queue_type_t; 116 117 typedef uint32_t hsa_queue_type32_t; 118 119 typedef struct hsa_queue_s { 120 hsa_queue_type32_t type; 121 uint32_t features; 122 123 #ifdef HSA_LARGE_MODEL 124 void *base_address; 125 #elif defined HSA_LITTLE_ENDIAN 126 void *base_address; 127 uint32_t reserved0; 128 #else 129 uint32_t reserved0; 130 void *base_address; 131 #endif 132 hsa_signal_t doorbell_signal; 133 uint32_t size; 134 uint32_t reserved1; 135 uint64_t id; 136 } hsa_queue_t; 137 138 hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size, 139 hsa_queue_type32_t type, 140 void (*callback)(hsa_status_t status, 141 hsa_queue_t *source, void *data), 142 void *data, uint32_t private_segment_size, 143 uint32_t group_segment_size, hsa_queue_t **queue); 144 145 uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue); 146 147 uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue, 148 uint64_t value); 149 150 typedef enum { 151 HSA_PACKET_TYPE_KERNEL_DISPATCH = 2, 152 } hsa_packet_type_t; 153 154 typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t; 155 156 typedef enum { 157 HSA_PACKET_HEADER_TYPE = 0, 158 HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9, 159 HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11 160 } hsa_packet_header_t; 161 162 typedef enum { 163 HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0 164 } hsa_kernel_dispatch_packet_setup_t; 165 166 typedef enum { 167 HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2 168 } hsa_kernel_dispatch_packet_setup_width_t; 169 170 typedef struct hsa_kernel_dispatch_packet_s { 171 uint16_t header; 172 uint16_t setup; 173 uint16_t workgroup_size_x; 174 uint16_t workgroup_size_y; 175 uint16_t workgroup_size_z; 176 uint16_t reserved0; 177 uint32_t grid_size_x; 178 uint32_t grid_size_y; 179 uint32_t grid_size_z; 180 uint32_t private_segment_size; 181 uint32_t group_segment_size; 182 uint64_t kernel_object; 183 #ifdef HSA_LARGE_MODEL 184 void *kernarg_address; 185 #elif defined HSA_LITTLE_ENDIAN 186 void *kernarg_address; 187 uint32_t reserved1; 188 #else 189 uint32_t reserved1; 190 void *kernarg_address; 191 #endif 192 uint64_t reserved2; 193 hsa_signal_t completion_signal; 194 } hsa_kernel_dispatch_packet_t; 195 196 typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t; 197 198 typedef enum { 199 HSA_EXECUTABLE_STATE_UNFROZEN = 0, 200 HSA_EXECUTABLE_STATE_FROZEN = 1 201 } hsa_executable_state_t; 202 203 typedef struct hsa_executable_s { 204 uint64_t handle; 205 } hsa_executable_t; 206 207 typedef struct hsa_executable_symbol_s { 208 uint64_t handle; 209 } hsa_executable_symbol_t; 210 211 typedef enum { 212 HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0, 213 HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1, 214 HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2, 215 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21, 216 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9, 217 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22, 218 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11, 219 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13, 220 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14, 221 } hsa_executable_symbol_info_t; 222 223 typedef struct hsa_code_object_s { 224 uint64_t handle; 225 } hsa_code_object_t; 226 227 typedef enum { 228 HSA_SYMBOL_KIND_VARIABLE = 0, 229 HSA_SYMBOL_KIND_KERNEL = 1, 230 HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2 231 } hsa_symbol_kind_t; 232 233 hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size); 234 235 hsa_status_t hsa_executable_create(hsa_profile_t profile, 236 hsa_executable_state_t executable_state, 237 const char *options, 238 hsa_executable_t *executable); 239 240 hsa_status_t hsa_executable_destroy(hsa_executable_t executable); 241 242 hsa_status_t hsa_executable_freeze(hsa_executable_t executable, 243 const char *options); 244 245 hsa_status_t 246 hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol, 247 hsa_executable_symbol_info_t attribute, 248 void *value); 249 250 hsa_status_t hsa_executable_iterate_symbols( 251 hsa_executable_t executable, 252 hsa_status_t (*callback)(hsa_executable_t exec, 253 hsa_executable_symbol_t symbol, void *data), 254 void *data); 255 256 hsa_status_t hsa_code_object_deserialize(void *serialized_code_object, 257 size_t serialized_code_object_size, 258 const char *options, 259 hsa_code_object_t *code_object); 260 261 hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable, 262 hsa_agent_t agent, 263 hsa_code_object_t code_object, 264 const char *options); 265 266 #ifdef __cplusplus 267 } 268 #endif 269 270 #endif 271