Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_Morton.H
Go to the documentation of this file.
1#ifndef AMREX_MORTON_H_
2#define AMREX_MORTON_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_REAL.H>
6#include <AMReX_Dim3.H>
7#include <AMReX_Algorithm.H>
8#include <AMReX_Array.H>
10
11#include <cstdint>
12#include <limits>
13
14namespace amrex::Morton {
15
33std::uint32_t makeSpace (std::uint32_t x) noexcept {
34#if (AMREX_SPACEDIM == 3)
35 // x : 0000,0000,0000,0000,0000,00a9,8765,4321
36 x = (x | (x << 16)) & 0x030000FF;
37 // x << 16 : 0000,00a9,8765,4321,0000,0000,0000,0000
38 // x | (x << 16): 0000,00a9,8765,4321,0000,00a9,8765,4321
39 // 0x030000FF : 0000,0011,0000,0000,0000,0000,1111,1111
40 // x : 0000,00a9,0000,0000,0000,0000,8765,4321
41 x = (x | (x << 8)) & 0x0300F00F;
42 // x << 8 : 0000,0000,0000,0000,8765,4321,0000,0000
43 // x | (x << 8) : 0000,00a9,0000,0000,8765,4321,8765,4321
44 // 0x0300F00F : 0000,0011,0000,0000,1111,0000,0000,1111
45 // x : 0000,00a9,0000,0000,8765,0000,0000,4321
46 x = (x | (x << 4)) & 0x030C30C3;
47 // x << 4 : 00a9,0000,0000,8765,0000,0000,4321,0000
48 // x | (x << 4) : 00a9,00a9,0000,8765,8765,0000,4321,4321
49 // 0x030C30C3 : 0000,0011,0000,1100,0011,0000,1100,0011
50 // x : 0000,00a9,0000,8700,0065,0000,4300,0021
51 x = (x | (x << 2)) & 0x09249249;
52 // x << 2 : 0000,a900,0087,0000,6500,0043,0000,2100
53 // x | (x << 2) : 0000,a9a9,0087,8700,6565,0043,4300,2121
54 // 0x09249249 : 0000,1001,0010,0100,1001,0010,0100,1001
55 // x : 0000,a009,0080,0700,6005,0040,0300,2001
56 return x;
57#elif (AMREX_SPACEDIM == 2)
58 // x : 0000,0000,0000,0000,gfed,cba9,8765,4321
59 x = (x | (x << 8)) & 0x00FF00FF;
60 // x << 8 : 0000,0000,gfed,cba9,8765,4321,0000,0000
61 // x | (x << 8): 0000,0000,gfed,cba9,????,????,8765,4321
62 // 0x00FF00FF : 0000,0000,1111,1111,0000,0000,1111,1111
63 // x : 0000,0000,gfed,cba9,0000,0000,8765,4321
64 x = (x | (x << 4)) & 0x0F0F0F0F;
65 // x << 4 : 0000,gfed,cba9,0000,0000,8765,4321,0000
66 // x | (x << 4): 0000,gfed,????,cba9,0000,8765,????,4321
67 // 0x0F0F0F0F : 0000,1111,0000,1111,0000,1111,0000,1111
68 // x : 0000,gfed,0000,cba9,0000,8765,0000,4321
69 x = (x | (x << 2)) & 0x33333333;
70 // x << 2 : 00gf,ed00,00cb,a900,0087,6500,0043,2100
71 // x | (x << 2): 00gf,??ed,00cb,??a9,0087,??65,0043,??21
72 // 0x33333333 : 0011,0011,0011,0011,0011,0011,0011,0011
73 // x : 00gf,00ed,00cb,00a9,0087,0065,0043,0021
74 x = (x | (x << 1)) & 0x55555555;
75 // x << 1 : 0gf0,0ed0,0cb0,0a90,0870,0650,0430,0210
76 // x | (x << 1): 0g?f,0e?d,0c?b,0a?9,08?7,06?5,04?3,02?1
77 // 0x55555555 : 0101,0101,0101,0101,0101,0101,0101,0101
78 // x : 0g0f,0e0d,0c0b,0a09,0807,0605,0403,0201
79 return x;
80#elif (AMREX_SPACEDIM == 1)
81 return x;
82#endif
83}
84
100std::uint32_t toUInt10 (amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept {
101 AMREX_ASSERT(x >= xmin && x < xmax);
102 constexpr std::uint32_t code_offset = (1U << 10);
103 constexpr Real rmax = Real(1.0) - std::numeric_limits<Real>::epsilon();
104 Real r = amrex::Clamp((x - xmin)/(xmax - xmin), Real(0.0), rmax);
105 return (std::uint32_t)(r * code_offset);
106}
107
123std::uint32_t toUInt16 (amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept {
124 AMREX_ASSERT(x >= xmin && x < xmax);
125 constexpr std::uint32_t code_offset = (1U << 16);
126 constexpr Real rmax = Real(1.0) - std::numeric_limits<Real>::epsilon();
127 Real r = amrex::Clamp((x - xmin)/(xmax - xmin), Real(0.0), rmax);
128 return (std::uint32_t)(r * code_offset);
129}
130
144std::uint32_t toUInt32 (amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept {
145 AMREX_ASSERT(x >= xmin && x < xmax);
146 constexpr unsigned long long code_offset = (1ULL << 32);
147 constexpr Real rmax = Real(1.0) - std::numeric_limits<Real>::epsilon();
148 Real r = amrex::Clamp((x - xmin)/(xmax - xmin), Real(0.0), rmax);
149 return (std::uint32_t)(r * Real(code_offset));
150}
151
173 const GpuArray<amrex::Real,AMREX_SPACEDIM>& phi) noexcept {
174#if (AMREX_SPACEDIM == 3)
175 std::uint32_t a = toUInt10(x, plo[0], phi[0]);
176 std::uint32_t b = toUInt10(y, plo[1], phi[1]);
177 std::uint32_t c = toUInt10(z, plo[2], phi[2]);
178#elif (AMREX_SPACEDIM == 2)
179 std::uint32_t a = toUInt16(x, plo[0], phi[0]);
180 std::uint32_t b = toUInt16(y, plo[1], phi[1]);
181#elif (AMREX_SPACEDIM == 1)
182 std::uint32_t a = toUInt32(x, plo[0], phi[0]);
183#endif
184 return AMREX_D_TERM(makeSpace(a), | (makeSpace(b) << 1), | (makeSpace(c) << 2));
185}
186
198#if (AMREX_SPACEDIM == 3)
200std::uint32_t get32BitCode (const XDim3& p,
202 const GpuArray<amrex::Real,AMREX_SPACEDIM>& phi) noexcept {
203 return get32BitCode(p.x, p.y, p.z, plo, phi);
204}
205#endif
206
207}
208#endif
General-purpose algorithm utilities available on both host and device.
Fixed-size array types for use on GPU and CPU.
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
Simple structs for 3-component integer and real-valued tuples.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:80
__host__ __device__ constexpr const T & Clamp(const T &v, const T &lo, const T &hi)
Definition AMReX_Algorithm.H:122
Definition AMReX_Morton.H:14
__host__ __device__ std::uint32_t toUInt16(amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept
Convert a Real to a uint32, keeping only 16 significant bits. This puts the first 16 bits of x into r...
Definition AMReX_Morton.H:123
__host__ __device__ std::uint32_t makeSpace(std::uint32_t x) noexcept
This makes space in the input 32-bit integer by splitting the bits so they can be interleaved.
Definition AMReX_Morton.H:33
__host__ __device__ std::uint32_t toUInt10(amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept
Convert a Real to a uint32, keeping only 10 significant bits. This puts the first 10 bits of x into r...
Definition AMReX_Morton.H:100
__host__ __device__ std::uint32_t toUInt32(amrex::Real x, amrex::Real xmin, amrex::Real xmax) noexcept
Convert a Real to a uint32, keeping all significant bits. It is assumed that xmin <= x < xmax....
Definition AMReX_Morton.H:144
__host__ __device__ std::uint32_t get32BitCode(Real x, Real y, Real z, const GpuArray< amrex::Real, 3 > &plo, const GpuArray< amrex::Real, 3 > &phi) noexcept
Given a 3D real-space coordinate, returns a Morton code stored in an unsigned 32 bit integer.
Definition AMReX_Morton.H:171
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
A simple struct holding 3 Real values for a 3D coordinate.
Definition AMReX_Dim3.H:32