Block-Structured AMR Software Framework
 
•All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Modules Pages
Loading...
Searching...
No Matches
AMReX_GpuAtomic.H
Go to the documentation of this file.
1#ifndef AMREX_GPU_ATOMIC_H_
2#define AMREX_GPU_ATOMIC_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_GpuQualifiers.H>
6#include <AMReX_Functional.H>
7#include <AMReX_INT.H>
8
9#include <utility>
10
11namespace amrex {
12
13namespace Gpu::Atomic {
14
15// For Add, Min and Max, we support int, unsigned int, long, unsigned long long, float and double.
16// For Multiply and Divide, we support generic types provided they are the same size as int or unsigned long long
17// and have *= and /= operators.
18// For LogicalOr and LogicalAnd, the data type is int.
19// For Exch and CAS, the data type is generic.
20// All these functions are non-atomic in host code!!!
21// If one needs them to be atomic in host code, use HostDevice::Atomic::*. Currently only
22// HostDevice::Atomic::Add is supported. We could certainly add more.
23
24namespace detail {
25
26#ifdef AMREX_USE_SYCL
27
28 template <typename R, typename I, typename F>
29 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
30 R atomic_op (R* const address, R const val, F const f) noexcept
31 {
32#if defined(__SYCL_DEVICE_ONLY__)
33 constexpr auto mo = sycl::memory_order::relaxed;
34 constexpr auto ms = sycl::memory_scope::device;
35 constexpr auto as = sycl::access::address_space::global_space;
36 static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
37 I* const add_as_I = reinterpret_cast<I*>(address);
38 sycl::atomic_ref<I,mo,ms,as> a{*add_as_I};
39 I old_I = a.load(), new_I;
40 do {
41 R const new_R = f(*(reinterpret_cast<R const*>(&old_I)), val);
42 new_I = *(reinterpret_cast<I const*>(&new_R));
43 } while (! a.compare_exchange_strong(old_I, new_I));
44 return *(reinterpret_cast<R const*>(&old_I));
45#else
46 R old = *address;
47 *address = f(*address, val);
48 return old;
49#endif
50 }
51
52 template <typename R, typename I, typename Op, typename Cond>
53 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
54 bool atomic_op_if (R* const address, R const val, Op&& op, Cond&& cond) noexcept
55 {
56#if defined(__SYCL_DEVICE_ONLY__)
57 constexpr auto mo = sycl::memory_order::relaxed;
58 constexpr auto ms = sycl::memory_scope::device;
59 constexpr auto as = sycl::access::address_space::global_space;
60 static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
61 I* const add_as_I = reinterpret_cast<I*>(address);
62 sycl::atomic_ref<I,mo,ms,as> a{*add_as_I};
63 I old_I = a.load(), new_I;
64 bool test_success;
65 do {
66 R const tmp = op(*(reinterpret_cast<R const*>(&old_I)), val);
67 new_I = *(reinterpret_cast<I const*>(&tmp));
68 test_success = cond(tmp);
69 } while (test_success && ! a.compare_exchange_strong(old_I, new_I));
70 return test_success;
71#else
72 R old = *address;
73 R const tmp = op(*(reinterpret_cast<R const*>(&old)), val);
74 if (cond(tmp)) {
75 *address = tmp;
76 return true;
77 } else {
78 return false;
79 }
80#endif
81 }
82
83#elif defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
84
85 template <typename R, typename I, typename F>
86 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
87 R atomic_op (R* const address, R const val, F const f) noexcept
88 {
89 static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
90 I* const add_as_I = reinterpret_cast<I*>(address);
91 I old_I = *add_as_I, assumed_I;
92 do {
93 assumed_I = old_I;
94 R const new_R = f(*(reinterpret_cast<R const*>(&assumed_I)), val);
95 old_I = atomicCAS(add_as_I, assumed_I, *(reinterpret_cast<I const*>(&new_R)));
96 } while (assumed_I != old_I);
97 return *(reinterpret_cast<R const*>(&old_I));
98 }
99
100 template <typename R, typename I, typename Op, typename Cond>
101 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
102 bool atomic_op_if (R* const address, R const val, Op&& op, Cond&& cond) noexcept
103 {
104 static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
105 I* const add_as_I = reinterpret_cast<I*>(address);
106 I old_I = *add_as_I, assumed_I;
107 bool test_success;
108 do {
109 assumed_I = old_I;
110 R const new_R = op(*(reinterpret_cast<R const*>(&assumed_I)), val);
111 test_success = cond(new_R);
112 if (test_success) {
113 old_I = atomicCAS(add_as_I, assumed_I, *(reinterpret_cast<I const*>(&new_R)));
114 }
115 } while (test_success && assumed_I != old_I);
116 return test_success;
117 }
118
119#endif
120
121}
122
124// Add
126
127#ifdef AMREX_USE_GPU
128
129#ifdef AMREX_USE_SYCL
130 template<class T, sycl::access::address_space AS = sycl::access::address_space::global_space>
131#else
132 template<class T>
133#endif
134 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
135 T Add_device (T* const sum, T const value) noexcept
136 {
137#if defined(__SYCL_DEVICE_ONLY__)
138 constexpr auto mo = sycl::memory_order::relaxed;
139 constexpr auto ms = sycl::memory_scope::device;
140 sycl::atomic_ref<T,mo,ms,AS> a{*sum};
141 return a.fetch_add(value);
142#else
143 AMREX_IF_ON_DEVICE(( return atomicAdd(sum, value); ))
144 AMREX_IF_ON_HOST((
145 amrex::ignore_unused(sum, value);
146 return T(); // should never get here, but have to return something
147 ))
148#endif
149 }
150
151#if defined(AMREX_USE_HIP) && defined(__gfx90a__)
152 // https://github.com/ROCm-Developer-Tools/hipamd/blob/rocm-4.5.x/include/hip/amd_detail/amd_hip_unsafe_atomics.h
153 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
154 float Add_device (float* const sum, float const value) noexcept
155 {
156 return unsafeAtomicAdd(sum, value);
157 }
158
159 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
160 double Add_device (double* const sum, double const value) noexcept
161 {
162 return unsafeAtomicAdd(sum, value);
163 }
164#endif
165
166#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
167
168 // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd
169 // https://rocmdocs.amd.com/en/latest/Programming_Guides/Kernel_language.html?#atomic-functions
170 // CUDA and HIP support int, unsigned int, and unsigned long long.
171
172 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
173 Long Add_device (Long* const sum, Long const value) noexcept
174 {
175 return detail::atomic_op<Long, unsigned long long>(sum, value, amrex::Plus<Long>());
176 }
177
178#endif
179
180#if defined(AMREX_USE_CUDA) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 600)
181
182 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
183 double Add_device (double* const sum, double const value) noexcept
184 {
185 return detail::atomic_op<double, unsigned long long>(sum, value, amrex::Plus<double>());
186 }
187
188#endif
189
190#endif
191
192#ifdef AMREX_USE_SYCL
193 template<class T, sycl::access::address_space AS = sycl::access::address_space::global_space>
194#else
195 template<class T>
196#endif
197 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
198 T Add (T* sum, T value) noexcept
199 {
200#ifdef AMREX_USE_SYCL
201 AMREX_IF_ON_DEVICE((return Add_device<T,AS>(sum, value);))
202#else
203 AMREX_IF_ON_DEVICE((return Add_device(sum, value);))
204#endif
205 AMREX_IF_ON_HOST((
206 auto old = *sum;
207 *sum += value;
208 return old;
209 ))
210 }
211
213// If
215
216#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) || defined(AMREX_USE_SYCL)
217 template <typename T, typename Op, typename Cond,
218 std::enable_if_t<sizeof(T) == sizeof(unsigned int), int> foo = 0>
219 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
220 bool If_device (T* const sum, T const value, Op&& op, Cond&& cond) noexcept
221 {
222 return detail::atomic_op_if<T, unsigned int>(sum, value,
223 std::forward<Op>(op), std::forward<Cond>(cond));
224 }
225
226 template <typename T, typename Op, typename Cond,
227 std::enable_if_t<sizeof(T) == sizeof(unsigned long long), int> foo = 0>
228 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
229 bool If_device (T* const sum, T const value, Op&& op, Cond&& cond) noexcept
230 {
231 return detail::atomic_op_if<T, unsigned long long>(sum, value,
232 std::forward<Op>(op), std::forward<Cond>(cond));
233 }
234#endif
235
252 template<class T, class Op, class Cond>
253 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
254 bool If (T* const add, T const value, Op&& op, Cond&& cond) noexcept
255 {
256 AMREX_IF_ON_DEVICE((
257 return If_device(add, value, std::forward<Op>(op), std::forward<Cond>(cond));
258 ))
259 AMREX_IF_ON_HOST((
260 T old = *add;
261 T const tmp = std::forward<Op>(op)(old, value);
262 if (std::forward<Cond>(cond)(tmp)) {
263 *add = tmp;
264 return true;
265 } else {
266 return false;
267 }
268 ))
269 }
270
272// AddNoRet
274
275#ifdef AMREX_USE_SYCL
276 template<class T, sycl::access::address_space AS = sycl::access::address_space::global_space>
277#else
278 template<class T>
279#endif
280 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
281 void AddNoRet (T* sum, T value) noexcept
282 {
283#if defined(__SYCL_DEVICE_ONLY__)
284 Add_device<T,AS>(sum, value);
285#else
286 AMREX_IF_ON_DEVICE((Add_device(sum, value);))
287 AMREX_IF_ON_HOST((*sum += value;))
288#endif
289 }
290
291#if defined(AMREX_USE_HIP) && defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR < 5)
292 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
293 void AddNoRet (float* const sum, float const value) noexcept
294 {
295#pragma clang diagnostic push
296#pragma clang diagnostic ignored "-Wdeprecated-declarations"
297 AMREX_IF_ON_DEVICE((atomicAddNoRet(sum, value);))
298#pragma clang diagnostic pop
299 AMREX_IF_ON_HOST((*sum += value;))
300 }
301#endif
302
304// Min
306
307#ifdef AMREX_USE_GPU
308
309 template<class T>
310 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
311 T Min_device (T* const m, T const value) noexcept
312 {
313#if defined(__SYCL_DEVICE_ONLY__)
314 constexpr auto mo = sycl::memory_order::relaxed;
315 constexpr auto ms = sycl::memory_scope::device;
316 constexpr auto as = sycl::access::address_space::global_space;
317 sycl::atomic_ref<T,mo,ms,as> a{*m};
318 return a.fetch_min(value);
319#else
320 AMREX_IF_ON_DEVICE(( return atomicMin(m, value); ))
321 AMREX_IF_ON_HOST((
322 amrex::ignore_unused(m,value);
323 return T(); // should never get here, but have to return something
324 ))
325#endif
326 }
327
328 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
329 float Min_device (float* const m, float const value) noexcept
330 {
331 return detail::atomic_op<float,int>(m,value,amrex::Minimum<float>());
332 }
333
334 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
335 double Min_device (double* const m, double const value) noexcept
336 {
337 return detail::atomic_op<double, unsigned long long int>(m,value,amrex::Minimum<double>());
338 }
339
340#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
341
342 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
343 Long Min_device (Long* const m, Long const value) noexcept
344 {
345 return detail::atomic_op<Long, unsigned long long int>(m,value,amrex::Minimum<Long>());
346 }
347
348#endif
349
350#endif
351
352 template<class T>
353 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
354 T Min (T* const m, T const value) noexcept
355 {
356 AMREX_IF_ON_DEVICE((
357 return Min_device(m, value);
358 ))
359 AMREX_IF_ON_HOST((
360 auto const old = *m;
361 *m = (*m) < value ? (*m) : value;
362 return old;
363 ))
364 }
365
367// Max
369
370#ifdef AMREX_USE_GPU
371
372 template<class T>
373 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
374 T Max_device (T* const m, T const value) noexcept
375 {
376#if defined(__SYCL_DEVICE_ONLY__)
377 constexpr auto mo = sycl::memory_order::relaxed;
378 constexpr auto ms = sycl::memory_scope::device;
379 constexpr auto as = sycl::access::address_space::global_space;
380 sycl::atomic_ref<T,mo,ms,as> a{*m};
381 return a.fetch_max(value);
382#else
383 AMREX_IF_ON_DEVICE(( return atomicMax(m, value); ))
384 AMREX_IF_ON_HOST((
385 amrex::ignore_unused(m,value);
386 return T(); // should never get here, but have to return something
387 ))
388#endif
389 }
390
391 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
392 float Max_device (float* const m, float const value) noexcept
393 {
394 return detail::atomic_op<float,int>(m,value,amrex::Maximum<float>());
395 }
396
397 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
398 double Max_device (double* const m, double const value) noexcept
399 {
400 return detail::atomic_op<double, unsigned long long int>(m,value,amrex::Maximum<double>());
401 }
402
403#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
404
405 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
406 Long Max_device (Long* const m, Long const value) noexcept
407 {
408 return detail::atomic_op<Long, unsigned long long int>(m,value,amrex::Maximum<Long>());
409 }
410
411#endif
412
413#endif
414
415 template<class T>
416 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
417 T Max (T* const m, T const value) noexcept
418 {
419 AMREX_IF_ON_DEVICE((
420 return Max_device(m, value);
421 ))
422 AMREX_IF_ON_HOST((
423 auto const old = *m;
424 *m = (*m) > value ? (*m) : value;
425 return old;
426 ))
427 }
428
430// LogicalOr
432
433 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
434 int LogicalOr (int* const m, int const value) noexcept
435 {
436#if defined(__SYCL_DEVICE_ONLY__)
437 constexpr auto mo = sycl::memory_order::relaxed;
438 constexpr auto ms = sycl::memory_scope::device;
439 constexpr auto as = sycl::access::address_space::global_space;
440 sycl::atomic_ref<int,mo,ms,as> a{*m};
441 return a.fetch_or(value);
442#else
443 AMREX_IF_ON_DEVICE((
444 return atomicOr(m, value);
445 ))
446 AMREX_IF_ON_HOST((
447 int const old = *m;
448 *m = (*m) || value;
449 return old;
450 ))
451#endif
452 }
453
455// LogicalAnd
457
458 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
459 int LogicalAnd (int* const m, int const value) noexcept
460 {
461#if defined(__SYCL_DEVICE_ONLY__)
462 constexpr auto mo = sycl::memory_order::relaxed;
463 constexpr auto ms = sycl::memory_scope::device;
464 constexpr auto as = sycl::access::address_space::global_space;
465 sycl::atomic_ref<int,mo,ms,as> a{*m};
466 return a.fetch_and(value ? ~0x0 : 0);
467#else
468 AMREX_IF_ON_DEVICE((
469 return atomicAnd(m, value ? ~0x0 : 0);
470 ))
471 AMREX_IF_ON_HOST((
472 int const old = *m;
473 *m = (*m) && value;
474 return old;
475 ))
476#endif
477 }
478
480// Exch
482
483 template <typename T>
484 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
485 T Exch (T* address, T val) noexcept
486 {
487#if defined(__SYCL_DEVICE_ONLY__)
488 constexpr auto mo = sycl::memory_order::relaxed;
489 constexpr auto ms = sycl::memory_scope::device;
490 constexpr auto as = sycl::access::address_space::global_space;
491 sycl::atomic_ref<T,mo,ms,as> a{*address};
492 return a.exchange(val);
493#else
494 AMREX_IF_ON_DEVICE((
495 return atomicExch(address, val);
496 ))
497 AMREX_IF_ON_HOST((
498 auto const old = *address;
499 *address = val;
500 return old;
501 ))
502#endif
503 }
504
506// CAS
508
509 template <typename T>
510 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
511 T CAS (T* const address, T compare, T const val) noexcept
512 { // cannot be T const compare because of compare_exchange_strong
513#if defined(__SYCL_DEVICE_ONLY__)
514 constexpr auto mo = sycl::memory_order::relaxed;
515 constexpr auto ms = sycl::memory_scope::device;
516 constexpr auto as = sycl::access::address_space::global_space;
517 sycl::atomic_ref<T,mo,ms,as> a{*address};
518 a.compare_exchange_strong(compare, val);
519 return compare;
520#else
521 AMREX_IF_ON_DEVICE((
522 return atomicCAS(address, compare, val);
523 ))
524 AMREX_IF_ON_HOST((
525 auto const old = *address;
526 *address = (old == compare ? val : old);
527 return old;
528 ))
529#endif
530 }
531
533// Multiply
535
536#ifdef AMREX_USE_GPU
537
538 template <typename T, std::enable_if_t<sizeof(T) == sizeof(int), int> = 0>
539 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
540 T Multiply_device (T* const prod, T const value) noexcept
541 {
542 return detail::atomic_op<T, int>(prod,value,amrex::Multiplies<T>());
543 }
544
545 template <typename T, std::enable_if_t<sizeof(T) == sizeof(unsigned long long), int> = 0>
546 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
547 T Multiply_device (T* const prod, T const value) noexcept
548 {
549 return detail::atomic_op<T, unsigned long long>(prod,value,amrex::Multiplies<T>());
550 }
551
552#endif
553
554 template<class T>
555 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
556 T Multiply (T* const prod, T const value) noexcept
557 {
558 AMREX_IF_ON_DEVICE((
559 return Multiply_device(prod, value);
560 ))
561 AMREX_IF_ON_HOST((
562 auto const old = *prod;
563 *prod *= value;
564 return old;
565 ))
566 }
567
569// Divide
571
572#ifdef AMREX_USE_GPU
573
574 template <typename T, std::enable_if_t<sizeof(T) == sizeof(int), int> = 0>
575 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
576 T Divide_device (T* const quot, T const value) noexcept
577 {
578 return detail::atomic_op<T, int>(quot,value,amrex::Divides<T>());
579 }
580
581 template <typename T, std::enable_if_t<sizeof(T) == sizeof(unsigned long long), int> = 0>
582 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
583 T Divide_device (T* const quot, T const value) noexcept
584 {
585 return detail::atomic_op<T, unsigned long long>(quot,value,amrex::Divides<T>());
586 }
587
588#endif
589
590 template<class T>
591 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
592 T Divide (T* const quot, T const value) noexcept
593 {
594 AMREX_IF_ON_DEVICE((
595 return Divide_device(quot, value);
596 ))
597 AMREX_IF_ON_HOST((
598 auto const old = *quot;
599 *quot /= value;
600 return old;
601 ))
602 }
603}
604
605namespace HostDevice::Atomic {
606
607 template <class T>
608 AMREX_FORCE_INLINE
609 void Add_Host (T* const sum, T const value) noexcept
610 {
611#ifdef AMREX_USE_OMP
612#pragma omp atomic update
613#endif
614 *sum += value;
615 }
616
617 template <class T>
618 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
619 void Add (T* const sum, T const value) noexcept
620 {
621 AMREX_IF_ON_DEVICE((Gpu::Atomic::AddNoRet(sum,value);))
622 AMREX_IF_ON_HOST((Add_Host(sum,value);))
623 }
624
625}
626
627#ifdef AMREX_USE_GPU
628// functors
629namespace Gpu {
630 template <typename T>
631 struct AtomicAdd
632 {
633 AMREX_GPU_DEVICE void operator() (T* const dest, T const source) noexcept {
634 Gpu::Atomic::AddNoRet(dest, source);
635 }
636 };
637
638 template <typename T>
639 struct AtomicMin
640 {
641 AMREX_GPU_DEVICE void operator() (T* const dest, T const source) noexcept {
642 Gpu::Atomic::Min(dest, source);
643 }
644 };
645
646 template <typename T>
647 struct AtomicMax
648 {
649 AMREX_GPU_DEVICE void operator() (T* const dest, T const source) noexcept {
650 Gpu::Atomic::Max(dest, source);
651 }
652 };
653
654 template <typename T>
655 struct AtomicLogicalAnd
656 {
657 AMREX_GPU_DEVICE void operator() (T* const dest, T const source) noexcept {
658 Gpu::Atomic::LogicalAnd(dest, source);
659 }
660 };
661
662 template <typename T>
663 struct AtomicLogicalOr
664 {
665 AMREX_GPU_DEVICE void operator() (T* const dest, T const source) noexcept {
666 Gpu::Atomic::LogicalOr(dest, source);
667 }
668 };
669}
670#endif
671
672}
673#endif
AMREX_FORCE_INLINE
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
AMReX_Functional.H
AMReX_GpuQualifiers.H
AMREX_IF_ON_DEVICE
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
AMREX_IF_ON_HOST
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
AMREX_GPU_DEVICE
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
AMREX_GPU_HOST_DEVICE
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
AMReX_INT.H
amrex::Gpu::Atomic::detail::atomic_op
AMREX_GPU_DEVICE AMREX_FORCE_INLINE R atomic_op(R *const address, R const val, F const f) noexcept
Definition AMReX_GpuAtomic.H:87
amrex::Gpu::Atomic::detail::atomic_op_if
AMREX_GPU_DEVICE AMREX_FORCE_INLINE bool atomic_op_if(R *const address, R const val, Op &&op, Cond &&cond) noexcept
Definition AMReX_GpuAtomic.H:102
amrex::Gpu::Atomic::Exch
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Exch(T *address, T val) noexcept
Definition AMReX_GpuAtomic.H:485
amrex::Gpu::Atomic::If_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE bool If_device(T *const sum, T const value, Op &&op, Cond &&cond) noexcept
Definition AMReX_GpuAtomic.H:220
amrex::Gpu::Atomic::CAS
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T CAS(T *const address, T compare, T const val) noexcept
Definition AMReX_GpuAtomic.H:511
amrex::Gpu::Atomic::Add
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Add(T *sum, T value) noexcept
Definition AMReX_GpuAtomic.H:198
amrex::Gpu::Atomic::Max
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Max(T *const m, T const value) noexcept
Definition AMReX_GpuAtomic.H:417
amrex::Gpu::Atomic::LogicalOr
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE int LogicalOr(int *const m, int const value) noexcept
Definition AMReX_GpuAtomic.H:434
amrex::Gpu::Atomic::Multiply
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Multiply(T *const prod, T const value) noexcept
Definition AMReX_GpuAtomic.H:556
amrex::Gpu::Atomic::Divide_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE T Divide_device(T *const quot, T const value) noexcept
Definition AMReX_GpuAtomic.H:576
amrex::Gpu::Atomic::LogicalAnd
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE int LogicalAnd(int *const m, int const value) noexcept
Definition AMReX_GpuAtomic.H:459
amrex::Gpu::Atomic::Min
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Min(T *const m, T const value) noexcept
Definition AMReX_GpuAtomic.H:354
amrex::Gpu::Atomic::Divide
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Divide(T *const quot, T const value) noexcept
Definition AMReX_GpuAtomic.H:592
amrex::Gpu::Atomic::Add_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE T Add_device(T *const sum, T const value) noexcept
Definition AMReX_GpuAtomic.H:135
amrex::Gpu::Atomic::AddNoRet
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void AddNoRet(T *sum, T value) noexcept
Definition AMReX_GpuAtomic.H:281
amrex::Gpu::Atomic::If
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool If(T *const add, T const value, Op &&op, Cond &&cond) noexcept
Conditionally perform an atomic operation.
Definition AMReX_GpuAtomic.H:254
amrex::Gpu::Atomic::Min_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE T Min_device(T *const m, T const value) noexcept
Definition AMReX_GpuAtomic.H:311
amrex::Gpu::Atomic::Max_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE T Max_device(T *const m, T const value) noexcept
Definition AMReX_GpuAtomic.H:374
amrex::Gpu::Atomic::Multiply_device
AMREX_GPU_DEVICE AMREX_FORCE_INLINE T Multiply_device(T *const prod, T const value) noexcept
Definition AMReX_GpuAtomic.H:540
amrex::HostDevice::Atomic::Add
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void Add(T *const sum, T const value) noexcept
Definition AMReX_GpuAtomic.H:619
amrex::HostDevice::Atomic::Add_Host
AMREX_FORCE_INLINE void Add_Host(T *const sum, T const value) noexcept
Definition AMReX_GpuAtomic.H:609
amrex
Definition AMReX_Amr.cpp:49
amrex::ignore_unused
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:127
detail
Definition AMReX_FabArrayCommI.H:896
amrex::Divides
Definition AMReX_Functional.H:77
amrex::Gpu::AtomicAdd
Definition AMReX_GpuAtomic.H:632
amrex::Gpu::AtomicAdd::operator()
AMREX_GPU_DEVICE void operator()(T *const dest, T const source) noexcept
Definition AMReX_GpuAtomic.H:633
amrex::Gpu::AtomicLogicalAnd
Definition AMReX_GpuAtomic.H:656
amrex::Gpu::AtomicLogicalAnd::operator()
AMREX_GPU_DEVICE void operator()(T *const dest, T const source) noexcept
Definition AMReX_GpuAtomic.H:657
amrex::Gpu::AtomicLogicalOr
Definition AMReX_GpuAtomic.H:664
amrex::Gpu::AtomicLogicalOr::operator()
AMREX_GPU_DEVICE void operator()(T *const dest, T const source) noexcept
Definition AMReX_GpuAtomic.H:665
amrex::Gpu::AtomicMax
Definition AMReX_GpuAtomic.H:648
amrex::Gpu::AtomicMax::operator()
AMREX_GPU_DEVICE void operator()(T *const dest, T const source) noexcept
Definition AMReX_GpuAtomic.H:649
amrex::Gpu::AtomicMin
Definition AMReX_GpuAtomic.H:640
amrex::Gpu::AtomicMin::operator()
AMREX_GPU_DEVICE void operator()(T *const dest, T const source) noexcept
Definition AMReX_GpuAtomic.H:641
amrex::LogicalAnd
Definition AMReX_Functional.H:50
amrex::LogicalOr
Definition AMReX_Functional.H:59
amrex::Maximum
Definition AMReX_Functional.H:41
amrex::Minimum
Definition AMReX_Functional.H:32
amrex::Multiplies
Definition AMReX_Functional.H:68
amrex::Plus
Definition AMReX_Functional.H:14