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_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 = 2
66} hsa_device_type_t;
67
68typedef enum {
69 HSA_ISA_INFO_NAME_LENGTH = 0,
70 HSA_ISA_INFO_NAME = 1
71} hsa_isa_info_t;
72
73typedef 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
93typedef enum {
94 HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
95 HSA_SYSTEM_INFO_VERSION_MINOR = 1,
96} hsa_system_info_t;
97
98typedef enum {
99 HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
100 HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
101} hsa_agent_feature_t;
102
103typedef struct hsa_region_s {
104 uint64_t handle;
105} hsa_region_t;
106
107typedef struct hsa_isa_s {
108 uint64_t handle;
109} hsa_isa_t;
110
111hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);
112
113hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
114 void *value);
115
116hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,
117 void *value);
118
119hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
120 void *data),
121 void *data);
122
123hsa_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
128typedef struct hsa_signal_s {
129 uint64_t handle;
130} hsa_signal_t;
131
132#ifdef HSA_LARGE_MODEL
133typedef int64_t hsa_signal_value_t;
134#else
135typedef int32_t hsa_signal_value_t;
136#endif
137
138hsa_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
143hsa_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
148hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
149
150void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
151
152void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
153
154hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);
155
156void hsa_signal_subtract_screlease(hsa_signal_t signal,
157 hsa_signal_value_t value);
158
159typedef enum {
160 HSA_SIGNAL_CONDITION_EQ = 0,
161 HSA_SIGNAL_CONDITION_NE = 1,
162} hsa_signal_condition_t;
163
164typedef enum {
165 HSA_WAIT_STATE_BLOCKED = 0,
166 HSA_WAIT_STATE_ACTIVE = 1
167} hsa_wait_state_t;
168
169hsa_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
175typedef enum {
176 HSA_QUEUE_TYPE_MULTI = 0,
177 HSA_QUEUE_TYPE_SINGLE = 1,
178} hsa_queue_type_t;
179
180typedef enum {
181 HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
182 HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
183} hsa_queue_feature_t;
184
185typedef uint32_t hsa_queue_type32_t;
186
187typedef 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
206hsa_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
213hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);
214
215uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);
216
217uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
218 uint64_t value);
219
220typedef enum {
221 HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
222 HSA_PACKET_TYPE_BARRIER_AND = 3,
223} hsa_packet_type_t;
224
225typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;
226
227typedef enum {
228 HSA_PACKET_HEADER_TYPE = 0,
229 HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
230 HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
231} hsa_packet_header_t;
232
233typedef enum {
234 HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
235} hsa_kernel_dispatch_packet_setup_t;
236
237typedef enum {
238 HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
239} hsa_kernel_dispatch_packet_setup_width_t;
240
241typedef struct hsa_kernel_dispatch_packet_s {
242 uint16_t header;
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
267typedef 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
276typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
277
278typedef enum {
279 HSA_EXECUTABLE_STATE_UNFROZEN = 0,
280 HSA_EXECUTABLE_STATE_FROZEN = 1
281} hsa_executable_state_t;
282
283typedef struct hsa_executable_s {
284 uint64_t handle;
285} hsa_executable_t;
286
287typedef struct hsa_executable_symbol_s {
288 uint64_t handle;
289} hsa_executable_symbol_t;
290
291typedef 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
304typedef struct hsa_code_object_s {
305 uint64_t handle;
306} hsa_code_object_t;
307
308typedef 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
314typedef 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
320hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);
321
322hsa_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
327hsa_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
332hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
333
334hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
335 const char *options);
336
337hsa_status_t hsa_executable_validate(hsa_executable_t executable,
338 uint32_t *result);
339
340hsa_status_t
341hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
342 hsa_executable_symbol_info_t attribute,
343 void *value);
344
345hsa_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
351hsa_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
356hsa_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
361hsa_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
366hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);
367
368typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);
369
370hsa_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
376hsa_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
380hsa_status_t
381hsa_code_object_reader_destroy(hsa_code_object_reader_t code_object_reader);
382
383hsa_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

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