Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
arch_device_host.h
Go to the documentation of this file.
1#ifndef ARCH_DEVICE_HOST_H
2#define ARCH_DEVICE_HOST_H
3
4/* Define architecture-specific macros */
5#define ARCH_LOOP_LAMBDA [=]
6#define ARCH_INNER_BODY2(i, j, aggregate) return [=](auto i, auto j, auto *aggregate)
7#define ARCH_INNER_BODY3(i, j, k, aggregate) return [=](auto i, auto j, auto k, auto *aggregate)
8#define ARCH_INNER_BODY4(i, j, k, l, aggregate) return [=](auto i, auto j, auto k, auto l, auto *aggregate)
9
10/* Namespace for architecture-specific functions */
11namespace arch{
12
13/* Buffer class for host compulation units */
14 template <typename T>
15 class buf {
16 private:
17 T *ptr;
18 T *d_ptr;
19 uint bytes;
20 uint is_copy = 0;
21
22 public:
23
24 void syncDeviceData(void) {}
25
26 void syncHostData(void) {}
27
28 buf(T * const _ptr, uint _bytes) : ptr(_ptr), bytes(_bytes) {}
29
30 buf(const buf& u) :
31 ptr(u.ptr), d_ptr(u.d_ptr), bytes(u.bytes), is_copy(1) {}
32
33 T &operator [] (uint i) const {
34 return ptr[i];
35 }
36 };
37
38/* Host function for memory allocation */
39 inline static void* allocate(size_t bytes) {
40 return malloc(bytes);
41 }
42
43/* Host function for memory deallocation */
44 inline static void free(void* ptr) {
45 ::free(ptr);
46 }
47
48/* Host-to-device memory copy */
49 template <typename T>
50 inline static void memcpy_h2d(T* dst, T* src, size_t bytes) {}
51
52/* Device-to-host memory copy */
53 template <typename T>
54 inline static void memcpy_d2h(T* dst, T* src, size_t bytes) {}
55
56/* Register, ie, page-lock existing host allocations */
57 template <typename T>
58 inline static void host_register(T* ptr, size_t bytes) {}
59
60/* Unregister page-locked host allocations */
61 template <typename T>
62 inline static void host_unregister(T* ptr) {}
63
64/* Parallel reduce driver function - specialization for 1D case */
65 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T>
66 inline static void parallel_reduce_driver(const uint (&limits)[1], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
67
68 if(Op == reduce_op::sum) {
69 #pragma omp parallel for reduction(+:sum[:n_redu_dynamic])
70 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
71 loop_body(idx0, sum);
72 }
73 } else if (Op == reduce_op::max) {
74 #pragma omp parallel for reduction(max:sum[:n_redu_dynamic])
75 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
76 loop_body(idx0, sum);
77 }
78 } else if (Op == reduce_op::min) {
79 #pragma omp parallel for reduction(min:sum[:n_redu_dynamic])
80 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
81 loop_body(idx0, sum);
82 }
83 } else if (Op == reduce_op::null) {
84 #pragma omp parallel for
85 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
86 loop_body(idx0, sum);
87 }
88 } else {
89 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
90 }
91 }
92
93/* Parallel reduce driver function - specialization for 2D case */
94 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type>
95 inline static void parallel_reduce_driver(const uint (&limits)[2], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
96
97 if(Op == reduce_op::sum) {
98 #pragma omp parallel for collapse(2) reduction(+:sum[:n_redu_dynamic])
99 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
100 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
101 loop_body(idx0, idx1, sum);
102 }
103 }
104 } else if (Op == reduce_op::max) {
105 #pragma omp parallel for collapse(2) reduction(max:sum[:n_redu_dynamic])
106 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
107 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
108 loop_body(idx0, idx1, sum);
109 }
110 }
111 } else if (Op == reduce_op::min) {
112 #pragma omp parallel for collapse(2) reduction(min:sum[:n_redu_dynamic])
113 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
114 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
115 loop_body(idx0, idx1, sum);
116 }
117 }
118 } else if (Op == reduce_op::null) {
119 #pragma omp parallel for collapse(2)
120 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
121 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
122 loop_body(idx0, idx1, sum);
123 }
124 }
125 } else {
126 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
127 }
128 }
129
130/* Parallel reduce driver function - specialization for 2D case with nested bodies */
131 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type, typename = void>
132 inline static void parallel_reduce_driver(const uint (&limits)[2], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
133
134 if(Op == reduce_op::sum) {
135 #pragma omp parallel for reduction(+:sum[:n_redu_dynamic])
136 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
137 auto inner_loop = loop_body(idx1, idx1, sum);
138 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
139 inner_loop(idx0, idx1, sum);
140 }
141 }
142 } else if (Op == reduce_op::max) {
143 #pragma omp parallel for reduction(max:sum[:n_redu_dynamic])
144 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
145 auto inner_loop = loop_body(idx1, idx1, sum);
146 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
147 inner_loop(idx0, idx1, sum);
148 }
149 }
150 } else if (Op == reduce_op::min) {
151 #pragma omp parallel for reduction(min:sum[:n_redu_dynamic])
152 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
153 auto inner_loop = loop_body(idx1, idx1, sum);
154 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
155 inner_loop(idx0, idx1, sum);
156 }
157 }
158 } else if (Op == reduce_op::null) {
159 #pragma omp parallel for
160 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
161 auto inner_loop = loop_body(idx1, idx1, sum);
162 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
163 inner_loop(idx0, idx1, sum);
164 }
165 }
166 } else {
167 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
168 }
169 }
170
171/* Parallel reduce driver function - specialization for 3D case */
172 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type>
173 inline static void parallel_reduce_driver(const uint (&limits)[3], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
174
175 if(Op == reduce_op::sum) {
176 #pragma omp parallel for collapse(3) reduction(+:sum[:n_redu_dynamic])
177 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
178 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
179 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
180 loop_body(idx0, idx1, idx2, sum);
181 }
182 }
183 }
184 } else if (Op == reduce_op::max) {
185 #pragma omp parallel for collapse(3) reduction(max:sum[:n_redu_dynamic])
186 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
187 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
188 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
189 loop_body(idx0, idx1, idx2, sum);
190 }
191 }
192 }
193 } else if (Op == reduce_op::min) {
194 #pragma omp parallel for collapse(3) reduction(min:sum[:n_redu_dynamic])
195 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
196 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
197 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
198 loop_body(idx0, idx1, idx2, sum);
199 }
200 }
201 }
202 } else if (Op == reduce_op::null) {
203 #pragma omp parallel for collapse(3)
204 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
205 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
206 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
207 loop_body(idx0, idx1, idx2, sum);
208 }
209 }
210 }
211 } else {
212 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
213 }
214 }
215
216/* Parallel reduce driver function - specialization for 3D case with nested bodies */
217 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type, typename = void>
218 inline static void parallel_reduce_driver(const uint (&limits)[3], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
219
220 if(Op == reduce_op::sum) {
221 #pragma omp parallel for reduction(+:sum[:n_redu_dynamic])
222 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
223 auto inner_loop = loop_body(idx2, idx2, idx2, sum);
224 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
225 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
226 inner_loop(idx0, idx1, idx2, sum);
227 }
228 }
229 }
230 } else if (Op == reduce_op::max) {
231 #pragma omp parallel for reduction(max:sum[:n_redu_dynamic])
232 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
233 auto inner_loop = loop_body(idx2, idx2, idx2, sum);
234 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
235 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
236 inner_loop(idx0, idx1, idx2, sum);
237 }
238 }
239 }
240 } else if (Op == reduce_op::min) {
241 #pragma omp parallel for reduction(min:sum[:n_redu_dynamic])
242 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
243 auto inner_loop = loop_body(idx2, idx2, idx2, sum);
244 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
245 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
246 inner_loop(idx0, idx1, idx2, sum);
247 }
248 }
249 }
250 } else if (Op == reduce_op::null) {
251 #pragma omp parallel for
252 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
253 auto inner_loop = loop_body(idx2, idx2, idx2, sum);
254 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
255 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
256 inner_loop(idx0, idx1, idx2, sum);
257 }
258 }
259 }
260 } else {
261 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
262 }
263 }
264
265/* Parallel reduce driver function - specialization for 4D case */
266 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type>
267 inline static void parallel_reduce_driver(const uint (&limits)[4], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
268
269 if(Op == reduce_op::sum) {
270 #pragma omp parallel for collapse(4) reduction(+:sum[:n_redu_dynamic])
271 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
272 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
273 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
274 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
275 loop_body(idx0, idx1, idx2, idx3, sum);
276 }
277 }
278 }
279 }
280 } else if (Op == reduce_op::max) {
281 #pragma omp parallel for collapse(4) reduction(max:sum[:n_redu_dynamic])
282 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
283 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
284 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
285 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
286 loop_body(idx0, idx1, idx2, idx3, sum);
287 }
288 }
289 }
290 }
291 } else if (Op == reduce_op::min) {
292 #pragma omp parallel for collapse(4) reduction(min:sum[:n_redu_dynamic])
293 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
294 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
295 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
296 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
297 loop_body(idx0, idx1, idx2, idx3, sum);
298 }
299 }
300 }
301 }
302 } else if (Op == reduce_op::null) {
303 #pragma omp parallel for collapse(4)
304 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
305 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
306 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
307 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
308 loop_body(idx0, idx1, idx2, idx3, sum);
309 }
310 }
311 }
312 }
313 } else {
314 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
315 }
316 }
317
318/* Parallel reduce driver function - specialization for 4D case with nested bodies */
319 template <reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type, typename = void>
320 inline static void parallel_reduce_driver(const uint (&limits)[4], Lambda loop_body, T *sum, const uint n_redu_dynamic) {
321
322 if(Op == reduce_op::sum) {
323 #pragma omp parallel for reduction(+:sum[:n_redu_dynamic])
324 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
325 auto inner_loop = loop_body(idx3, idx3, idx3, idx3, sum);
326 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
327 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
328 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
329 inner_loop(idx0, idx1, idx2, idx3, sum);
330 }
331 }
332 }
333 }
334 } else if (Op == reduce_op::max) {
335 #pragma omp parallel for reduction(max:sum[:n_redu_dynamic])
336 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
337 auto inner_loop = loop_body(idx3, idx3, idx3, idx3, sum);
338 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
339 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
340 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
341 inner_loop(idx0, idx1, idx2, idx3, sum);
342 }
343 }
344 }
345 }
346 } else if (Op == reduce_op::min) {
347 #pragma omp parallel for reduction(min:sum[:n_redu_dynamic])
348 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
349 auto inner_loop = loop_body(idx3, idx3, idx3, idx3, sum);
350 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
351 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
352 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
353 inner_loop(idx0, idx1, idx2, idx3, sum);
354 }
355 }
356 }
357 }
358 } else if (Op == reduce_op::null) {
359 #pragma omp parallel for
360 for (uint idx3 = 0; idx3 < limits[3]; ++idx3) {
361 auto inner_loop = loop_body(idx3, idx3, idx3, idx3, sum);
362 for (uint idx2 = 0; idx2 < limits[2]; ++idx2) {
363 for (uint idx1 = 0; idx1 < limits[1]; ++idx1) {
364 for (uint idx0 = 0; idx0 < limits[0]; ++idx0) {
365 inner_loop(idx0, idx1, idx2, idx3, sum);
366 }
367 }
368 }
369 }
370 } else {
371 printf("ERROR at %s:%d: Invalid reduction identifier \"Op\".", __FILE__, __LINE__);
372 }
373 }
374}
375#endif // !ARCH_DEVICE_HOST_H
i
for i
Definition Dispersion.m:24
arch::buf::syncDeviceData
void syncDeviceData(void)
Definition arch_device_host.h:24
arch::buf::operator[]
__host__ __device__ T & operator[](uint i) const
Definition arch_device_cuda.h:220
arch::buf::syncHostData
void syncHostData(void)
Definition arch_device_host.h:26
arch::buf::buf
buf(const buf &u)
Definition arch_device_host.h:30
arch::buf::buf
buf(T *const _ptr, uint _bytes)
Definition arch_device_host.h:28
arch::n_redu_dynamic
static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint n_redu_dynamic
Definition arch_device_cuda.h:359
arch::memcpy_d2h
static __forceinline__ void memcpy_d2h(T *dst, T *src, size_t bytes)
Definition arch_device_cuda.h:295
arch::parallel_reduce_driver
static __forceinline__ void parallel_reduce_driver(const uint(&limits)[NDim], Lambda loop_body, T *sum, const uint n_redu_dynamic)
Definition arch_device_cuda.h:442
arch::host_unregister
static __forceinline__ void host_unregister(T *ptr)
Definition arch_device_cuda.h:317
arch::allocate
__host__ static __forceinline__ void * allocate(size_t bytes)
Definition arch_device_cuda.h:243
arch::host_register
static __forceinline__ void host_register(T *ptr, size_t bytes)
Definition arch_device_cuda.h:311
arch::memcpy_h2d
static __forceinline__ void memcpy_h2d(T *dst, T *src, size_t bytes)
Definition arch_device_cuda.h:279
arch::uint
uint32_t uint
Definition arch_device_api.h:17
arch::free
__host__ static __forceinline__ void free(T *ptr)
Definition arch_device_cuda.h:264