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