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