1 /*===--------------------------------------------------------------------------
2  *              ATMI (Asynchronous Task and Memory Interface)
3  *
4  * This file is distributed under the MIT License. See LICENSE.txt for details.
5  *===------------------------------------------------------------------------*/
6 #ifndef SRC_RUNTIME_INCLUDE_INTERNAL_H_
7 #define SRC_RUNTIME_INCLUDE_INTERNAL_H_
8 #include <inttypes.h>
9 #include <pthread.h>
10 #include <stddef.h>
11 #include <stdint.h>
12 #include <stdio.h>
13 #include <stdlib.h>
14 
15 #include <atomic>
16 #include <cstring>
17 #include <deque>
18 #include <map>
19 #include <queue>
20 #include <string>
21 #include <utility>
22 #include <vector>
23 
24 #include "hsa.h"
25 #include "hsa_ext_amd.h"
26 #include "hsa_ext_finalize.h"
27 
28 #include "atmi.h"
29 #include "atmi_runtime.h"
30 #include "rt.h"
31 
32 #define MAX_NUM_KERNELS (1024 * 16)
33 
34 typedef struct atmi_implicit_args_s {
35   unsigned long offset_x;
36   unsigned long offset_y;
37   unsigned long offset_z;
38   unsigned long hostcall_ptr;
39   char num_gpu_queues;
40   unsigned long gpu_queue_ptr;
41   char num_cpu_queues;
42   unsigned long cpu_worker_signals;
43   unsigned long cpu_queue_ptr;
44   unsigned long kernarg_template_ptr;
45 } atmi_implicit_args_t;
46 
47 #ifdef __cplusplus
48 extern "C" {
49 #endif
50 
51 #define check(msg, status)                                                     \
52   if (status != HSA_STATUS_SUCCESS) {                                          \
53     printf("%s failed.\n", #msg);                                              \
54     exit(1);                                                                   \
55   }
56 
57 #ifdef DEBUG
58 #define DEBUG_PRINT(fmt, ...)                                                  \
59   if (core::Runtime::getInstance().getDebugMode()) {                           \
60     fprintf(stderr, "[%s:%d] " fmt, __FILE__, __LINE__, ##__VA_ARGS__);        \
61   }
62 #else
63 #define DEBUG_PRINT(...)                                                       \
64   do {                                                                         \
65   } while (false)
66 #endif
67 
68 #ifndef HSA_RUNTIME_INC_HSA_H_
69 typedef struct hsa_signal_s {
70   uint64_t handle;
71 } hsa_signal_t;
72 #endif
73 
74 /*  All global values go in this global structure */
75 typedef struct atl_context_s {
76   bool struct_initialized;
77   bool g_hsa_initialized;
78   bool g_gpu_initialized;
79   bool g_tasks_initialized;
80 } atl_context_t;
81 extern atl_context_t atlc;
82 extern atl_context_t *atlc_p;
83 
84 #ifdef __cplusplus
85 }
86 #endif
87 
88 /* ---------------------------------------------------------------------------------
89  * Simulated CPU Data Structures and API
90  * ---------------------------------------------------------------------------------
91  */
92 
93 #define ATMI_WAIT_STATE HSA_WAIT_STATE_BLOCKED
94 
95 // ---------------------- Kernel Start -------------
96 typedef struct atl_kernel_info_s {
97   uint64_t kernel_object;
98   uint32_t group_segment_size;
99   uint32_t private_segment_size;
100   uint32_t sgpr_count;
101   uint32_t vgpr_count;
102   uint32_t sgpr_spill_count;
103   uint32_t vgpr_spill_count;
104   uint32_t kernel_segment_size;
105   uint32_t num_args;
106   std::vector<uint64_t> arg_alignments;
107   std::vector<uint64_t> arg_offsets;
108   std::vector<uint64_t> arg_sizes;
109 } atl_kernel_info_t;
110 
111 typedef struct atl_symbol_info_s {
112   uint64_t addr;
113   uint32_t size;
114 } atl_symbol_info_t;
115 
116 extern std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
117 extern std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
118 
119 // ---------------------- Kernel End -------------
120 
121 extern struct timespec context_init_time;
122 
123 namespace core {
124 class TaskgroupImpl;
125 class TaskImpl;
126 class Kernel;
127 class KernelImpl;
128 } // namespace core
129 
130 struct SignalPoolT {
131   SignalPoolT() {
132     // If no signals are created, and none can be created later,
133     // will ultimately fail at pop()
134 
135     unsigned N = 1024; // default max pool size from atmi
136     for (unsigned i = 0; i < N; i++) {
137       hsa_signal_t new_signal;
138       hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
139       if (err != HSA_STATUS_SUCCESS) {
140         break;
141       }
142       state.push(new_signal);
143     }
144     DEBUG_PRINT("Signal Pool Initial Size: %lu\n", state.size());
145   }
146   SignalPoolT(const SignalPoolT &) = delete;
147   SignalPoolT(SignalPoolT &&) = delete;
148   ~SignalPoolT() {
149     size_t N = state.size();
150     for (size_t i = 0; i < N; i++) {
151       hsa_signal_t signal = state.front();
152       state.pop();
153       hsa_status_t rc = hsa_signal_destroy(signal);
154       if (rc != HSA_STATUS_SUCCESS) {
155         DEBUG_PRINT("Signal pool destruction failed\n");
156       }
157     }
158   }
159   size_t size() {
160     lock l(&mutex);
161     return state.size();
162   }
163   void push(hsa_signal_t s) {
164     lock l(&mutex);
165     state.push(s);
166   }
167   hsa_signal_t pop(void) {
168     lock l(&mutex);
169     if (!state.empty()) {
170       hsa_signal_t res = state.front();
171       state.pop();
172       return res;
173     }
174 
175     // Pool empty, attempt to create another signal
176     hsa_signal_t new_signal;
177     hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
178     if (err == HSA_STATUS_SUCCESS) {
179       return new_signal;
180     }
181 
182     // Fail
183     return {0};
184   }
185 
186 private:
187   static pthread_mutex_t mutex;
188   std::queue<hsa_signal_t> state;
189   struct lock {
190     lock(pthread_mutex_t *m) : m(m) { pthread_mutex_lock(m); }
191     ~lock() { pthread_mutex_unlock(m); }
192     pthread_mutex_t *m;
193   };
194 };
195 
196 extern std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
197 
198 namespace core {
199 atmi_status_t atl_init_gpu_context();
200 
201 hsa_status_t init_hsa();
202 hsa_status_t finalize_hsa();
203 /*
204  * Generic utils
205  */
206 template <typename T> inline T alignDown(T value, size_t alignment) {
207   return (T)(value & ~(alignment - 1));
208 }
209 
210 template <typename T> inline T *alignDown(T *value, size_t alignment) {
211   return reinterpret_cast<T *>(alignDown((intptr_t)value, alignment));
212 }
213 
214 template <typename T> inline T alignUp(T value, size_t alignment) {
215   return alignDown((T)(value + alignment - 1), alignment);
216 }
217 
218 template <typename T> inline T *alignUp(T *value, size_t alignment) {
219   return reinterpret_cast<T *>(
220       alignDown((intptr_t)(value + alignment - 1), alignment));
221 }
222 
223 extern void register_allocation(void *addr, size_t size,
224                                 atmi_mem_place_t place);
225 extern hsa_amd_memory_pool_t
226 get_memory_pool_by_mem_place(atmi_mem_place_t place);
227 extern bool atl_is_atmi_initialized();
228 
229 bool handle_group_signal(hsa_signal_value_t value, void *arg);
230 
231 void packet_store_release(uint32_t *packet, uint16_t header, uint16_t rest);
232 uint16_t
233 create_header(hsa_packet_type_t type, int barrier,
234               atmi_task_fence_scope_t acq_fence = ATMI_FENCE_SCOPE_SYSTEM,
235               atmi_task_fence_scope_t rel_fence = ATMI_FENCE_SCOPE_SYSTEM);
236 
237 void allow_access_to_all_gpu_agents(void *ptr);
238 } // namespace core
239 
240 const char *get_error_string(hsa_status_t err);
241 const char *get_atmi_error_string(atmi_status_t err);
242 
243 #define ATMIErrorCheck(msg, status)                                            \
244   if (status != ATMI_STATUS_SUCCESS) {                                         \
245     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, #msg,                \
246            get_atmi_error_string(status));                                     \
247     exit(1);                                                                   \
248   } else {                                                                     \
249     /*  printf("%s succeeded.\n", #msg);*/                                     \
250   }
251 
252 #define ErrorCheck(msg, status)                                                \
253   if (status != HSA_STATUS_SUCCESS) {                                          \
254     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, #msg,                \
255            get_error_string(status));                                          \
256     exit(1);                                                                   \
257   } else {                                                                     \
258     /*  printf("%s succeeded.\n", #msg);*/                                     \
259   }
260 
261 #define ErrorCheckAndContinue(msg, status)                                     \
262   if (status != HSA_STATUS_SUCCESS) {                                          \
263     DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, #msg,           \
264                 get_error_string(status));                                     \
265     continue;                                                                  \
266   } else {                                                                     \
267     /*  printf("%s succeeded.\n", #msg);*/                                     \
268   }
269 
270 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_
271