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