1 //===--- amdgpu/impl/internal.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 #ifndef SRC_RUNTIME_INCLUDE_INTERNAL_H_
9 #define SRC_RUNTIME_INCLUDE_INTERNAL_H_
10 #include <inttypes.h>
11 #include <pthread.h>
12 #include <stddef.h>
13 #include <stdint.h>
14 #include <stdio.h>
15 #include <stdlib.h>
16 
17 #include <cstring>
18 #include <map>
19 #include <queue>
20 #include <string>
21 #include <utility>
22 #include <vector>
23 
24 #include "hsa_api.h"
25 
26 #include "impl_runtime.h"
27 
28 #ifndef TARGET_NAME
29 #error "Missing TARGET_NAME macro"
30 #endif
31 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
32 #include "Debug.h"
33 
34 #define MAX_NUM_KERNELS (1024 * 16)
35 
36 typedef struct impl_implicit_args_s {
37   unsigned long offset_x;
38   unsigned long offset_y;
39   unsigned long offset_z;
40   unsigned long hostcall_ptr;
41   char num_gpu_queues;
42   unsigned long gpu_queue_ptr;
43   char num_cpu_queues;
44   unsigned long cpu_worker_signals;
45   unsigned long cpu_queue_ptr;
46   unsigned long kernarg_template_ptr;
47 } impl_implicit_args_t;
48 
49 // ---------------------- Kernel Start -------------
50 typedef struct atl_kernel_info_s {
51   uint64_t kernel_object;
52   uint32_t group_segment_size;
53   uint32_t private_segment_size;
54   uint32_t sgpr_count;
55   uint32_t vgpr_count;
56   uint32_t sgpr_spill_count;
57   uint32_t vgpr_spill_count;
58   uint32_t kernel_segment_size;
59   uint32_t num_args;
60   std::vector<uint64_t> arg_alignments;
61   std::vector<uint64_t> arg_offsets;
62   std::vector<uint64_t> arg_sizes;
63 } atl_kernel_info_t;
64 
65 typedef struct atl_symbol_info_s {
66   uint64_t addr;
67   uint32_t size;
68 } atl_symbol_info_t;
69 
70 // ---------------------- Kernel End -------------
71 
72 namespace core {
73 class TaskgroupImpl;
74 class TaskImpl;
75 class Kernel;
76 class KernelImpl;
77 } // namespace core
78 
79 struct SignalPoolT {
80   SignalPoolT() {}
81   SignalPoolT(const SignalPoolT &) = delete;
82   SignalPoolT(SignalPoolT &&) = delete;
83   ~SignalPoolT() {
84     size_t N = state.size();
85     for (size_t i = 0; i < N; i++) {
86       hsa_signal_t signal = state.front();
87       state.pop();
88       hsa_status_t rc = hsa_signal_destroy(signal);
89       if (rc != HSA_STATUS_SUCCESS) {
90         DP("Signal pool destruction failed\n");
91       }
92     }
93   }
94   size_t size() {
95     lock l(&mutex);
96     return state.size();
97   }
98   void push(hsa_signal_t s) {
99     lock l(&mutex);
100     state.push(s);
101   }
102   hsa_signal_t pop(void) {
103     lock l(&mutex);
104     if (!state.empty()) {
105       hsa_signal_t res = state.front();
106       state.pop();
107       return res;
108     }
109 
110     // Pool empty, attempt to create another signal
111     hsa_signal_t new_signal;
112     hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
113     if (err == HSA_STATUS_SUCCESS) {
114       return new_signal;
115     }
116 
117     // Fail
118     return {0};
119   }
120 
121 private:
122   static pthread_mutex_t mutex;
123   std::queue<hsa_signal_t> state;
124   struct lock {
125     lock(pthread_mutex_t *m) : m(m) { pthread_mutex_lock(m); }
126     ~lock() { pthread_mutex_unlock(m); }
127     pthread_mutex_t *m;
128   };
129 };
130 
131 namespace core {
132 hsa_status_t atl_init_gpu_context();
133 
134 hsa_status_t init_hsa();
135 hsa_status_t finalize_hsa();
136 /*
137  * Generic utils
138  */
139 template <typename T> inline T alignDown(T value, size_t alignment) {
140   return (T)(value & ~(alignment - 1));
141 }
142 
143 template <typename T> inline T *alignDown(T *value, size_t alignment) {
144   return reinterpret_cast<T *>(alignDown((intptr_t)value, alignment));
145 }
146 
147 template <typename T> inline T alignUp(T value, size_t alignment) {
148   return alignDown((T)(value + alignment - 1), alignment);
149 }
150 
151 template <typename T> inline T *alignUp(T *value, size_t alignment) {
152   return reinterpret_cast<T *>(
153       alignDown((intptr_t)(value + alignment - 1), alignment));
154 }
155 
156 extern bool atl_is_impl_initialized();
157 
158 bool handle_group_signal(hsa_signal_value_t value, void *arg);
159 
160 hsa_status_t allow_access_to_all_gpu_agents(void *ptr);
161 } // namespace core
162 
163 inline const char *get_error_string(hsa_status_t err) {
164   const char *res;
165   hsa_status_t rc = hsa_status_string(err, &res);
166   return (rc == HSA_STATUS_SUCCESS) ? res : "HSA_STATUS UNKNOWN.";
167 }
168 
169 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_
170