Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_GpuLaunchFunctsC.H
Go to the documentation of this file.
1#ifndef AMREX_GPU_LAUNCH_FUNCTS_C_H_
2#define AMREX_GPU_LAUNCH_FUNCTS_C_H_
3#include <AMReX_Config.H>
4
5namespace amrex {
6
15template<int WIDTH, class N=int>
17{
19 static constexpr int width = WIDTH;
20
22 N index = 0;
23};
24
25namespace detail {
26
27 // call_f_scalar_handler
28
29 template <typename F, typename N>
31 auto call_f_scalar_handler (F const& f, N i)
32 noexcept -> decltype(f(0))
33 {
34 f(i);
35 }
36
37 template <typename F, typename N>
39 auto call_f_scalar_handler (F const& f, N i)
40 noexcept -> decltype(f(0,Gpu::Handler{}))
41 {
42 f(i, Gpu::Handler{});
43 }
44
45 // call_f_intvect_inner
46
47 template <typename F, std::size_t...Ns, class...Args>
49 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<1> iv, Args...args)
50 noexcept -> decltype(f(0, 0, 0, args...))
51 {
52 f(iv[0], 0, 0, args...);
53 }
54
55 template <typename F, std::size_t...Ns, class...Args>
57 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<2> iv, Args...args)
58 noexcept -> decltype(f(0, 0, 0, args...))
59 {
60 f(iv[0], iv[1], 0, args...);
61 }
62
63 template <typename F, int dim, std::size_t...Ns, class...Args>
65 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<dim> iv, Args...args)
66 noexcept -> decltype(f(iv, args...))
67 {
68 f(iv, args...);
69 }
70
71 template <typename F, int dim, std::size_t...Ns, class...Args>
73 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<dim> iv, Args...args)
74 noexcept -> decltype(f(iv[Ns]..., args...))
75 {
76 f(iv[Ns]..., args...);
77 }
78
79 // call_f_intvect_engine
80
81 template <typename F, int dim>
84 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, engine))
85 {
86 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, engine);
87 }
88
89 // call_f_intvect_handler
90
91 template <typename F, int dim>
94 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv))
95 {
96 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv);
97 }
98
99 template <typename F, int dim>
101 auto call_f_intvect_handler (F const& f, IntVectND<dim> iv)
102 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, Gpu::Handler{}))
103 {
104 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, Gpu::Handler{});
105 }
106
107 // call_f_intvect_ncomp_engine
108
109 template <typename F, typename T, int dim>
112 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, engine))
113 {
114 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, engine);
115 }
116
117 // call_f_intvect_ncomp_handler
118
119 template <typename F, typename T, int dim>
122 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n))
123 {
124 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n);
125 }
126
127 template <typename F, typename T, int dim>
129 auto call_f_intvect_ncomp_handler (F const& f, IntVectND<dim> iv, T n)
130 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, Gpu::Handler{}))
131 {
132 call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, Gpu::Handler{});
133 }
134
135}
136
137template<typename T, typename L>
138void launch (T const& n, L&& f) noexcept
139{
140 std::forward<L>(f)(n);
141}
142
143template<int MT, typename T, typename L>
144void launch (T const& n, L&& f) noexcept
145{
147 std::forward<L>(f)(n);
148}
149
150template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
152void For (T n, L const& f) noexcept
153{
154 for (T i = 0; i < n; ++i) {
156 }
157}
158
159template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
160void For (T n, L&& f) noexcept
161{
163 For(n, std::forward<L>(f));
164}
165
166template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
167void For (Gpu::KernelInfo const&, T n, L&& f) noexcept
168{
169 For(n, std::forward<L>(f));
170}
171
172template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
173void For (Gpu::KernelInfo const&, T n, L&& f) noexcept
174{
176 For(n, std::forward<L>(f));
177}
178
179template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
181void ParallelFor (T n, L const& f) noexcept
182{
184 for (T i = 0; i < n; ++i) {
186 }
187}
188
189template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
190void ParallelFor (T n, L&& f) noexcept
191{
193 ParallelFor(n, std::forward<L>(f));
194}
195
204template <int WIDTH, typename N, typename L, typename M=std::enable_if_t<std::is_integral_v<N>> >
206void ParallelForSIMD (N n, L const& f) noexcept
207{
208 N i = 0;
209 // vectorize full lanes
210 for (; i + WIDTH <= n; i+=WIDTH) {
212 }
213 // scalar handling of the remainder
214 // note: we could make the remainder calls faster, by repeatedly
215 // decreasing the SIMD width by 2 until we reach 1
216 for (; i < n; ++i) {
217 f(SIMDindex<1, N>{i});
218 }
219}
220
221template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
222void ParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
223{
224 ParallelFor(n, std::forward<L>(f));
225}
226
227template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
228void ParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
229{
231 ParallelFor(n, std::forward<L>(f));
232}
233
234namespace detail {
235
236template <int idim, typename L, int dim>
238void For_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
239{
240 if constexpr (idim == 1) {
241 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
243 }
244 } else if constexpr (idim == 2) {
245 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
246 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
248 }}
249 } else if constexpr (idim == 3) {
250 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
251 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
252 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
254 }}}
255 } else {
256 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
257 For_impND<idim-1>(f, lo, hi, iv);
258 }
259 }
260}
261
262}
263
264template <typename L, int dim>
266void For (BoxND<dim> const& box, L const& f) noexcept
267{
268 const auto lo = amrex::lbound_iv(box);
269 const auto hi = amrex::ubound_iv(box);
271 detail::For_impND<dim>(f, lo, hi, iv);
272}
273
274template <int MT, typename L, int dim>
275void For (BoxND<dim> const& box, L&& f) noexcept
276{
278 For(box, std::forward<L>(f));
279}
280
281template <typename L, int dim>
282void For (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
283{
284 For(box, std::forward<L>(f));
285}
286
287template <int MT, typename L, int dim>
288void For (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
289{
291 For(box, std::forward<L>(f));
292}
293
294namespace detail {
295
296template <int idim, typename L, int dim>
298void ParallelFor_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
299{
300 if constexpr (idim == 1) {
302 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
304 }
305 } else if constexpr (idim == 2) {
306 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
308 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
310 }}
311 } else if constexpr (idim == 3) {
312 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
313 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
315 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
317 }}}
318 } else {
319 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
320 ParallelFor_impND<idim-1>(f, lo, hi, iv);
321 }
322 }
323}
324
325}
326
327template <typename L, int dim>
329void ParallelFor (BoxND<dim> const& box, L const& f) noexcept
330{
331 const auto lo = amrex::lbound_iv(box);
332 const auto hi = amrex::ubound_iv(box);
334 detail::ParallelFor_impND<dim>(f, lo, hi, iv);
335}
336
337template <int MT, typename L, int dim>
338void ParallelFor (BoxND<dim> const& box, L&& f) noexcept
339{
341 ParallelFor(box, std::forward<L>(f));
342}
343
344template <typename L, int dim>
345void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
346{
347 ParallelFor(box, std::forward<L>(f));
348}
349
350template <int MT, typename L, int dim>
351void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
352{
354 ParallelFor(box, std::forward<L>(f));
355}
356
357namespace detail {
358
359template <int idim, typename L, typename T, int dim>
361void For_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
362{
363 if constexpr (idim == 1) {
364 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
366 }
367 } else if constexpr (idim == 2) {
368 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
369 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
371 }}
372 } else if constexpr (idim == 3) {
373 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
374 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
375 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
377 }}}
378 } else {
379 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
380 For_impND<idim-1>(f, lo, hi, iv, n);
381 }
382 }
383}
384
385}
386
387template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
389void For (BoxND<dim> const& box, T ncomp, L const& f) noexcept
390{
391 const auto lo = amrex::lbound_iv(box);
392 const auto hi = amrex::ubound_iv(box);
394 for (T n = 0; n < ncomp; ++n) {
395 detail::For_impND<dim>(f, lo, hi, iv, n);
396 }
397}
398
399template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
400void For (BoxND<dim> const& box, T ncomp, L&& f) noexcept
401{
403 For(box, ncomp, std::forward<L>(f));
404}
405
406template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
407void For (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
408{
409 For(box, ncomp, std::forward<L>(f));
410}
411
412template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
413void For (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
414{
416 For(box, ncomp, std::forward<L>(f));
417}
418
419namespace detail {
420
421template <int idim, typename L, typename T, int dim>
423void ParallelFor_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
424{
425 if constexpr (idim == 1) {
427 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
429 }
430 } else if constexpr (idim == 2) {
431 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
433 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
435 }}
436 } else if constexpr (idim == 3) {
437 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
438 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
440 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
442 }}}
443 } else {
444 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
445 ParallelFor_impND<idim-1>(f, lo, hi, iv, n);
446 }
447 }
448}
449
450}
451
452template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
454void ParallelFor (BoxND<dim> const& box, T ncomp, L const& f) noexcept
455{
456 const auto lo = amrex::lbound_iv(box);
457 const auto hi = amrex::ubound_iv(box);
459 for (T n = 0; n < ncomp; ++n) {
460 detail::ParallelFor_impND<dim>(f, lo, hi, iv, n);
461 }
462}
463
464template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
465void ParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
466{
468 ParallelFor(box, ncomp, std::forward<L>(f));
469}
470
471template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
472void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
473{
474 ParallelFor(box, ncomp, std::forward<L>(f));
475}
476
477template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
478void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
479{
481 ParallelFor(box, ncomp, std::forward<L>(f));
482}
483
484template <typename L1, typename L2, int dim>
485void For (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
486{
487 For(box1, std::forward<L1>(f1));
488 For(box2, std::forward<L2>(f2));
489}
490
491template <int MT, typename L1, typename L2, int dim>
492void For (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
493{
495 For(box1, std::forward<L1>(f1));
496 For(box2, std::forward<L2>(f2));
497}
498
499template <typename L1, typename L2, int dim>
500void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
501{
502 For (box1, box2, std::forward<L1>(f1), std::forward<L2>(f2));
503}
504
505template <int MT, typename L1, typename L2, int dim>
506void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
507{
509 For (box1, box2, std::forward<L1>(f1), std::forward<L2>(f2));
510}
511
512template <typename L1, typename L2, typename L3, int dim>
513void For (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
514{
515 For(box1, std::forward<L1>(f1));
516 For(box2, std::forward<L2>(f2));
517 For(box3, std::forward<L3>(f3));
518}
519
520template <int MT, typename L1, typename L2, typename L3, int dim>
521void For (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
522{
524 For(box1, std::forward<L1>(f1));
525 For(box2, std::forward<L2>(f2));
526 For(box3, std::forward<L3>(f3));
527}
528
529template <typename L1, typename L2, typename L3, int dim>
530void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
531{
532 For(box1, box2, box3, std::forward<L1>(f1), std::forward<L2>(f2), std::forward<L3>(f3));
533}
534
535template <int MT, typename L1, typename L2, typename L3, int dim>
536void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
537{
539 For(box1, box2, box3, std::forward<L1>(f1), std::forward<L2>(f2), std::forward<L3>(f3));
540}
541
542template <typename T1, typename T2, typename L1, typename L2, int dim,
543 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
544 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
545void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
546 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
547{
548 For(box1, ncomp1, std::forward<L1>(f1));
549 For(box2, ncomp2, std::forward<L2>(f2));
550}
551
552template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
553 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
554 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
555void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
556 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
557{
559 For(box1, ncomp1, std::forward<L1>(f1));
560 For(box2, ncomp2, std::forward<L2>(f2));
561}
562
563template <typename T1, typename T2, typename L1, typename L2, int dim,
564 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
565 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
566void For (Gpu::KernelInfo const&,
567 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
568 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
569{
570 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
571}
572
573template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
574 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
575 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
576void For (Gpu::KernelInfo const&,
577 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
578 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
579{
581 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
582}
583
584template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
585 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
586 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
587 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
588void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
589 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
590 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
591{
592 For(box1, ncomp1, std::forward<L1>(f1));
593 For(box2, ncomp2, std::forward<L2>(f2));
594 For(box3, ncomp3, std::forward<L3>(f3));
595}
596
597template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
598 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
599 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
600 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
601void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
602 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
603 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
604{
606 For(box1, ncomp1, std::forward<L1>(f1));
607 For(box2, ncomp2, std::forward<L2>(f2));
608 For(box3, ncomp3, std::forward<L3>(f3));
609}
610
611template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
612 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
613 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
614 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
615void For (Gpu::KernelInfo const&,
616 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
617 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
618 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
619{
620 For(box1,ncomp1,std::forward<L1>(f1),
621 box2,ncomp2,std::forward<L2>(f2),
622 box3,ncomp3,std::forward<L3>(f3));
623}
624
625template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
626 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
627 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
628 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
629void For (Gpu::KernelInfo const&,
630 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
631 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
632 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
633{
635 For(box1,ncomp1,std::forward<L1>(f1),
636 box2,ncomp2,std::forward<L2>(f2),
637 box3,ncomp3,std::forward<L3>(f3));
638}
639
640template <typename L1, typename L2, int dim>
641void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
642{
643 ParallelFor(box1, std::forward<L1>(f1));
644 ParallelFor(box2, std::forward<L2>(f2));
645}
646
647template <int MT, typename L1, typename L2, int dim>
648void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
649{
651 ParallelFor(box1, std::forward<L1>(f1));
652 ParallelFor(box2, std::forward<L2>(f2));
653}
654
655template <typename L1, typename L2, int dim>
656void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
657{
658 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
659}
660
661template <int MT, typename L1, typename L2, int dim>
662void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
663{
665 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
666}
667
668template <typename L1, typename L2, typename L3, int dim>
669void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
670{
671 ParallelFor(box1, std::forward<L1>(f1));
672 ParallelFor(box2, std::forward<L2>(f2));
673 ParallelFor(box3, std::forward<L3>(f3));
674}
675
676template <int MT, typename L1, typename L2, typename L3, int dim>
677void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
678{
680 ParallelFor(box1, std::forward<L1>(f1));
681 ParallelFor(box2, std::forward<L2>(f2));
682 ParallelFor(box3, std::forward<L3>(f3));
683}
684
685template <typename L1, typename L2, typename L3, int dim>
686void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
687{
688 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
689}
690
691template <int MT, typename L1, typename L2, typename L3, int dim>
692void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
693{
695 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
696}
697
698template <typename T1, typename T2, typename L1, typename L2, int dim,
699 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
700 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
701void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
702 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
703{
704 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
705 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
706}
707
708template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
709 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
710 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
711void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
712 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
713{
715 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
716 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
717}
718
719template <typename T1, typename T2, typename L1, typename L2, int dim,
720 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
721 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
723 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
724 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
725{
726 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
727 box2,ncomp2,std::forward<L2>(f2));
728}
729
730template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
731 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
732 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
734 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
735 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
736{
738 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
739 box2,ncomp2,std::forward<L2>(f2));
740}
741
742template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
743 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
744 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
745 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
746void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
747 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
748 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
749{
750 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
751 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
752 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
753}
754
755template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
756 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
757 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
758 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
759void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
760 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
761 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
762{
764 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
765 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
766 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
767}
768
769template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
770 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
771 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
772 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
774 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
775 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
776 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
777{
778 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
779 box2, ncomp2, std::forward<L2>(f2),
780 box3, ncomp3, std::forward<L3>(f3));
781}
782
783template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
784 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
785 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
786 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
788 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
789 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
790 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
791{
793 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
794 box2, ncomp2, std::forward<L2>(f2),
795 box3, ncomp3, std::forward<L3>(f3));
796}
797
798template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
799void HostDeviceParallelFor (T n, L&& f) noexcept
800{
801 ParallelFor(n,std::forward<L>(f));
802}
803
804template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
805void HostDeviceParallelFor (T n, L&& f) noexcept
806{
808 ParallelFor(n,std::forward<L>(f));
809}
810
811template <typename L, int dim>
812void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
813{
814 ParallelFor(box,std::forward<L>(f));
815}
816
817template <int MT, typename L, int dim>
818void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
819{
821 ParallelFor(box,std::forward<L>(f));
822}
823
824template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
825void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
826{
827 ParallelFor(box,ncomp,std::forward<L>(f));
828}
829
830template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
831void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
832{
834 ParallelFor(box,ncomp,std::forward<L>(f));
835}
836
837template <typename L1, typename L2, int dim>
838void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
839{
840 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
841}
842
843template <int MT, typename L1, typename L2, int dim>
844void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
845{
847 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
848}
849
850template <typename L1, typename L2, typename L3, int dim>
851void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
852 L1&& f1, L2&& f2, L3&& f3) noexcept
853{
854 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
855}
856
857template <int MT, typename L1, typename L2, typename L3, int dim>
858void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
859 L1&& f1, L2&& f2, L3&& f3) noexcept
860{
862 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
863}
864
865template <typename T1, typename T2, typename L1, typename L2, int dim,
866 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
867 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
868void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
869 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
870{
871 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
872}
873
874template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
875 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
876 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
877void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
878 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
879{
881 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
882}
883
884template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
885 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
886 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
887 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
888void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
889 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
890 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
891{
892 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
893 box2,ncomp2,std::forward<L2>(f2),
894 box3,ncomp3,std::forward<L3>(f3));
895}
896
897template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
898 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
899 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
900 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
901void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
902 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
903 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
904{
906 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
907 box2,ncomp2,std::forward<L2>(f2),
908 box3,ncomp3,std::forward<L3>(f3));
909}
910
911template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
912void HostDeviceFor (T n, L&& f) noexcept
913{
914 For(n,std::forward<L>(f));
915}
916
917template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
918void HostDeviceFor (T n, L&& f) noexcept
919{
921 For(n,std::forward<L>(f));
922}
923
924template <typename L, int dim>
925void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
926{
927 For(box,std::forward<L>(f));
928}
929
930template <int MT, typename L, int dim>
931void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
932{
934 For(box,std::forward<L>(f));
935}
936
937template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
938void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
939{
940 For(box,ncomp,std::forward<L>(f));
941}
942
943template <int MT, typename T, int dim, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
944void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
945{
947 For(box,ncomp,std::forward<L>(f));
948}
949
950template <typename L1, typename L2, int dim>
951void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
952{
953 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
954}
955
956template <int MT, typename L1, typename L2, int dim>
957void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
958{
960 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
961}
962
963template <typename L1, typename L2, typename L3, int dim>
964void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
965 L1&& f1, L2&& f2, L3&& f3) noexcept
966{
967 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
968}
969
970template <int MT, typename L1, typename L2, typename L3, int dim>
971void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
972 L1&& f1, L2&& f2, L3&& f3) noexcept
973{
975 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
976}
977
978template <typename T1, typename T2, typename L1, typename L2, int dim,
979 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
980 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
981void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
982 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
983{
984 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
985}
986
987template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
988 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
989 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
990void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
991 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
992{
994 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
995}
996
997template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
998 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
999 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1000 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1001void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1002 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1003 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1004{
1005 For(box1,ncomp1,std::forward<L1>(f1),
1006 box2,ncomp2,std::forward<L2>(f2),
1007 box3,ncomp3,std::forward<L3>(f3));
1008}
1009
1010template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1011 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1012 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1013 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1014void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1015 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1016 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1017{
1019 For(box1,ncomp1,std::forward<L1>(f1),
1020 box2,ncomp2,std::forward<L2>(f2),
1021 box3,ncomp3,std::forward<L3>(f3));
1022}
1023
1024template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1025void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1026{
1027 ParallelFor(n,std::forward<L>(f));
1028}
1029
1030template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1031void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1032{
1034 ParallelFor(n,std::forward<L>(f));
1035}
1036
1037template <typename L, int dim>
1038void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1039{
1040 ParallelFor(box,std::forward<L>(f));
1041}
1042
1043template <int MT, typename L, int dim>
1044void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1045{
1047 ParallelFor(box,std::forward<L>(f));
1048}
1049
1050template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1051void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1052{
1053 ParallelFor(box,ncomp,std::forward<L>(f));
1054}
1055
1056template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1057void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1058{
1060 ParallelFor(box,ncomp,std::forward<L>(f));
1061}
1062
1063template <typename L1, typename L2, int dim>
1064void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1065{
1066 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1067}
1068
1069template <int MT, typename L1, typename L2, int dim>
1070void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1071{
1073 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1074}
1075
1076template <typename L1, typename L2, typename L3, int dim>
1078 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1079 L1&& f1, L2&& f2, L3&& f3) noexcept
1080{
1081 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1082}
1083
1084template <int MT, typename L1, typename L2, typename L3, int dim>
1086 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1087 L1&& f1, L2&& f2, L3&& f3) noexcept
1088{
1090 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1091}
1092
1093template <typename T1, typename T2, typename L1, typename L2, int dim,
1094 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1095 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1097 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1098 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1099{
1100 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1101}
1102
1103template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
1104 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1105 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1107 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1108 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1109{
1111 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1112}
1113
1114template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1115 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1116 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1117 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1119 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1120 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1121 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1122{
1123 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1124 box2,ncomp2,std::forward<L2>(f2),
1125 box3,ncomp3,std::forward<L3>(f3));
1126}
1127
1128template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1129 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1130 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1131 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1133 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1134 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1135 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1136{
1138 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1139 box2,ncomp2,std::forward<L2>(f2),
1140 box3,ncomp3,std::forward<L3>(f3));
1141}
1142
1143template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1144void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1145{
1146 For(n,std::forward<L>(f));
1147}
1148
1149template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1150void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1151{
1153 For(n,std::forward<L>(f));
1154}
1155
1156template <typename L, int dim>
1157void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1158{
1159 For(box,std::forward<L>(f));
1160}
1161
1162template <int MT, typename L, int dim>
1163void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1164{
1166 For(box,std::forward<L>(f));
1167}
1168
1169template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1170void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1171{
1172 For(box,ncomp,std::forward<L>(f));
1173}
1174
1175template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1176void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1177{
1179 For(box,ncomp,std::forward<L>(f));
1180}
1181
1182template <typename L1, typename L2, int dim>
1183void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1184{
1185 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1186}
1187
1188template <int MT, typename L1, typename L2, int dim>
1189void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1190{
1192 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1193}
1194
1195template <typename L1, typename L2, typename L3, int dim>
1197 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1198 L1&& f1, L2&& f2, L3&& f3) noexcept
1199{
1200 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1201}
1202
1203template <int MT, typename L1, typename L2, typename L3, int dim>
1205 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1206 L1&& f1, L2&& f2, L3&& f3) noexcept
1207{
1209 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1210}
1211
1212template <typename T1, typename T2, typename L1, typename L2, int dim,
1213 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1214 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1216 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1217 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1218{
1219 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1220}
1221
1222template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
1223 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1224 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1226 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1227 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1228{
1230 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1231}
1232
1233template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1234 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1235 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1236 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1238 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1239 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1240 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1241{
1242 For(box1,ncomp1,std::forward<L1>(f1),
1243 box2,ncomp2,std::forward<L2>(f2),
1244 box3,ncomp3,std::forward<L3>(f3));
1245}
1246
1247template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1248 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1249 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1250 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1252 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1253 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1254 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1255{
1257 For(box1,ncomp1,std::forward<L1>(f1),
1258 box2,ncomp2,std::forward<L2>(f2),
1259 box3,ncomp3,std::forward<L3>(f3));
1260}
1261
1262template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1264void ParallelForRNG (T n, L const& f) noexcept
1265{
1266 for (T i = 0; i < n; ++i) {
1267 f(i,RandomEngine{});
1268 }
1269}
1270
1271namespace detail {
1272
1273template <int idim, typename L, int dim>
1275void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
1276{
1277 if constexpr (idim == 1) {
1278 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1280 }
1281 } else if constexpr (idim == 2) {
1282 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1283 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1285 }}
1286 } else if constexpr (idim == 3) {
1287 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1288 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1289 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1291 }}}
1292 } else {
1293 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1294 ParallelForRNG_impND<idim-1>(f, lo, hi, iv);
1295 }
1296 }
1297}
1298
1299template <int idim, typename L, typename T, int dim>
1301void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
1302{
1303 if constexpr (idim == 1) {
1304 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1306 }
1307 } else if constexpr (idim == 2) {
1308 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1309 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1311 }}
1312 } else if constexpr (idim == 3) {
1313 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1314 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1315 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1317 }}}
1318 } else {
1319 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1320 ParallelForRNG_impND<idim-1>(f, lo, hi, iv, n);
1321 }
1322 }
1323}
1324
1325}
1326
1327template <typename L, int dim>
1329void ParallelForRNG (BoxND<dim> const& box, L const& f) noexcept
1330{
1331 const auto lo = amrex::lbound_iv(box);
1332 const auto hi = amrex::ubound_iv(box);
1333 IntVectND<dim> iv;
1334 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv);
1335}
1336
1337template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1339void ParallelForRNG (BoxND<dim> const& box, T ncomp, L const& f) noexcept
1340{
1341 const auto lo = amrex::lbound_iv(box);
1342 const auto hi = amrex::ubound_iv(box);
1343 IntVectND<dim> iv;
1344 for (T n = 0; n < ncomp; ++n) {
1345 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv, n);
1346 }
1347}
1348
1349template <typename L>
1350void single_task (L&& f) noexcept
1351{
1352 std::forward<L>(f)();
1353}
1354
1355}
1356
1357#endif
#define AMREX_PRAGMA_SIMD
Definition AMReX_Extension.H:80
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_ATTRIBUTE_FLATTEN_FOR
Definition AMReX_Extension.H:151
A Rectangular Domain on an Integer Lattice.
Definition AMReX_Box.H:46
Definition AMReX_GpuKernelInfo.H:8
Definition AMReX_IntVect.H:55
AMREX_FORCE_INLINE auto call_f_intvect_handler(F const &f, IntVectND< dim > iv) noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence< dim >(), f, iv))
Definition AMReX_GpuLaunchFunctsC.H:93
AMREX_FORCE_INLINE auto call_f_intvect_ncomp_handler(F const &f, IntVectND< dim > iv, T n) noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence< dim >(), f, iv, n))
Definition AMReX_GpuLaunchFunctsC.H:121
AMREX_FORCE_INLINE void For_impND(L const &f, IntVectND< dim > const lo, IntVectND< dim > const hi, IntVectND< dim > iv) noexcept
Definition AMReX_GpuLaunchFunctsC.H:238
AMREX_FORCE_INLINE auto call_f_scalar_handler(F const &f, N i) noexcept -> decltype(f(0))
Definition AMReX_GpuLaunchFunctsC.H:31
AMREX_FORCE_INLINE void ParallelForRNG_impND(L const &f, IntVectND< dim > const lo, IntVectND< dim > const hi, IntVectND< dim > iv) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1275
AMREX_FORCE_INLINE void ParallelFor_impND(L const &f, IntVectND< dim > const lo, IntVectND< dim > const hi, IntVectND< dim > iv) noexcept
Definition AMReX_GpuLaunchFunctsC.H:298
AMREX_FORCE_INLINE auto call_f_intvect_engine(F const &f, IntVectND< dim > iv, RandomEngine engine) noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence< dim >(), f, iv, engine))
Definition AMReX_GpuLaunchFunctsC.H:83
AMREX_FORCE_INLINE auto call_f_intvect_ncomp_engine(F const &f, IntVectND< dim > iv, T n, RandomEngine engine) noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence< dim >(), f, iv, n, engine))
Definition AMReX_GpuLaunchFunctsC.H:111
AMREX_FORCE_INLINE auto call_f_intvect_inner(std::index_sequence< Ns... >, F const &f, IntVectND< 1 > iv, Args...args) noexcept -> decltype(f(0, 0, 0, args...))
Definition AMReX_GpuLaunchFunctsC.H:49
Definition AMReX_Amr.cpp:49
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:138
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:191
__host__ __device__ IntVectND< dim > lbound_iv(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1798
void launch(T const &n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:138
void HostDeviceFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:912
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForSIMD(N n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:206
void HostDeviceParallelFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:799
__host__ __device__ IntVectND< dim > ubound_iv(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1807
void single_task(L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1350
AMREX_ATTRIBUTE_FLATTEN_FOR void For(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:152
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1264
Definition AMReX_FabArrayCommI.H:1000
Definition AMReX_GpuTypes.H:86
Definition AMReX_RandomEngine.H:72
Definition AMReX_GpuLaunchFunctsC.H:17
static constexpr int width
Definition AMReX_GpuLaunchFunctsC.H:19
N index
Definition AMReX_GpuLaunchFunctsC.H:22