Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_Algorithm.H
Go to the documentation of this file.
1#ifndef AMREX_ALGORITHM_H_
2#define AMREX_ALGORITHM_H_
3#include <AMReX_Config.H>
4
6#include <AMReX_Extension.H>
7#include <AMReX_Dim3.H>
8#include <AMReX_BLassert.H>
9#include <AMReX_Math.H>
10
11#include <algorithm>
12#include <limits>
13#include <type_traits>
14#include <cstdint>
15#include <climits>
16
17namespace amrex
18{
22 template <class T>
24 AMREX_FORCE_INLINE constexpr const T& min (const T& a, const T& b) noexcept
25 {
26 return std::min(a,b);
27 }
28
32 template <class T, class ... Ts>
34 AMREX_FORCE_INLINE constexpr const T& min (const T& a, const T& b, const Ts& ... c) noexcept
35 {
36 return min(min(a,b),c...);
37 }
38
42 template <class T>
44 AMREX_FORCE_INLINE constexpr const T& max (const T& a, const T& b) noexcept
45 {
46 return std::max(a,b);
47 }
48
52 template <class T, class ... Ts>
54 AMREX_FORCE_INLINE constexpr const T& max (const T& a, const T& b, const Ts& ... c) noexcept
55 {
56 return max(max(a,b),c...);
57 }
58
60 template <class T>
62 T elemwiseMin (T const& a, T const& b) noexcept {
63 return T{amrex::min(a.x,b.x),amrex::min(a.y,b.y),amrex::min(a.z,b.z)};
64 }
65
67 template <class T, class ... Ts>
69 T elemwiseMin (const T& a, const T& b, const Ts& ... c) noexcept
70 {
71 return elemwiseMin(elemwiseMin(a,b),c...);
72 }
73
75 template <class T>
77 T elemwiseMax (T const& a, T const& b) noexcept {
78 return T{amrex::max(a.x,b.x),amrex::max(a.y,b.y),amrex::max(a.z,b.z)};
79 }
80
82 template <class T, class ... Ts>
84 T elemwiseMax (const T& a, const T& b, const Ts& ... c) noexcept
85 {
86 return elemwiseMax(elemwiseMax(a,b),c...);
87 }
88
91 template<typename T>
93 void Swap (T& t1, T& t2) noexcept
94 {
95 T temp = std::move(t1);
96 t1 = std::move(t2);
97 t2 = std::move(temp);
98 }
99
104 template <typename T>
106 constexpr const T& Clamp (const T& v, const T& lo, const T& hi)
107 {
108 if (v < lo) {
109 return lo; // NOLINT(bugprone-return-const-ref-from-parameter)
110 } else if (hi < v) {
111 return hi; // NOLINT(bugprone-return-const-ref-from-parameter)
112 } else {
113 return v; // NOLINT(bugprone-return-const-ref-from-parameter)
114 }
115 }
116
119 template <typename T>
121 std::enable_if_t<std::is_floating_point_v<T>,bool>
122 almostEqual (T x, T y, int ulp = 2)
123 {
124 // Reference: https://en.cppreference.com/w/cpp/types/numeric_limits/epsilon
125 // The machine epsilon has to be scaled to the magnitude of the values used
126 // and multiplied by the desired precision in ULPs (units in the last place)
127 return std::abs(x-y) <= std::numeric_limits<T>::epsilon() * std::abs(x+y) * ulp
128 // unless the result is subnormal
129 || std::abs(x-y) < std::numeric_limits<T>::min();
130 }
131
148 template <class T, class F,
149 std::enable_if_t<std::is_floating_point_v<T>,int>FOO = 0>
151 T bisect (T lo, T hi, F f, T tol=1e-12, int max_iter=100)
152 {
154 "Error - calling bisect but lo and hi don't describe a reasonable interval.");
155
156 T flo = f(lo);
157 T fhi = f(hi);
158
159 if (flo == T(0)) { return flo; }
160 if (fhi == T(0)) { return fhi; }
161
162 AMREX_ASSERT_WITH_MESSAGE(flo * fhi <= T(0),
163 "Error - calling bisect but lo and hi don't bracket a root.");
164
165 T mi = (lo + hi) / T(2);
166 T fmi = T(0);
167 int n = 1;
168 while (n <= max_iter)
169 {
170 if (hi - lo < tol || almostEqual(lo,hi)) { break; }
171 mi = (lo + hi) / T(2);
172 fmi = f(mi);
173 if (fmi == T(0)) { break; }
174 fmi*flo < T(0) ? hi = mi : lo = mi;
175 flo = f(lo);
176 fhi = f(hi);
177 ++n;
178 }
179
180 AMREX_ASSERT_WITH_MESSAGE(n < max_iter,
181 "Error - maximum number of iterations reached in bisect.");
182
183 return mi;
184 }
185
203 template <typename T, typename I,
204 std::enable_if_t<std::is_integral_v<I>,int> = 0>
206 I bisect (T const* d, I lo, I hi, T const& v) {
207 while (lo <= hi) {
208 int mid = lo + (hi-lo)/2;
209 if (v >= d[mid] && v < d[mid+1]) {
210 return mid;
211 } else if (v < d[mid]) {
212 hi = mid-1;
213 } else {
214 lo = mid+1;
215 }
216 };
217 return hi;
218 }
219
234 template<typename ItType, typename ValType>
236 ItType upper_bound (ItType first, ItType last, const ValType& val)
237 {
239 std::ptrdiff_t count = last-first;
240 while(count>0){
241 auto it = first;
242 const auto step = count/2;
243 it += step;
244 if (!(val < *it)){
245 first = ++it;
246 count -= step + 1;
247 }
248 else{
249 count = step;
250 }
251 }
252 return first;
253 ))
255 return std::upper_bound(first, last, val);
256 ))
257 }
258
273 template<typename ItType, typename ValType>
275 ItType lower_bound (ItType first, ItType last, const ValType& val)
276 {
278 std::ptrdiff_t count = last-first;
279 while(count>0)
280 {
281 auto it = first;
282 const auto step = count/2;
283 it += step;
284 if (*it < val){
285 first = ++it;
286 count -= step + 1;
287 }
288 else{
289 count = step;
290 }
291 }
292
293 return first;
294 ))
296 return std::lower_bound(first, last, val);
297 ))
298 }
299
317 template<typename ItType, typename ValType,
318 std::enable_if_t<
319 std::is_floating_point_v<typename std::iterator_traits<ItType>::value_type> &&
320 std::is_floating_point_v<ValType>,
321 int> = 0>
323 void linspace (ItType first, const ItType& last, const ValType& start, const ValType& stop)
324 {
325 const std::ptrdiff_t count = last-first;
326 if (count >= 2){
327 const auto delta = (stop - start)/(count - 1);
328 for (std::ptrdiff_t i = 0; i < count-1; ++i){
329 *(first++) = start + i*delta;
330 }
331 *first = stop;
332 }
333 }
334
353 template<typename ItType, typename ValType,
354 std::enable_if_t<
355 std::is_floating_point_v<typename std::iterator_traits<ItType>::value_type> &&
356 std::is_floating_point_v<ValType>,
357 int> = 0>
359 void logspace (ItType first, const ItType& last,
360 const ValType& start, const ValType& stop, const ValType& base)
361 {
362 const std::ptrdiff_t count = last-first;
363 if (count >= 2){
364 const auto delta = (stop - start)/(count - 1);
365 for (std::ptrdiff_t i = 0; i < count-1; ++i){
366 *(first++) = std::pow(base, start + i*delta);
367 }
368 *first = std::pow(base, stop);
369 }
370 }
371
373namespace detail {
374
375struct clzll_tag {};
376struct clzl_tag : clzll_tag {};
377struct clz_tag : clzl_tag {};
378
379// in gcc and clang, there are three versions of __builtin_clz taking unsigned int,
380// unsigned long, and unsigned long long inputs. Because the sizes of these data types
381// vary on different platforms, we work with fixed-width integer types.
382// these tags and overloads select the smallest version of __builtin_clz that will hold the input type
383template <typename T, typename = std::enable_if_t<sizeof(T) <= sizeof(unsigned int)>>
384AMREX_FORCE_INLINE
385int builtin_clz_wrapper (clz_tag, T x) noexcept
386{
387 return static_cast<int>(__builtin_clz(x) - (sizeof(unsigned int) * CHAR_BIT - sizeof(T) * CHAR_BIT));
388}
389
390template <typename T, typename = std::enable_if_t<sizeof(T) <= sizeof(unsigned long)>>
391AMREX_FORCE_INLINE
392int builtin_clz_wrapper (clzl_tag, T x) noexcept
393{
394 return static_cast<int>(__builtin_clzl(x) - (sizeof(unsigned long) * CHAR_BIT - sizeof(T) * CHAR_BIT));
395}
396
397template <typename T, typename = std::enable_if_t<sizeof(T) <= sizeof(unsigned long long)>>
398AMREX_FORCE_INLINE
399int builtin_clz_wrapper (clzll_tag, T x) noexcept
400{
401 return static_cast<int>(__builtin_clzll(x) - (sizeof(unsigned long long) * CHAR_BIT - sizeof(T) * CHAR_BIT));
402}
403
404}
406
408template <class T, std::enable_if_t<std::is_same_v<std::decay_t<T>,std::uint8_t> ||
409 std::is_same_v<std::decay_t<T>,std::uint16_t> ||
410 std::is_same_v<std::decay_t<T>,std::uint32_t> ||
411 std::is_same_v<std::decay_t<T>,std::uint64_t>, int> = 0>
413int clz (T x) noexcept;
414
416
418int clz_generic (std::uint8_t x) noexcept
419{
420#if !defined(__NVCOMPILER)
421 static constexpr int clz_lookup[16] = { 4, 3, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0 };
422#else
423 constexpr int clz_lookup[16] = { 4, 3, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0 };
424#endif
425 auto upper = x >> 4;
426 auto lower = x & 0xF;
427 return upper ? clz_lookup[upper] : 4 + clz_lookup[lower];
428}
429
431int clz_generic (std::uint16_t x) noexcept
432{
433 auto upper = std::uint8_t(x >> 8);
434 auto lower = std::uint8_t(x & 0xFF);
435 return upper ? clz(upper) : 8 + clz(lower);
436}
437
439int clz_generic (std::uint32_t x) noexcept
440{
441 auto upper = std::uint16_t(x >> 16);
442 auto lower = std::uint16_t(x & 0xFFFF);
443 return upper ? clz(upper) : 16 + clz(lower);
444}
445
447int clz_generic (std::uint64_t x) noexcept
448{
449 auto upper = std::uint32_t(x >> 32);
450 auto lower = std::uint32_t(x & 0xFFFFFFFF);
451 return upper ? clz(upper) : 32 + clz(lower);
452}
453
455
456#if defined AMREX_USE_CUDA
457
459namespace detail {
460 // likewise with CUDA, there are __clz functions that take (signed) int and long long int
461 template <typename T, typename = std::enable_if_t<sizeof(T) <= sizeof(int)> >
462 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
463 int clz_wrapper (clz_tag, T x) noexcept
464 {
465 return __clz((int) x) - (sizeof(int) * CHAR_BIT - sizeof(T) * CHAR_BIT);
466 }
467
468 template <typename T, typename = std::enable_if_t<sizeof(T) <= sizeof(long long int)> >
469 AMREX_GPU_DEVICE AMREX_FORCE_INLINE
470 int clz_wrapper (clzll_tag, T x) noexcept
471 {
472 return __clzll((long long int) x) - (sizeof(long long int) * CHAR_BIT - sizeof(T) * CHAR_BIT);
473 }
474}
476
477template <class T, std::enable_if_t<std::is_same_v<std::decay_t<T>,std::uint8_t> ||
478 std::is_same_v<std::decay_t<T>,std::uint16_t> ||
479 std::is_same_v<std::decay_t<T>,std::uint32_t> ||
480 std::is_same_v<std::decay_t<T>,std::uint64_t>, int> >
482int clz (T x) noexcept
483{
484 AMREX_IF_ON_DEVICE((return detail::clz_wrapper(detail::clz_tag{}, x);))
485#if AMREX_HAS_BUILTIN_CLZ
486 AMREX_IF_ON_HOST((return detail::builtin_clz_wrapper(detail::clz_tag{}, x);))
487#else
488 AMREX_IF_ON_HOST((return clz_generic(x);))
489#endif
490}
491
492#else // !defined AMREX_USE_CUDA
493
494template <class T, std::enable_if_t<std::is_same_v<std::decay_t<T>,std::uint8_t> ||
495 std::is_same_v<std::decay_t<T>,std::uint16_t> ||
496 std::is_same_v<std::decay_t<T>,std::uint32_t> ||
497 std::is_same_v<std::decay_t<T>,std::uint64_t>, int> >
499int clz (T x) noexcept
500{
501#if (!AMREX_DEVICE_COMPILE && AMREX_HAS_BUILTIN_CLZ)
502 return detail::builtin_clz_wrapper(detail::clz_tag{}, x);
503#else
504 return clz_generic(x);
505#endif
506}
507
508#endif // defined AMREX_USE_CUDA
509
510}
511
512#endif
#define AMREX_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:37
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
Definition AMReX_Amr.cpp:49
__host__ __device__ T bisect(T lo, T hi, F f, T tol=1e-12, int max_iter=100)
Find a root of a scalar function on a bracketing interval using bisection.
Definition AMReX_Algorithm.H:151
__host__ __device__ ItType upper_bound(ItType first, ItType last, const ValType &val)
Return an iterator to the first element greater than a given value.
Definition AMReX_Algorithm.H:236
__host__ __device__ void Swap(T &t1, T &t2) noexcept
Definition AMReX_Algorithm.H:93
__host__ __device__ constexpr T elemwiseMax(T const &a, T const &b) noexcept
Return the element-wise maximum of the given values for types like XDim3.
Definition AMReX_Algorithm.H:77
__host__ __device__ constexpr const T & Clamp(const T &v, const T &lo, const T &hi)
Definition AMReX_Algorithm.H:106
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:24
__host__ __device__ ItType lower_bound(ItType first, ItType last, const ValType &val)
Return an iterator to the first element not less than a given value.
Definition AMReX_Algorithm.H:275
__host__ __device__ std::enable_if_t< std::is_floating_point_v< T >, bool > almostEqual(T x, T y, int ulp=2)
Definition AMReX_Algorithm.H:122
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:44
__host__ __device__ void logspace(ItType first, const ItType &last, const ValType &start, const ValType &stop, const ValType &base)
Fill a range with logarithmically spaced values over a closed interval.
Definition AMReX_Algorithm.H:359
__host__ __device__ void linspace(ItType first, const ItType &last, const ValType &start, const ValType &stop)
Fill a range with linearly spaced values over a closed interval.
Definition AMReX_Algorithm.H:323
__host__ __device__ int clz(T x) noexcept
Return the number of leading zeros of the given integer.
Definition AMReX_Algorithm.H:482
__host__ __device__ constexpr T elemwiseMin(T const &a, T const &b) noexcept
Return the element-wise minimum of the given values for types like XDim3.
Definition AMReX_Algorithm.H:62