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