1/*
2 * zmorton.hpp
3 *
4 * Created on: Jul 31, 2019
5 * Author: i-bird
6 */
7
8#ifndef ZMORTON_HPP_
9#define ZMORTON_HPP_
10
11#include "Grid/grid_key.hpp"
12
13template<typename T>
14inline __device__ __host__ size_t lin_zid(const grid_key_dx<1,T> & key)
15{
16 return key.get(0);
17}
18
19template<typename T>
20inline __device__ __host__ void invlin_zid(size_t lin, grid_key_dx<1,T> & key)
21{
22 return key.set_d(0,lin);
23}
24
25
26template<typename T>
27inline __device__ __host__ size_t lin_zid(const grid_key_dx<2,T> & key)
28{
29 size_t x = key.get(0);
30 size_t y = key.get(1);
31
32
33 x = (x | (x << 16)) & 0x0000FFFF0000FFFF;
34 x = (x | (x << 8)) & 0x00FF00FF00FF00FF;
35 x = (x | (x << 4)) & 0x0F0F0F0F0F0F0F0F;
36 x = (x | (x << 2)) & 0x3333333333333333;
37 x = (x | (x << 1)) & 0x5555555555555555;
38
39 y = (y | (y << 16)) & 0x0000FFFF0000FFFF;
40 y = (y | (y << 8)) & 0x00FF00FF00FF00FF;
41 y = (y | (y << 4)) & 0x0F0F0F0F0F0F0F0F;
42 y = (y | (y << 2)) & 0x3333333333333333;
43 y = (y | (y << 1)) & 0x5555555555555555;
44
45 return x | (y << 1);
46}
47
48template<typename T>
49inline __device__ __host__ void invlin_zid(size_t lin, grid_key_dx<2,T> & key)
50{
51 size_t x = lin & 0x5555555555555555;
52 size_t y = (lin & 0xAAAAAAAAAAAAAAAA) >> 1;
53
54 x = (x | (x >> 1)) & 0x3333333333333333;
55 x = (x | (x >> 2)) & 0x0F0F0F0F0F0F0F0F;
56 x = (x | (x >> 4)) & 0x00FF00FF00FF00FF;
57 x = (x | (x >> 8)) & 0x0000FFFF0000FFFF;
58 x = (x | (x >> 16)) & 0x00000000FFFFFFFF;
59
60 y = (y | (y >> 1)) & 0x3333333333333333;
61 y = (y | (y >> 2)) & 0x0F0F0F0F0F0F0F0F;
62 y = (y | (y >> 4)) & 0x00FF00FF00FF00FF;
63 y = (y | (y >> 8)) & 0x0000FFFF0000FFFF;
64 y = (y | (y >> 16)) & 0x00000000FFFFFFFF;
65
66 key.set_d(0,x);
67 key.set_d(1,y);
68}
69
70static const size_t S3[] = {2, 4, 8, 16, 32};
71
72template<typename T>
73inline __device__ __host__ size_t lin_zid(const grid_key_dx<3,T> & key)
74{
75 size_t x = key.get(0);
76 size_t z = key.get(2);
77 size_t y = key.get(1);
78
79 x = (x | (x << 32)) & 0xFFFF0000FFFFFFFF;
80 x = (x | (x << 16)) & 0x0FFF000FFF000FFF;
81 x = (x | (x << 8)) & 0xF00F00F00F00F00F;
82 x = (x | (x << 4)) & 0x30C30C30C30C30C3;
83 x = (x | (x << 2)) & 0x9249249249249249;
84
85 y = (y | (y << 32)) & 0xFFFF0000FFFFFFFF;
86 y = (y | (y << 16)) & 0x0FFF000FFF000FFF;
87 y = (y | (y << 8)) & 0xF00F00F00F00F00F;
88 y = (y | (y << 4)) & 0x30C30C30C30C30C3;
89 y = (y | (y << 2)) & 0x9249249249249249;
90
91 z = (z | (z << 32)) & 0xFFFF0000FFFFFFFF;
92 z = (z | (z << 16)) & 0x0FFF000FFF000FFF;
93 z = (z | (z << 8)) & 0xF00F00F00F00F00F;
94 z = (z | (z << 4)) & 0x30C30C30C30C30C3;
95 z = (z | (z << 2)) & 0x9249249249249249;
96
97 return x | (y << 1) | (z << 2);
98}
99
100template<typename T>
101inline __device__ __host__ void invlin_zid(size_t lin, grid_key_dx<3,T> & key)
102{
103 size_t x = lin & 0x9249249249249249;
104 size_t y = (lin >> 1) & 0x9249249249249249;
105 size_t z = (lin >> 2) & 0x9249249249249249;
106
107 x = (x | (x >> 2)) & 0x30C30C30C30C30C3;
108 x = (x | (x >> 4)) & 0xF00F00F00F00F00F;
109 x = (x | (x >> 8)) & 0x00FF0000FF0000FF;
110 x = (x | (x >> 16)) & 0x00000FF0000FFFF;
111 x = (x | x >> 16) & 0xFFFFFF;
112
113 y = (y | (y >> 2)) & 0x30C30C30C30C30C3;
114 y = (y | (y >> 4)) & 0xF00F00F00F00F00F;
115 y = (y | (y >> 8)) & 0x00FF0000FF0000FF;
116 y = (y | (y >> 16)) & 0x00000FF0000FFFF;
117 y = (y | y >> 16) & 0xFFFFFF;
118
119 z = (z | (z >> 2)) & 0x30C30C30C30C30C3;
120 z = (z | (z >> 4)) & 0xF00F00F00F00F00F;
121 z = (z | (z >> 8)) & 0x00FF0000FF0000FF;
122 z = (z | (z >> 16)) & 0x00000FF0000FFFF;
123 z = (z | z >> 16) & 0xFFFFFF;
124
125 key.set_d(0,x);
126 key.set_d(1,y);
127 key.set_d(2,z);
128}
129
130#endif /* ZMORTON_HPP_ */
131