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 #ifdef DEBUG
52 #define DEBUG_PRINT(fmt, ...)                                                  \
53   if (core::Runtime::getInstance().getDebugMode()) {                           \
54     fprintf(stderr, "[%s:%d] " fmt, __FILE__, __LINE__, ##__VA_ARGS__);        \
55   }
56 #else
57 #define DEBUG_PRINT(...)                                                       \
58   do {                                                                         \
59   } while (false)
60 #endif
61 
62 #ifndef HSA_RUNTIME_INC_HSA_H_
63 typedef struct hsa_signal_s {
64   uint64_t handle;
65 } hsa_signal_t;
66 #endif
67 
68 /*  All global values go in this global structure */
69 typedef struct atl_context_s {
70   bool struct_initialized;
71   bool g_hsa_initialized;
72   bool g_gpu_initialized;
73   bool g_tasks_initialized;
74 } atl_context_t;
75 extern atl_context_t atlc;
76 
77 #ifdef __cplusplus
78 }
79 #endif
80 
81 /* ---------------------------------------------------------------------------------
82  * Simulated CPU Data Structures and API
83  * ---------------------------------------------------------------------------------
84  */
85 
86 #define ATMI_WAIT_STATE HSA_WAIT_STATE_BLOCKED
87 
88 // ---------------------- Kernel Start -------------
89 typedef struct atl_kernel_info_s {
90   uint64_t kernel_object;
91   uint32_t group_segment_size;
92   uint32_t private_segment_size;
93   uint32_t sgpr_count;
94   uint32_t vgpr_count;
95   uint32_t sgpr_spill_count;
96   uint32_t vgpr_spill_count;
97   uint32_t kernel_segment_size;
98   uint32_t num_args;
99   std::vector<uint64_t> arg_alignments;
100   std::vector<uint64_t> arg_offsets;
101   std::vector<uint64_t> arg_sizes;
102 } atl_kernel_info_t;
103 
104 typedef struct atl_symbol_info_s {
105   uint64_t addr;
106   uint32_t size;
107 } atl_symbol_info_t;
108 
109 // ---------------------- Kernel End -------------
110 
111 namespace core {
112 class TaskgroupImpl;
113 class TaskImpl;
114 class Kernel;
115 class KernelImpl;
116 } // namespace core
117 
118 struct SignalPoolT {
119   SignalPoolT() {
120     // If no signals are created, and none can be created later,
121     // will ultimately fail at pop()
122 
123     unsigned N = 1024; // default max pool size from atmi
124     for (unsigned i = 0; i < N; i++) {
125       hsa_signal_t new_signal;
126       hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
127       if (err != HSA_STATUS_SUCCESS) {
128         break;
129       }
130       state.push(new_signal);
131     }
132     DEBUG_PRINT("Signal Pool Initial Size: %lu\n", state.size());
133   }
134   SignalPoolT(const SignalPoolT &) = delete;
135   SignalPoolT(SignalPoolT &&) = delete;
136   ~SignalPoolT() {
137     size_t N = state.size();
138     for (size_t i = 0; i < N; i++) {
139       hsa_signal_t signal = state.front();
140       state.pop();
141       hsa_status_t rc = hsa_signal_destroy(signal);
142       if (rc != HSA_STATUS_SUCCESS) {
143         DEBUG_PRINT("Signal pool destruction failed\n");
144       }
145     }
146   }
147   size_t size() {
148     lock l(&mutex);
149     return state.size();
150   }
151   void push(hsa_signal_t s) {
152     lock l(&mutex);
153     state.push(s);
154   }
155   hsa_signal_t pop(void) {
156     lock l(&mutex);
157     if (!state.empty()) {
158       hsa_signal_t res = state.front();
159       state.pop();
160       return res;
161     }
162 
163     // Pool empty, attempt to create another signal
164     hsa_signal_t new_signal;
165     hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
166     if (err == HSA_STATUS_SUCCESS) {
167       return new_signal;
168     }
169 
170     // Fail
171     return {0};
172   }
173 
174 private:
175   static pthread_mutex_t mutex;
176   std::queue<hsa_signal_t> state;
177   struct lock {
178     lock(pthread_mutex_t *m) : m(m) { pthread_mutex_lock(m); }
179     ~lock() { pthread_mutex_unlock(m); }
180     pthread_mutex_t *m;
181   };
182 };
183 
184 extern std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
185 
186 namespace core {
187 hsa_status_t atl_init_gpu_context();
188 
189 hsa_status_t init_hsa();
190 hsa_status_t finalize_hsa();
191 /*
192  * Generic utils
193  */
194 template <typename T> inline T alignDown(T value, size_t alignment) {
195   return (T)(value & ~(alignment - 1));
196 }
197 
198 template <typename T> inline T *alignDown(T *value, size_t alignment) {
199   return reinterpret_cast<T *>(alignDown((intptr_t)value, alignment));
200 }
201 
202 template <typename T> inline T alignUp(T value, size_t alignment) {
203   return alignDown((T)(value + alignment - 1), alignment);
204 }
205 
206 template <typename T> inline T *alignUp(T *value, size_t alignment) {
207   return reinterpret_cast<T *>(
208       alignDown((intptr_t)(value + alignment - 1), alignment));
209 }
210 
211 hsa_status_t register_allocation(void *addr, size_t size,
212                                  atmi_devtype_t DeviceType);
213 
214 extern bool atl_is_atmi_initialized();
215 
216 bool handle_group_signal(hsa_signal_value_t value, void *arg);
217 
218 hsa_status_t allow_access_to_all_gpu_agents(void *ptr);
219 } // namespace core
220 
221 const char *get_error_string(hsa_status_t err);
222 const char *get_atmi_error_string(hsa_status_t err);
223 
224 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_
225