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 <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
144void For (T n, L&& f) noexcept
145{
146 amrex::ignore_unused(MT);
147 For(n, std::forward<L>(f));
148}
149
150template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
151void For (Gpu::KernelInfo const&, T n, L&& f) noexcept
152{
153 For(n, std::forward<L>(f));
154}
155
156template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
174void ParallelFor (T n, L&& f) noexcept
175{
176 amrex::ignore_unused(MT);
177 ParallelFor(n, std::forward<L>(f));
178}
179
180template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
181void ParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
182{
183 ParallelFor(n, std::forward<L>(f));
184}
185
186template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
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 <typename T1, typename T2, typename L1, typename L2, int dim,
510 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
511 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
512void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
513 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
514{
515 For(box1, ncomp1, std::forward<L1>(f1));
516 For(box2, ncomp2, std::forward<L2>(f2));
517}
518
519template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
520 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
521 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
522void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
523 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
524{
525 amrex::ignore_unused(MT);
526 For(box1, ncomp1, std::forward<L1>(f1));
527 For(box2, ncomp2, std::forward<L2>(f2));
528}
529
530template <typename T1, typename T2, typename L1, typename L2, int dim,
531 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
532 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
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, typename T1, typename T2, typename L1, typename L2, int dim,
541 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
542 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
543void For (Gpu::KernelInfo const&,
544 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
545 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
546{
547 amrex::ignore_unused(MT);
548 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
549}
550
551template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
552 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
553 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
554 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
555void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
556 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
557 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
558{
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 <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
565 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
566 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
567 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
568void For (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
569 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
570 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
571{
572 amrex::ignore_unused(MT);
573 For(box1, ncomp1, std::forward<L1>(f1));
574 For(box2, ncomp2, std::forward<L2>(f2));
575 For(box3, ncomp3, std::forward<L3>(f3));
576}
577
578template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
579 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
580 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
581 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
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 For(box1,ncomp1,std::forward<L1>(f1),
588 box2,ncomp2,std::forward<L2>(f2),
589 box3,ncomp3,std::forward<L3>(f3));
590}
591
592template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
593 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
594 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
595 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
596void For (Gpu::KernelInfo const&,
597 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
598 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
599 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
600{
601 amrex::ignore_unused(MT);
602 For(box1,ncomp1,std::forward<L1>(f1),
603 box2,ncomp2,std::forward<L2>(f2),
604 box3,ncomp3,std::forward<L3>(f3));
605}
606
607template <typename L1, typename L2, int dim>
608void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
609{
610 ParallelFor(box1, std::forward<L1>(f1));
611 ParallelFor(box2, std::forward<L2>(f2));
612}
613
614template <int MT, typename L1, typename L2, int dim>
615void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
616{
617 amrex::ignore_unused(MT);
618 ParallelFor(box1, std::forward<L1>(f1));
619 ParallelFor(box2, std::forward<L2>(f2));
620}
621
622template <typename L1, typename L2, int dim>
623void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
624{
625 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
626}
627
628template <int MT, typename L1, typename L2, int dim>
629void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
630{
631 amrex::ignore_unused(MT);
632 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
633}
634
635template <typename L1, typename L2, typename L3, int dim>
636void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
637{
638 ParallelFor(box1, std::forward<L1>(f1));
639 ParallelFor(box2, std::forward<L2>(f2));
640 ParallelFor(box3, std::forward<L3>(f3));
641}
642
643template <int MT, typename L1, typename L2, typename L3, int dim>
644void ParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
645{
646 amrex::ignore_unused(MT);
647 ParallelFor(box1, std::forward<L1>(f1));
648 ParallelFor(box2, std::forward<L2>(f2));
649 ParallelFor(box3, std::forward<L3>(f3));
650}
651
652template <typename L1, typename L2, typename L3, int dim>
653void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
654{
655 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
656}
657
658template <int MT, typename L1, typename L2, typename L3, int dim>
659void ParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3, L1&& f1, L2&& f2, L3&& f3) noexcept
660{
661 amrex::ignore_unused(MT);
662 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
663}
664
665template <typename T1, typename T2, typename L1, typename L2, int dim,
666 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
667 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
668void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
669 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
670{
671 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
672 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
673}
674
675template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
676 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
677 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
678void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
679 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
680{
681 amrex::ignore_unused(MT);
682 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
683 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
684}
685
686template <typename T1, typename T2, typename L1, typename L2, int dim,
687 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
688 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
689void ParallelFor (Gpu::KernelInfo const&,
690 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
691 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
692{
693 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
694 box2,ncomp2,std::forward<L2>(f2));
695}
696
697template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
698 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
699 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
700void ParallelFor (Gpu::KernelInfo const&,
701 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
702 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
703{
704 amrex::ignore_unused(MT);
705 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
706 box2,ncomp2,std::forward<L2>(f2));
707}
708
709template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
710 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
711 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
712 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
713void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
714 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
715 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
716{
717 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
718 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
719 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
720}
721
722template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
723 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
724 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
725 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
726void ParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
727 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
728 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
729{
730 amrex::ignore_unused(MT);
731 ParallelFor(box1, ncomp1, std::forward<L1>(f1));
732 ParallelFor(box2, ncomp2, std::forward<L2>(f2));
733 ParallelFor(box3, ncomp3, std::forward<L3>(f3));
734}
735
736template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
737 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
738 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
739 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
740void ParallelFor (Gpu::KernelInfo const&,
741 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
742 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
743 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
744{
745 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
746 box2, ncomp2, std::forward<L2>(f2),
747 box3, ncomp3, std::forward<L3>(f3));
748}
749
750template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
751 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
752 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
753 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
754void ParallelFor (Gpu::KernelInfo const&,
755 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
756 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
757 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
758{
759 amrex::ignore_unused(MT);
760 ParallelFor(box1, ncomp1, std::forward<L1>(f1),
761 box2, ncomp2, std::forward<L2>(f2),
762 box3, ncomp3, std::forward<L3>(f3));
763}
764
765template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
766void HostDeviceParallelFor (T n, L&& f) noexcept
767{
768 ParallelFor(n,std::forward<L>(f));
769}
770
771template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
772void HostDeviceParallelFor (T n, L&& f) noexcept
773{
774 amrex::ignore_unused(MT);
775 ParallelFor(n,std::forward<L>(f));
776}
777
778template <typename L, int dim>
779void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
780{
781 ParallelFor(box,std::forward<L>(f));
782}
783
784template <int MT, typename L, int dim>
785void HostDeviceParallelFor (BoxND<dim> const& box, L&& f) noexcept
786{
787 amrex::ignore_unused(MT);
788 ParallelFor(box,std::forward<L>(f));
789}
790
791template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
792void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
793{
794 ParallelFor(box,ncomp,std::forward<L>(f));
795}
796
797template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
798void HostDeviceParallelFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
799{
800 amrex::ignore_unused(MT);
801 ParallelFor(box,ncomp,std::forward<L>(f));
802}
803
804template <typename L1, typename L2, int dim>
805void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
806{
807 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
808}
809
810template <int MT, typename L1, typename L2, int dim>
811void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
812{
813 amrex::ignore_unused(MT);
814 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
815}
816
817template <typename L1, typename L2, typename L3, int dim>
818void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
819 L1&& f1, L2&& f2, L3&& f3) noexcept
820{
821 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
822}
823
824template <int MT, typename L1, typename L2, typename L3, int dim>
825void HostDeviceParallelFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
826 L1&& f1, L2&& f2, L3&& f3) noexcept
827{
828 amrex::ignore_unused(MT);
829 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
830}
831
832template <typename T1, typename T2, typename L1, typename L2, int dim,
833 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
834 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
835void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
836 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
837{
838 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
839}
840
841template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
842 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
843 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
844void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
845 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
846{
847 amrex::ignore_unused(MT);
848 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
849}
850
851template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
852 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
853 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
854 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
855void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
856 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
857 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
858{
859 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
860 box2,ncomp2,std::forward<L2>(f2),
861 box3,ncomp3,std::forward<L3>(f3));
862}
863
864template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
865 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
866 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
867 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
868void HostDeviceParallelFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
869 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
870 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
871{
872 amrex::ignore_unused(MT);
873 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
874 box2,ncomp2,std::forward<L2>(f2),
875 box3,ncomp3,std::forward<L3>(f3));
876}
877
878template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
879void HostDeviceFor (T n, L&& f) noexcept
880{
881 For(n,std::forward<L>(f));
882}
883
884template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
885void HostDeviceFor (T n, L&& f) noexcept
886{
887 amrex::ignore_unused(MT);
888 For(n,std::forward<L>(f));
889}
890
891template <typename L, int dim>
892void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
893{
894 For(box,std::forward<L>(f));
895}
896
897template <int MT, typename L, int dim>
898void HostDeviceFor (BoxND<dim> const& box, L&& f) noexcept
899{
900 amrex::ignore_unused(MT);
901 For(box,std::forward<L>(f));
902}
903
904template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
905void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
906{
907 For(box,ncomp,std::forward<L>(f));
908}
909
910template <int MT, typename T, int dim, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
911void HostDeviceFor (BoxND<dim> const& box, T ncomp, L&& f) noexcept
912{
913 amrex::ignore_unused(MT);
914 For(box,ncomp,std::forward<L>(f));
915}
916
917template <typename L1, typename L2, int dim>
918void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
919{
920 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
921}
922
923template <int MT, typename L1, typename L2, int dim>
924void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
925{
926 amrex::ignore_unused(MT);
927 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
928}
929
930template <typename L1, typename L2, typename L3, int dim>
931void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
932 L1&& f1, L2&& f2, L3&& f3) noexcept
933{
934 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
935}
936
937template <int MT, typename L1, typename L2, typename L3, int dim>
938void HostDeviceFor (BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
939 L1&& f1, L2&& f2, L3&& f3) noexcept
940{
941 amrex::ignore_unused(MT);
942 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
943}
944
945template <typename T1, typename T2, typename L1, typename L2, int dim,
946 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
947 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
948void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
949 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
950{
951 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
952}
953
954template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
955 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
956 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
957void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
958 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
959{
960 amrex::ignore_unused(MT);
961 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
962}
963
964template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
965 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
966 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
967 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
968void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
969 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
970 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
971{
972 For(box1,ncomp1,std::forward<L1>(f1),
973 box2,ncomp2,std::forward<L2>(f2),
974 box3,ncomp3,std::forward<L3>(f3));
975}
976
977template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
978 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
979 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
980 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
981void HostDeviceFor (BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
982 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
983 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
984{
985 amrex::ignore_unused(MT);
986 For(box1,ncomp1,std::forward<L1>(f1),
987 box2,ncomp2,std::forward<L2>(f2),
988 box3,ncomp3,std::forward<L3>(f3));
989}
990
991template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
992void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
993{
994 ParallelFor(n,std::forward<L>(f));
995}
996
997template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
998void HostDeviceParallelFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
999{
1000 amrex::ignore_unused(MT);
1001 ParallelFor(n,std::forward<L>(f));
1002}
1003
1004template <typename L, int dim>
1005void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1006{
1007 ParallelFor(box,std::forward<L>(f));
1008}
1009
1010template <int MT, typename L, int dim>
1011void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1012{
1013 amrex::ignore_unused(MT);
1014 ParallelFor(box,std::forward<L>(f));
1015}
1016
1017template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1018void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1019{
1020 ParallelFor(box,ncomp,std::forward<L>(f));
1021}
1022
1023template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1024void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1025{
1026 amrex::ignore_unused(MT);
1027 ParallelFor(box,ncomp,std::forward<L>(f));
1028}
1029
1030template <typename L1, typename L2, int dim>
1031void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1032{
1033 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1034}
1035
1036template <int MT, typename L1, typename L2, int dim>
1037void HostDeviceParallelFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1038{
1039 amrex::ignore_unused(MT);
1040 ParallelFor(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1041}
1042
1043template <typename L1, typename L2, typename L3, int dim>
1044void HostDeviceParallelFor (Gpu::KernelInfo const&,
1045 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1046 L1&& f1, L2&& f2, L3&& f3) noexcept
1047{
1048 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1049}
1050
1051template <int MT, typename L1, typename L2, typename L3, int dim>
1052void HostDeviceParallelFor (Gpu::KernelInfo const&,
1053 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1054 L1&& f1, L2&& f2, L3&& f3) noexcept
1055{
1056 amrex::ignore_unused(MT);
1057 ParallelFor(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1058}
1059
1060template <typename T1, typename T2, typename L1, typename L2, int dim,
1061 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1062 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1063void HostDeviceParallelFor (Gpu::KernelInfo const&,
1064 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1065 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1066{
1067 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1068}
1069
1070template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
1071 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1072 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1073void HostDeviceParallelFor (Gpu::KernelInfo const&,
1074 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1075 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1076{
1077 amrex::ignore_unused(MT);
1078 ParallelFor(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1079}
1080
1081template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1082 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1083 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1084 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1085void HostDeviceParallelFor (Gpu::KernelInfo const&,
1086 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1087 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1088 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1089{
1090 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1091 box2,ncomp2,std::forward<L2>(f2),
1092 box3,ncomp3,std::forward<L3>(f3));
1093}
1094
1095template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1096 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1097 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1098 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1099void HostDeviceParallelFor (Gpu::KernelInfo const&,
1100 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1101 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1102 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1103{
1104 amrex::ignore_unused(MT);
1105 ParallelFor(box1,ncomp1,std::forward<L1>(f1),
1106 box2,ncomp2,std::forward<L2>(f2),
1107 box3,ncomp3,std::forward<L3>(f3));
1108}
1109
1110template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1111void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1112{
1113 For(n,std::forward<L>(f));
1114}
1115
1116template <int MT, typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1117void HostDeviceFor (Gpu::KernelInfo const&, T n, L&& f) noexcept
1118{
1119 amrex::ignore_unused(MT);
1120 For(n,std::forward<L>(f));
1121}
1122
1123template <typename L, int dim>
1124void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1125{
1126 For(box,std::forward<L>(f));
1127}
1128
1129template <int MT, typename L, int dim>
1130void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, L&& f) noexcept
1131{
1132 amrex::ignore_unused(MT);
1133 For(box,std::forward<L>(f));
1134}
1135
1136template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1137void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1138{
1139 For(box,ncomp,std::forward<L>(f));
1140}
1141
1142template <int MT, typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1143void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box, T ncomp, L&& f) noexcept
1144{
1145 amrex::ignore_unused(MT);
1146 For(box,ncomp,std::forward<L>(f));
1147}
1148
1149template <typename L1, typename L2, int dim>
1150void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1151{
1152 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1153}
1154
1155template <int MT, typename L1, typename L2, int dim>
1156void HostDeviceFor (Gpu::KernelInfo const&, BoxND<dim> const& box1, BoxND<dim> const& box2, L1&& f1, L2&& f2) noexcept
1157{
1158 amrex::ignore_unused(MT);
1159 For(box1,box2,std::forward<L1>(f1),std::forward<L2>(f2));
1160}
1161
1162template <typename L1, typename L2, typename L3, int dim>
1163void HostDeviceFor (Gpu::KernelInfo const&,
1164 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1165 L1&& f1, L2&& f2, L3&& f3) noexcept
1166{
1167 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1168}
1169
1170template <int MT, typename L1, typename L2, typename L3, int dim>
1171void HostDeviceFor (Gpu::KernelInfo const&,
1172 BoxND<dim> const& box1, BoxND<dim> const& box2, BoxND<dim> const& box3,
1173 L1&& f1, L2&& f2, L3&& f3) noexcept
1174{
1175 amrex::ignore_unused(MT);
1176 For(box1,box2,box3,std::forward<L1>(f1),std::forward<L2>(f2),std::forward<L3>(f3));
1177}
1178
1179template <typename T1, typename T2, typename L1, typename L2, int dim,
1180 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1181 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1182void HostDeviceFor (Gpu::KernelInfo const&,
1183 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1184 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1185{
1186 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1187}
1188
1189template <int MT, typename T1, typename T2, typename L1, typename L2, int dim,
1190 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1191 typename M2=std::enable_if_t<std::is_integral_v<T2>> >
1192void HostDeviceFor (Gpu::KernelInfo const&,
1193 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1194 BoxND<dim> const& box2, T2 ncomp2, L2&& f2) noexcept
1195{
1196 amrex::ignore_unused(MT);
1197 For(box1,ncomp1,std::forward<L1>(f1),box2,ncomp2,std::forward<L2>(f2));
1198}
1199
1200template <typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1201 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1202 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1203 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1204void HostDeviceFor (Gpu::KernelInfo const&,
1205 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1206 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1207 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1208{
1209 For(box1,ncomp1,std::forward<L1>(f1),
1210 box2,ncomp2,std::forward<L2>(f2),
1211 box3,ncomp3,std::forward<L3>(f3));
1212}
1213
1214template <int MT, typename T1, typename T2, typename T3, typename L1, typename L2, typename L3, int dim,
1215 typename M1=std::enable_if_t<std::is_integral_v<T1>>,
1216 typename M2=std::enable_if_t<std::is_integral_v<T2>>,
1217 typename M3=std::enable_if_t<std::is_integral_v<T3>> >
1218void HostDeviceFor (Gpu::KernelInfo const&,
1219 BoxND<dim> const& box1, T1 ncomp1, L1&& f1,
1220 BoxND<dim> const& box2, T2 ncomp2, L2&& f2,
1221 BoxND<dim> const& box3, T3 ncomp3, L3&& f3) noexcept
1222{
1223 amrex::ignore_unused(MT);
1224 For(box1,ncomp1,std::forward<L1>(f1),
1225 box2,ncomp2,std::forward<L2>(f2),
1226 box3,ncomp3,std::forward<L3>(f3));
1227}
1228
1229template <typename T, typename L, typename M=std::enable_if_t<std::is_integral_v<T>> >
1230AMREX_ATTRIBUTE_FLATTEN_FOR
1231void ParallelForRNG (T n, L const& f) noexcept
1232{
1233 for (T i = 0; i < n; ++i) {
1234 f(i,RandomEngine{});
1235 }
1236}
1237
1239namespace detail {
1240
1241template <int idim, typename L, int dim>
1242AMREX_FORCE_INLINE
1243void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv) noexcept
1244{
1245 if constexpr (idim == 1) {
1246 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1247 call_f_intvect_engine(f,iv,RandomEngine{});
1248 }
1249 } else if constexpr (idim == 2) {
1250 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1251 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1252 call_f_intvect_engine(f,iv,RandomEngine{});
1253 }}
1254 } else if constexpr (idim == 3) {
1255 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1256 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1257 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1258 call_f_intvect_engine(f,iv,RandomEngine{});
1259 }}}
1260 } else {
1261 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1262 ParallelForRNG_impND<idim-1>(f, lo, hi, iv);
1263 }
1264 }
1265}
1266
1267template <int idim, typename L, typename T, int dim>
1268AMREX_FORCE_INLINE
1269void ParallelForRNG_impND (L const& f, IntVectND<dim> const lo, IntVectND<dim> const hi, IntVectND<dim> iv, T n) noexcept
1270{
1271 if constexpr (idim == 1) {
1272 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1273 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1274 }
1275 } else if constexpr (idim == 2) {
1276 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1277 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1278 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1279 }}
1280 } else if constexpr (idim == 3) {
1281 for (int i2 = lo[2], h2 = hi[2]; i2 <= h2; ++i2) { iv[2] = i2;
1282 for (int i1 = lo[1], h1 = hi[1]; i1 <= h1; ++i1) { iv[1] = i1;
1283 for (int i0 = lo[0], h0 = hi[0]; i0 <= h0; ++i0) { iv[0] = i0;
1284 call_f_intvect_ncomp_engine(f,iv,n,RandomEngine{});
1285 }}}
1286 } else {
1287 for (int id = lo[idim-1], hd = hi[idim-1]; id <= hd; ++id) { iv[idim-1] = id;
1288 ParallelForRNG_impND<idim-1>(f, lo, hi, iv, n);
1289 }
1290 }
1291}
1292
1293}
1295
1296template <typename L, int dim>
1297AMREX_ATTRIBUTE_FLATTEN_FOR
1298void ParallelForRNG (BoxND<dim> const& box, L const& f) noexcept
1299{
1300 const auto lo = amrex::lbound_iv(box);
1301 const auto hi = amrex::ubound_iv(box);
1302 IntVectND<dim> iv;
1303 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv);
1304}
1305
1306template <typename T, typename L, int dim, typename M=std::enable_if_t<std::is_integral_v<T>> >
1307AMREX_ATTRIBUTE_FLATTEN_FOR
1308void ParallelForRNG (BoxND<dim> const& box, T ncomp, L const& f) noexcept
1309{
1310 const auto lo = amrex::lbound_iv(box);
1311 const auto hi = amrex::ubound_iv(box);
1312 IntVectND<dim> iv;
1313 for (T n = 0; n < ncomp; ++n) {
1314 detail::ParallelForRNG_impND<dim>(f, lo, hi, iv, n);
1315 }
1316}
1317
1318template <typename L>
1319void single_task (L&& f) noexcept
1320{
1321 std::forward<L>(f)();
1322}
1323
1324}
1325
1326#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:57
amrex
Definition AMReX_Amr.cpp:49
amrex::ignore_unused
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:139
amrex::ParallelFor
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:193
amrex::lbound_iv
__host__ __device__ IntVectND< dim > lbound_iv(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1905
amrex::launch
void launch(T const &n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:122
amrex::HostDeviceFor
void HostDeviceFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:879
amrex::HostDeviceParallelFor
void HostDeviceParallelFor(T n, L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:766
amrex::ubound_iv
__host__ __device__ IntVectND< dim > ubound_iv(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1914
amrex::single_task
void single_task(L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1319
amrex::For
AMREX_ATTRIBUTE_FLATTEN_FOR void For(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:136
amrex::ParallelForRNG
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1231
amrex::RandomEngine
Definition AMReX_RandomEngine.H:72