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