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
25extern "C" {
26#endif
27
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_NOT_INITIALIZED = 0x100B,
34} hsa_status_t;
35
36hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);
37
38typedef struct hsa_dim3_s {
39 uint32_t x;
40 uint32_t y;
41 uint32_t z;
42} hsa_dim3_t;
43
44hsa_status_t hsa_init();
45
46hsa_status_t hsa_shut_down();
47
48typedef struct hsa_agent_s {
49 uint64_t handle;
50} hsa_agent_t;
51
52typedef 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
58typedef enum {
59 HSA_ISA_INFO_NAME_LENGTH = 0,
60 HSA_ISA_INFO_NAME = 1
61} hsa_isa_info_t;
62
63typedef 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
83typedef enum {
84 HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
85 HSA_SYSTEM_INFO_VERSION_MINOR = 1,
86} hsa_system_info_t;
87
88typedef enum {
89 HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
90 HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
91} hsa_agent_feature_t;
92
93typedef struct hsa_region_s {
94 uint64_t handle;
95} hsa_region_t;
96
97typedef struct hsa_isa_s {
98 uint64_t handle;
99} hsa_isa_t;
100
101hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);
102
103hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
104 void *value);
105
106hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,
107 void *value);
108
109hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
110 void *data),
111 void *data);
112
113hsa_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
118typedef struct hsa_signal_s {
119 uint64_t handle;
120} hsa_signal_t;
121
122#ifdef HSA_LARGE_MODEL
123typedef int64_t hsa_signal_value_t;
124#else
125typedef int32_t hsa_signal_value_t;
126#endif
127
128hsa_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
133hsa_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
138hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
139
140void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
141
142void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
143
144hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);
145
146void hsa_signal_subtract_screlease(hsa_signal_t signal,
147 hsa_signal_value_t value);
148
149typedef enum {
150 HSA_SIGNAL_CONDITION_EQ = 0,
151 HSA_SIGNAL_CONDITION_NE = 1,
152} hsa_signal_condition_t;
153
154typedef enum {
155 HSA_WAIT_STATE_BLOCKED = 0,
156 HSA_WAIT_STATE_ACTIVE = 1
157} hsa_wait_state_t;
158
159hsa_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
165typedef enum {
166 HSA_QUEUE_TYPE_MULTI = 0,
167 HSA_QUEUE_TYPE_SINGLE = 1,
168} hsa_queue_type_t;
169
170typedef enum {
171 HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
172 HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
173} hsa_queue_feature_t;
174
175typedef uint32_t hsa_queue_type32_t;
176
177typedef 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
196hsa_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
203hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);
204
205uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);
206
207uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
208 uint64_t value);
209
210typedef enum {
211 HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
212 HSA_PACKET_TYPE_BARRIER_AND = 3,
213} hsa_packet_type_t;
214
215typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;
216
217typedef 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
223typedef enum {
224 HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
225} hsa_kernel_dispatch_packet_setup_t;
226
227typedef enum {
228 HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
229} hsa_kernel_dispatch_packet_setup_width_t;
230
231typedef 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
257typedef 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
266typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
267
268typedef enum {
269 HSA_EXECUTABLE_STATE_UNFROZEN = 0,
270 HSA_EXECUTABLE_STATE_FROZEN = 1
271} hsa_executable_state_t;
272
273typedef struct hsa_executable_s {
274 uint64_t handle;
275} hsa_executable_t;
276
277typedef struct hsa_executable_symbol_s {
278 uint64_t handle;
279} hsa_executable_symbol_t;
280
281typedef 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_KERNEL_DYNAMIC_CALLSTACK = 15,
292} hsa_executable_symbol_info_t;
293
294typedef struct hsa_code_object_s {
295 uint64_t handle;
296} hsa_code_object_t;
297
298typedef enum {
299 HSA_SYMBOL_KIND_VARIABLE = 0,
300 HSA_SYMBOL_KIND_KERNEL = 1,
301 HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
302} hsa_symbol_kind_t;
303
304typedef enum {
305 HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
306 HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
307 HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,
308} hsa_default_float_rounding_mode_t;
309
310hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);
311
312hsa_status_t hsa_executable_create(hsa_profile_t profile,
313 hsa_executable_state_t executable_state,
314 const char *options,
315 hsa_executable_t *executable);
316
317hsa_status_t hsa_executable_create_alt(
318 hsa_profile_t profile,
319 hsa_default_float_rounding_mode_t default_float_rounding_mode,
320 const char *options, hsa_executable_t *executable);
321
322hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
323
324hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
325 const char *options);
326
327hsa_status_t hsa_executable_validate(hsa_executable_t executable,
328 uint32_t *result);
329
330hsa_status_t
331hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
332 hsa_executable_symbol_info_t attribute,
333 void *value);
334
335hsa_status_t hsa_executable_iterate_symbols(
336 hsa_executable_t executable,
337 hsa_status_t (*callback)(hsa_executable_t exec,
338 hsa_executable_symbol_t symbol, void *data),
339 void *data);
340
341hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,
342 const char *symbol_name,
343 const hsa_agent_t *agent,
344 hsa_executable_symbol_t *symbol);
345
346hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
347 size_t serialized_code_object_size,
348 const char *options,
349 hsa_code_object_t *code_object);
350
351hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
352 hsa_agent_t agent,
353 hsa_code_object_t code_object,
354 const char *options);
355
356hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);
357
358typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);
359
360hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,
361 hsa_signal_condition_t cond,
362 hsa_signal_value_t value,
363 hsa_amd_signal_handler handler,
364 void *arg);
365
366#ifdef __cplusplus
367}
368#endif
369
370#endif
371

source code of offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h