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
8namespace detail {
9
10 // call_f_scalar_handler
11
12 template <typename F, typename N>
13 AMREX_FORCE_INLINE
14 auto call_f_scalar_handler (F const& f, N i)
15 noexcept -> decltype(f(0))
16 {
17 return f(i);
18 }
19
20 template <typename F, typename N>
21 AMREX_FORCE_INLINE
22 auto call_f_scalar_handler (F const& f, N i)
23 noexcept -> decltype(f(0,Gpu::Handler{}))
24 {
25 return f(i, Gpu::Handler{});
26 }
27
28 // call_f_intvect_inner
29
30 template <typename F, std::size_t...Ns, class...Args>
31 AMREX_FORCE_INLINE
32 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<1> iv, Args...args)
33 noexcept -> decltype(f(0, 0, 0, args...))
34 {
35 return f(iv[0], 0, 0, args...);
36 }
37
38 template <typename F, std::size_t...Ns, class...Args>
39 AMREX_FORCE_INLINE
40 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<2> iv, Args...args)
41 noexcept -> decltype(f(0, 0, 0, args...))
42 {
43 return f(iv[0], iv[1], 0, args...);
44 }
45
46 template <typename F, int dim, std::size_t...Ns, class...Args>
47 AMREX_FORCE_INLINE
48 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<dim> iv, Args...args)
49 noexcept -> decltype(f(iv, args...))
50 {
51 return f(iv, args...);
52 }
53
54 template <typename F, int dim, std::size_t...Ns, class...Args>
55 AMREX_FORCE_INLINE
56 auto call_f_intvect_inner (std::index_sequence<Ns...>, F const& f, IntVectND<dim> iv, Args...args)
57 noexcept -> decltype(f(iv[Ns]..., args...))
58 {
59 return f(iv[Ns]..., args...);
60 }
61
62 // call_f_intvect_engine
63
64 template <typename F, int dim>
65 AMREX_FORCE_INLINE
66 auto call_f_intvect_engine (F const& f, IntVectND<dim> iv, RandomEngine engine)
67 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, engine))
68 {
69 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, engine);
70 }
71
72 // call_f_intvect_handler
73
74 template <typename F, int dim>
75 AMREX_FORCE_INLINE
76 auto call_f_intvect_handler (F const& f, IntVectND<dim> iv)
77 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv))
78 {
79 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv);
80 }
81
82 template <typename F, int dim>
83 AMREX_FORCE_INLINE
84 auto call_f_intvect_handler (F const& f, IntVectND<dim> iv)
85 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, Gpu::Handler{}))
86 {
87 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, Gpu::Handler{});
88 }
89
90 // call_f_intvect_ncomp_engine
91
92 template <typename F, typename T, int dim>
93 AMREX_FORCE_INLINE
94 auto call_f_intvect_ncomp_engine (F const& f, IntVectND<dim> iv, T n, RandomEngine engine)
95 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, engine))
96 {
97 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, engine);
98 }
99
100 // call_f_intvect_ncomp_handler
101
102 template <typename F, typename T, int dim>
103 AMREX_FORCE_INLINE
104 auto call_f_intvect_ncomp_handler (F const& f, IntVectND<dim> iv, T n)
105 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n))
106 {
107 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n);
108 }
109
110 template <typename F, typename T, int dim>
111 AMREX_FORCE_INLINE
112 auto call_f_intvect_ncomp_handler (F const& f, IntVectND<dim> iv, T n)
113 noexcept -> decltype(call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, Gpu::Handler{}))
114 {
115 return call_f_intvect_inner(std::make_index_sequence<dim>(), f, iv, n, Gpu::Handler{});
116 }
117
118}
120
121template<typename T, typename L>
122void launch (T const& n, L const& f) noexcept
123{
124 // Like the GPU backends: f(i) for each i in [0,n) when n is integral,
125 // f(box) once when n is a BoxND.
126 for (auto const i : Gpu::Range(n)) {
127 f(i);
128 }
129}
130
131template<int MT, typename T, typename L>
132void launch (T const& n, L const& f) noexcept
133{
134 amrex::ignore_unused(MT);
135 for (auto const i : Gpu::Range(n)) {
136 f(i);
137 }
138}
139
140template <std::integral T, typename L >
141AMREX_ATTRIBUTE_FLATTEN_FOR
142void For (T n, L const& f) noexcept
143{
144 for (T i = 0; i < n; ++i) {
145 detail::call_f_scalar_handler(f,i);
146 }
147}
148
149template <int MT, std::integral T, typename L >
150void For (T n, L&& f) noexcept
151{
152 amrex::ignore_unused(MT);
153 For(n, std::forward<L>(f));
154}
155
156template <std::integral T, typename L >
157void For (Gpu::KernelInfo const&, T n, L&& f) noexcept
158{
159 For(n, std::forward<L>(f));
160}
161
162template <int MT, std::integral T, typename L >
163void For (Gpu::KernelInfo const&, T n, L&& f) noexcept
164{
165 amrex::ignore_unused(MT);
166 For(n, std::forward<L>(f));
167}
168
169template <std::integral T, typename L >
170AMREX_ATTRIBUTE_FLATTEN_FOR
171void ParallelFor (T n, L const& f) noexcept
172{
173 AMREX_PRAGMA_SIMD
174 for (T i = 0; i < n; ++i) {
175 detail::call_f_scalar_handler(f,i);
176 }
177}
178
179template <int MT, std::integral T, typename L >
180void ParallelFor (T n, L&& f) noexcept
181{
182 amrex::ignore_unused(MT);
183 ParallelFor(n, std::forward<L>(f));
184}
185
186template <std::integral T, typename L >
187void ParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
188{
189 ParallelFor(n, std::forward<L>(f));
190}
191
192template <int MT, std::integral T, typename L >
193void ParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
194{
195 amrex::ignore_unused(MT);
196 ParallelFor(n, std::forward<L>(f));
197}
198
200namespace detail {
201
202template <int idim, typename L, int dim>
203AMREX_FORCE_INLINE
204void For_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
205{
206 if constexpr (idim == 1) {
207 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
208 call_f_intvect_handler(f,iv);
209 }
210 } else if constexpr (idim == 2) {
211 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
212 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
213 call_f_intvect_handler(f,iv);
214 }}
215 } else if constexpr (idim == 3) {
216 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
217 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
218 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
219 call_f_intvect_handler(f,iv);
220 }}}
221 } else {
222 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
223 For_impND<idim-1>(f, lo, hi, iv);
224 }
225 }
226}
227
228}
230
231template <typename L, int dim>
232AMREX_ATTRIBUTE_FLATTEN_FOR
233void For (BoxND<dim> const& box, L const& f) noexcept
234{
235 const auto lo = amrex::lbound_iv(box);
236 const auto hi = amrex::ubound_iv(box);
237 IntVectND<dim> iv;
238 detail::For_impND<dim>(f, lo, hi, iv);
239}
240
241template <int MT, typename L, int dim>
242void For (BoxND<dim> const& box, L&& f) noexcept
243{
244 amrex::ignore_unused(MT);
245 For(box, std::forward<L>(f));
246}
247
248template <typename L, int dim>
249void For (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
250{
251 For(box, std::forward<L>(f));
252}
253
254template <int MT, typename L, int dim>
255void For (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
256{
257 amrex::ignore_unused(MT);
258 For(box, std::forward<L>(f));
259}
260
262namespace detail {
263
264template <int idim, typename L, int dim>
265AMREX_FORCE_INLINE
266void ParallelFor_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
267{
268 if constexpr (idim == 1) {
269 AMREX_PRAGMA_SIMD
270 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
271 call_f_intvect_handler(f,iv);
272 }
273 } else if constexpr (idim == 2) {
274 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
275 AMREX_PRAGMA_SIMD
276 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
277 call_f_intvect_handler(f,iv);
278 }}
279 } else if constexpr (idim == 3) {
280 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
281 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
282 AMREX_PRAGMA_SIMD
283 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
284 call_f_intvect_handler(f,iv);
285 }}}
286 } else {
287 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
288 ParallelFor_impND<idim-1>(f, lo, hi, iv);
289 }
290 }
291}
292
293}
295
296template <typename L, int dim>
297AMREX_ATTRIBUTE_FLATTEN_FOR
298void ParallelFor (BoxND<dim> const& box, L const& f) noexcept
299{
300 const auto lo = amrex::lbound_iv(box);
301 const auto hi = amrex::ubound_iv(box);
302 IntVectND<dim> iv;
303 detail::ParallelFor_impND<dim>(f, lo, hi, iv);
304}
305
306template <int MT, typename L, int dim>
307void ParallelFor (BoxND<dim> const& box, L&& f) noexcept
308{
309 amrex::ignore_unused(MT);
310 ParallelFor(box, std::forward<L>(f));
311}
312
313template <typename L, int dim>
314void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
315{
316 ParallelFor(box, std::forward<L>(f));
317}
318
319template <int MT, typename L, int dim>
320void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
321{
322 amrex::ignore_unused(MT);
323 ParallelFor(box, std::forward<L>(f));
324}
325
327namespace detail {
328
329template <int idim, typename L, typename T, int dim>
330AMREX_FORCE_INLINE
331void For_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
332{
333 if constexpr (idim == 1) {
334 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
335 call_f_intvect_ncomp_handler(f,iv,n);
336 }
337 } else if constexpr (idim == 2) {
338 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
339 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
340 call_f_intvect_ncomp_handler(f,iv,n);
341 }}
342 } else if constexpr (idim == 3) {
343 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
344 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
345 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
346 call_f_intvect_ncomp_handler(f,iv,n);
347 }}}
348 } else {
349 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
350 For_impND<idim-1>(f, lo, hi, iv, n);
351 }
352 }
353}
354
355}
357
358template <std::integral T, typename L, int dim >
359AMREX_ATTRIBUTE_FLATTEN_FOR
360void For (BoxND<dim> const& box, T ncomp, L const& f) noexcept
361{
362 const auto lo = amrex::lbound_iv(box);
363 const auto hi = amrex::ubound_iv(box);
364 IntVectND<dim> iv;
365 for (T n = 0; n < ncomp; ++n) {
366 detail::For_impND<dim>(f, lo, hi, iv, n);
367 }
368}
369
370template <int MT, std::integral T, typename L, int dim >
371void For (BoxND<dim> const& box, T ncomp, L&& f) noexcept
372{
373 amrex::ignore_unused(MT);
374 For(box, ncomp, std::forward<L>(f));
375}
376
377template <std::integral T, typename L, int dim >
378void For (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
379{
380 For(box, ncomp, std::forward<L>(f));
381}
382
383template <int MT, std::integral T, typename L, int dim >
384void For (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
385{
386 amrex::ignore_unused(MT);
387 For(box, ncomp, std::forward<L>(f));
388}
389
391namespace detail {
392
393template <int idim, typename L, typename T, int dim>
394AMREX_FORCE_INLINE
395void ParallelFor_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
396{
397 if constexpr (idim == 1) {
398 AMREX_PRAGMA_SIMD
399 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
400 call_f_intvect_ncomp_handler(f,iv,n);
401 }
402 } else if constexpr (idim == 2) {
403 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
404 AMREX_PRAGMA_SIMD
405 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
406 call_f_intvect_ncomp_handler(f,iv,n);
407 }}
408 } else if constexpr (idim == 3) {
409 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
410 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
411 AMREX_PRAGMA_SIMD
412 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
413 call_f_intvect_ncomp_handler(f,iv,n);
414 }}}
415 } else {
416 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
417 ParallelFor_impND<idim-1>(f, lo, hi, iv, n);
418 }
419 }
420}
421
422}
424
425template <std::integral T, typename L, int dim >
426AMREX_ATTRIBUTE_FLATTEN_FOR
427void ParallelFor (BoxND<dim> const& box, T ncomp, L const& f) noexcept
428{
429 const auto lo = amrex::lbound_iv(box);
430 const auto hi = amrex::ubound_iv(box);
431 IntVectND<dim> iv;
432 for (T n = 0; n < ncomp; ++n) {
433 detail::ParallelFor_impND<dim>(f, lo, hi, iv, n);
434 }
435}
436
437template <int MT, std::integral T, typename L, int dim >
438void ParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
439{
440 amrex::ignore_unused(MT);
441 ParallelFor(box, ncomp, std::forward<L>(f));
442}
443
444template <std::integral T, typename L, int dim >
445void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
446{
447 ParallelFor(box, ncomp, std::forward<L>(f));
448}
449
450template <int MT, std::integral T, typename L, int dim >
451void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
452{
453 amrex::ignore_unused(MT);
454 ParallelFor(box, ncomp, std::forward<L>(f));
455}
456
457template <typename L1, typename L2, int dim>
458void For (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
459{
460 For(box1, std::forward<L1>(f1));
461 For(box2, std::forward<L2>(f2));
462}
463
464template <int MT, typename L1, typename L2, int dim>
465void For (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
466{
467 amrex::ignore_unused(MT);
468 For(box1, std::forward<L1>(f1));
469 For(box2, std::forward<L2>(f2));
470}
471
472template <typename L1, typename L2, int dim>
473void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
474{
475 For (box1, box2, std::forward<L1>(f1), std::forward<L2>(f2));
476}
477
478template <int MT, typename L1, typename L2, int dim>
479void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
480{
481 amrex::ignore_unused(MT);
482 For (box1, box2, std::forward<L1>(f1), std::forward<L2>(f2));
483}
484
485template <typename L1, typename L2, typename L3, int dim>
486void For (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
487{
488 For(box1, std::forward<L1>(f1));
489 For(box2, std::forward<L2>(f2));
490 For(box3, std::forward<L3>(f3));
491}
492
493template <int MT, typename L1, typename L2, typename L3, int dim>
494void For (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
495{
496 amrex::ignore_unused(MT);
497 For(box1, std::forward<L1>(f1));
498 For(box2, std::forward<L2>(f2));
499 For(box3, std::forward<L3>(f3));
500}
501
502template <typename L1, typename L2, typename L3, int dim>
503void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
504{
505 For(box1, box2, box3, std::forward<L1>(f1), std::forward<L2>(f2), std::forward<L3>(f3));
506}
507
508template <int MT, typename L1, typename L2, typename L3, int dim>
509void For (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
510{
511 amrex::ignore_unused(MT);
512 For(box1, box2, box3, std::forward<L1>(f1), std::forward<L2>(f2), std::forward<L3>(f3));
513}
514
515template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
516void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
517 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
518{
519 For(box1, ncomp1, std::forward<L1>(f1));
520 For(box2, ncomp2, std::forward<L2>(f2));
521}
522
523template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
524void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
525 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
526{
527 amrex::ignore_unused(MT);
528 For(box1, ncomp1, std::forward<L1>(f1));
529 For(box2, ncomp2, std::forward<L2>(f2));
530}
531
532template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
533void For (Gpu::KernelInfo const&,
534 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
535 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
536{
537 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
538}
539
540template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
541void For (Gpu::KernelInfo const&,
542 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
543 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
544{
545 amrex::ignore_unused(MT);
546 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
547}
548
549template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
550void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
551 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
552 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
553{
554 For(box1, ncomp1, std::forward<L1>(f1));
555 For(box2, ncomp2, std::forward<L2>(f2));
556 For(box3, ncomp3, std::forward<L3>(f3));
557}
558
559template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
560void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
561 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
562 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
563{
564 amrex::ignore_unused(MT);
565 For(box1, ncomp1, std::forward<L1>(f1));
566 For(box2, ncomp2, std::forward<L2>(f2));
567 For(box3, ncomp3, std::forward<L3>(f3));
568}
569
570template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
571void For (Gpu::KernelInfo const&,
572 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
573 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
574 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
575{
576 For(box1,ncomp1,std::forward<L1>(f1),
577 box2,ncomp2,std::forward<L2>(f2),
578 box3,ncomp3,std::forward<L3>(f3));
579}
580
581template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
582void For (Gpu::KernelInfo const&,
583 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
584 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
585 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
586{
587 amrex::ignore_unused(MT);
588 For(box1,ncomp1,std::forward<L1>(f1),
589 box2,ncomp2,std::forward<L2>(f2),
590 box3,ncomp3,std::forward<L3>(f3));
591}
592
593template <typename L1, typename L2, int dim>
594void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
595{
596 ParallelFor(box1, std::forward<L1>(f1));
597 ParallelFor(box2, std::forward<L2>(f2));
598}
599
600template <int MT, typename L1, typename L2, int dim>
601void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
602{
603 amrex::ignore_unused(MT);
604 ParallelFor(box1, std::forward<L1>(f1));
605 ParallelFor(box2, std::forward<L2>(f2));
606}
607
608template <typename L1, typename L2, int dim>
609void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
610{
611 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
612}
613
614template <int MT, typename L1, typename L2, int dim>
615void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
616{
617 amrex::ignore_unused(MT);
618 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
619}
620
621template <typename L1, typename L2, typename L3, int dim>
622void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
623{
624 ParallelFor(box1, std::forward<L1>(f1));
625 ParallelFor(box2, std::forward<L2>(f2));
626 ParallelFor(box3, std::forward<L3>(f3));
627}
628
629template <int MT, typename L1, typename L2, typename L3, int dim>
630void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
631{
632 amrex::ignore_unused(MT);
633 ParallelFor(box1, std::forward<L1>(f1));
634 ParallelFor(box2, std::forward<L2>(f2));
635 ParallelFor(box3, std::forward<L3>(f3));
636}
637
638template <typename L1, typename L2, typename L3, int dim>
639void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
640{
641 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
642}
643
644template <int MT, typename L1, typename L2, typename L3, int dim>
645void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
646{
647 amrex::ignore_unused(MT);
648 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
649}
650
651template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
652void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
653 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
654{
655 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
656 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
657}
658
659template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
660void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
661 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
662{
663 amrex::ignore_unused(MT);
664 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
665 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
666}
667
668template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
669void ParallelFor (Gpu::KernelInfo const&,
670 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
671 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
672{
673 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
674 box2,ncomp2,std::forward<L2>(f2));
675}
676
677template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
678void ParallelFor (Gpu::KernelInfo const&,
679 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
680 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
681{
682 amrex::ignore_unused(MT);
683 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
684 box2,ncomp2,std::forward<L2>(f2));
685}
686
687template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
688void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
689 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
690 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
691{
692 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
693 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
694 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
695}
696
697template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
698void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
699 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
700 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
701{
702 amrex::ignore_unused(MT);
703 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
704 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
705 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
706}
707
708template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
709void ParallelFor (Gpu::KernelInfo const&,
710 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
711 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
712 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
713{
714 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
715 box2, ncomp2, std::forward<L2>(f2),
716 box3, ncomp3, std::forward<L3>(f3));
717}
718
719template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
720void ParallelFor (Gpu::KernelInfo const&,
721 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
722 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
723 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
724{
725 amrex::ignore_unused(MT);
726 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
727 box2, ncomp2, std::forward<L2>(f2),
728 box3, ncomp3, std::forward<L3>(f3));
729}
730
731template <std::integral T, typename L >
732void HostDeviceParallelFor (T n, L&& f) noexcept
733{
734 ParallelFor(n,std::forward<L>(f));
735}
736
737template <int MT, std::integral T, typename L >
738void HostDeviceParallelFor (T n, L&& f) noexcept
739{
740 amrex::ignore_unused(MT);
741 ParallelFor(n,std::forward<L>(f));
742}
743
744template <typename L, int dim>
745void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
746{
747 ParallelFor(box,std::forward<L>(f));
748}
749
750template <int MT, typename L, int dim>
751void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
752{
753 amrex::ignore_unused(MT);
754 ParallelFor(box,std::forward<L>(f));
755}
756
757template <std::integral T, typename L, int dim >
758void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
759{
760 ParallelFor(box,ncomp,std::forward<L>(f));
761}
762
763template <int MT, std::integral T, typename L, int dim >
764void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
765{
766 amrex::ignore_unused(MT);
767 ParallelFor(box,ncomp,std::forward<L>(f));
768}
769
770template <typename L1, typename L2, int dim>
771void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
772{
773 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
774}
775
776template <int MT, typename L1, typename L2, int dim>
777void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
778{
779 amrex::ignore_unused(MT);
780 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
781}
782
783template <typename L1, typename L2, typename L3, int dim>
784void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
785 L1&& f1, L2&& f2, L3&& f3) noexcept
786{
787 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
788}
789
790template <int MT, typename L1, typename L2, typename L3, int dim>
791void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
792 L1&& f1, L2&& f2, L3&& f3) noexcept
793{
794 amrex::ignore_unused(MT);
795 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
796}
797
798template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
799void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
800 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
801{
802 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
803}
804
805template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
806void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
807 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
808{
809 amrex::ignore_unused(MT);
810 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
811}
812
813template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
814void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
815 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
816 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
817{
818 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
819 box2,ncomp2,std::forward<L2>(f2),
820 box3,ncomp3,std::forward<L3>(f3));
821}
822
823template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
824void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
825 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
826 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
827{
828 amrex::ignore_unused(MT);
829 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
830 box2,ncomp2,std::forward<L2>(f2),
831 box3,ncomp3,std::forward<L3>(f3));
832}
833
834template <std::integral T, typename L >
835void HostDeviceFor (T n, L&& f) noexcept
836{
837 For(n,std::forward<L>(f));
838}
839
840template <int MT, std::integral T, typename L >
841void HostDeviceFor (T n, L&& f) noexcept
842{
843 amrex::ignore_unused(MT);
844 For(n,std::forward<L>(f));
845}
846
847template <typename L, int dim>
848void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
849{
850 For(box,std::forward<L>(f));
851}
852
853template <int MT, typename L, int dim>
854void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
855{
856 amrex::ignore_unused(MT);
857 For(box,std::forward<L>(f));
858}
859
860template <std::integral T, typename L, int dim >
861void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
862{
863 For(box,ncomp,std::forward<L>(f));
864}
865
866template <int MT, std::integral T, int dim, typename L >
867void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
868{
869 amrex::ignore_unused(MT);
870 For(box,ncomp,std::forward<L>(f));
871}
872
873template <typename L1, typename L2, int dim>
874void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
875{
876 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
877}
878
879template <int MT, typename L1, typename L2, int dim>
880void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
881{
882 amrex::ignore_unused(MT);
883 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
884}
885
886template <typename L1, typename L2, typename L3, int dim>
887void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
888 L1&& f1, L2&& f2, L3&& f3) noexcept
889{
890 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
891}
892
893template <int MT, typename L1, typename L2, typename L3, int dim>
894void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
895 L1&& f1, L2&& f2, L3&& f3) noexcept
896{
897 amrex::ignore_unused(MT);
898 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
899}
900
901template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
902void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
903 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
904{
905 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
906}
907
908template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
909void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
910 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
911{
912 amrex::ignore_unused(MT);
913 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
914}
915
916template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
917void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
918 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
919 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
920{
921 For(box1,ncomp1,std::forward<L1>(f1),
922 box2,ncomp2,std::forward<L2>(f2),
923 box3,ncomp3,std::forward<L3>(f3));
924}
925
926template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
927void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
928 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
929 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
930{
931 amrex::ignore_unused(MT);
932 For(box1,ncomp1,std::forward<L1>(f1),
933 box2,ncomp2,std::forward<L2>(f2),
934 box3,ncomp3,std::forward<L3>(f3));
935}
936
937template <std::integral T, typename L >
938void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
939{
940 ParallelFor(n,std::forward<L>(f));
941}
942
943template <int MT, std::integral T, typename L >
944void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
945{
946 amrex::ignore_unused(MT);
947 ParallelFor(n,std::forward<L>(f));
948}
949
950template <typename L, int dim>
951void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
952{
953 ParallelFor(box,std::forward<L>(f));
954}
955
956template <int MT, typename L, int dim>
957void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
958{
959 amrex::ignore_unused(MT);
960 ParallelFor(box,std::forward<L>(f));
961}
962
963template <std::integral T, typename L, int dim >
964void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
965{
966 ParallelFor(box,ncomp,std::forward<L>(f));
967}
968
969template <int MT, std::integral T, typename L, int dim >
970void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
971{
972 amrex::ignore_unused(MT);
973 ParallelFor(box,ncomp,std::forward<L>(f));
974}
975
976template <typename L1, typename L2, int dim>
977void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
978{
979 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
980}
981
982template <int MT, typename L1, typename L2, int dim>
983void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
984{
985 amrex::ignore_unused(MT);
986 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
987}
988
989template <typename L1, typename L2, typename L3, int dim>
990void HostDeviceParallelFor (Gpu::KernelInfo const&,
991 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
992 L1&& f1, L2&& f2, L3&& f3) noexcept
993{
994 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
995}
996
997template <int MT, typename L1, typename L2, typename L3, int dim>
998void HostDeviceParallelFor (Gpu::KernelInfo const&,
999 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1000 L1&& f1, L2&& f2, L3&& f3) noexcept
1001{
1002 amrex::ignore_unused(MT);
1003 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1004}
1005
1006template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
1007void HostDeviceParallelFor (Gpu::KernelInfo const&,
1008 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1009 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1010{
1011 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1012}
1013
1014template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
1015void HostDeviceParallelFor (Gpu::KernelInfo const&,
1016 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1017 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1018{
1019 amrex::ignore_unused(MT);
1020 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1021}
1022
1023template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
1024void HostDeviceParallelFor (Gpu::KernelInfo const&,
1025 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1026 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1027 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1028{
1029 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1030 box2,ncomp2,std::forward<L2>(f2),
1031 box3,ncomp3,std::forward<L3>(f3));
1032}
1033
1034template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
1035void HostDeviceParallelFor (Gpu::KernelInfo const&,
1036 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1037 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1038 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1039{
1040 amrex::ignore_unused(MT);
1041 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1042 box2,ncomp2,std::forward<L2>(f2),
1043 box3,ncomp3,std::forward<L3>(f3));
1044}
1045
1046template <std::integral T, typename L >
1047void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1048{
1049 For(n,std::forward<L>(f));
1050}
1051
1052template <int MT, std::integral T, typename L >
1053void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1054{
1055 amrex::ignore_unused(MT);
1056 For(n,std::forward<L>(f));
1057}
1058
1059template <typename L, int dim>
1060void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1061{
1062 For(box,std::forward<L>(f));
1063}
1064
1065template <int MT, typename L, int dim>
1066void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1067{
1068 amrex::ignore_unused(MT);
1069 For(box,std::forward<L>(f));
1070}
1071
1072template <std::integral T, typename L, int dim >
1073void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1074{
1075 For(box,ncomp,std::forward<L>(f));
1076}
1077
1078template <int MT, std::integral T, typename L, int dim >
1079void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1080{
1081 amrex::ignore_unused(MT);
1082 For(box,ncomp,std::forward<L>(f));
1083}
1084
1085template <typename L1, typename L2, int dim>
1086void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1087{
1088 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1089}
1090
1091template <int MT, typename L1, typename L2, int dim>
1092void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1093{
1094 amrex::ignore_unused(MT);
1095 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1096}
1097
1098template <typename L1, typename L2, typename L3, int dim>
1099void HostDeviceFor (Gpu::KernelInfo const&,
1100 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1101 L1&& f1, L2&& f2, L3&& f3) noexcept
1102{
1103 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1104}
1105
1106template <int MT, typename L1, typename L2, typename L3, int dim>
1107void HostDeviceFor (Gpu::KernelInfo const&,
1108 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1109 L1&& f1, L2&& f2, L3&& f3) noexcept
1110{
1111 amrex::ignore_unused(MT);
1112 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1113}
1114
1115template <std::integral T1, std::integral T2, typename L1, typename L2, int dim>
1116void HostDeviceFor (Gpu::KernelInfo const&,
1117 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1118 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1119{
1120 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1121}
1122
1123template <int MT, std::integral T1, std::integral T2, typename L1, typename L2, int dim>
1124void HostDeviceFor (Gpu::KernelInfo const&,
1125 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1126 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1127{
1128 amrex::ignore_unused(MT);
1129 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1130}
1131
1132template <std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
1133void HostDeviceFor (Gpu::KernelInfo const&,
1134 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1135 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1136 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1137{
1138 For(box1,ncomp1,std::forward<L1>(f1),
1139 box2,ncomp2,std::forward<L2>(f2),
1140 box3,ncomp3,std::forward<L3>(f3));
1141}
1142
1143template <int MT, std::integral T1, std::integral T2, std::integral T3, typename L1, typename L2, typename L3, int dim>
1144void HostDeviceFor (Gpu::KernelInfo const&,
1145 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1146 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1147 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1148{
1149 amrex::ignore_unused(MT);
1150 For(box1,ncomp1,std::forward<L1>(f1),
1151 box2,ncomp2,std::forward<L2>(f2),
1152 box3,ncomp3,std::forward<L3>(f3));
1153}
1154
1155template <std::integral T, typename L >
1156AMREX_ATTRIBUTE_FLATTEN_FOR
1157void ParallelForRNG (T n, L const& f) noexcept
1158{
1159 for (T i = 0; i < n; ++i) {
1160 f(i,RandomEngine{});
1161 }
1162}
1163
1165namespace detail {
1166
1167template <int idim, typename L, int dim>
1168AMREX_FORCE_INLINE
1169void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
1170{
1171 if constexpr (idim == 1) {
1172 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1173 call_f_intvect_engine(f,iv,RandomEngine{});
1174 }
1175 } else if constexpr (idim == 2) {
1176 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1177 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1178 call_f_intvect_engine(f,iv,RandomEngine{});
1179 }}
1180 } else if constexpr (idim == 3) {
1181 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1182 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1183 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1184 call_f_intvect_engine(f,iv,RandomEngine{});
1185 }}}
1186 } else {
1187 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1188 ParallelForRNG_impND<idim-1>(f, lo, hi, iv);
1189 }
1190 }
1191}
1192
1193template <int idim, typename L, typename T, int dim>
1194AMREX_FORCE_INLINE
1195void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
1196{
1197 if constexpr (idim == 1) {
1198 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1199 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1200 }
1201 } else if constexpr (idim == 2) {
1202 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1203 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1204 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1205 }}
1206 } else if constexpr (idim == 3) {
1207 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1208 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1209 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1210 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1211 }}}
1212 } else {
1213 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1214 ParallelForRNG_impND<idim-1>(f, lo, hi, iv, n);
1215 }
1216 }
1217}
1218
1219}
1221
1222template <typename L, int dim>
1223AMREX_ATTRIBUTE_FLATTEN_FOR
1224void ParallelForRNG (BoxND<dim> const& box, L const& f) noexcept
1225{
1226 const auto lo = amrex::lbound_iv(box);
1227 const auto hi = amrex::ubound_iv(box);
1228 IntVectND<dim> iv;
1229 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv);
1230}
1231
1232template <std::integral T, typename L, int dim >
1233AMREX_ATTRIBUTE_FLATTEN_FOR
1234void ParallelForRNG (BoxND<dim> const& box, T ncomp, L const& f) noexcept
1235{
1236 const auto lo = amrex::lbound_iv(box);
1237 const auto hi = amrex::ubound_iv(box);
1238 IntVectND<dim> iv;
1239 for (T n = 0; n < ncomp; ++n) {
1240 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv, n);
1241 }
1242}
1243
1244template <typename L>
1245void single_task (L&& f) noexcept
1246{
1247 std::forward<L>(f)();
1248}
1249
1250}
1251
1252#endif
AMREX_PRAGMA_SIMD
#define AMREX_PRAGMA_SIMD
Definition AMReX_Extension.H:85
AMREX_FORCE_INLINE
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
AMREX_ATTRIBUTE_FLATTEN_FOR
#define AMREX_ATTRIBUTE_FLATTEN_FOR
Definition AMReX_Extension.H:156
amrex::BoxND
A Rectangular Domain on an Integer Lattice.
Definition AMReX_Box.H:54
amrex::Gpu::KernelInfo
Definition AMReX_GpuKernelInfo.H:8
amrex::IntVectND
An Integer Vector in dim-Dimensional Space.
Definition AMReX_IntVect.H:149
amrex::Gpu::Range
__host__ __device__ range_detail::range_impl< T > Range(T const &b) noexcept
Definition AMReX_GpuRange.H:128
amrex
Definition AMReX_Amr.cpp:50
amrex::ignore_unused
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:259
amrex::HostDeviceParallelFor
void HostDeviceParallelFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:732
amrex::For
void For(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:400
amrex::lbound_iv
__host__ __device__ IntVectND< dim > lbound_iv(BoxND< dim > const &box) noexcept
Return the inclusive lower corner of box as an IntVectND.
Definition AMReX_Box.H:2065
amrex::launch
void launch(T const &n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:122
amrex::ParallelFor
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:202
amrex::HostDeviceFor
void HostDeviceFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:835
amrex::ubound_iv
__host__ __device__ IntVectND< dim > ubound_iv(BoxND< dim > const &box) noexcept
Return the inclusive upper corner of box as an IntVectND.
Definition AMReX_Box.H:2080
amrex::single_task
void single_task(L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1245
amrex::ParallelForRNG
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1157
amrex::RandomEngine
Definition AMReX_RandomEngine.H:72