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