26 static constexpr u32 splts_count = 1U << dim;
28 std::array<u64, dim> coord_min;
29 std::array<u64, dim> coord_max;
33 PatchCoord(std::array<u64, dim> coord_min, std::array<u64, dim> coord_max)
34 : coord_min(coord_min), coord_max(coord_max) {}
36 [[nodiscard]]
inline static auto get_split_coord(
37 std::array<u64, dim> coord_min, std::array<u64, dim> coord_max)
38 -> std::array<u64, dim> {
40 (((coord_max[0] - coord_min[0]) + 1) / 2) - 1 + coord_min[0],
41 (((coord_max[1] - coord_min[1]) + 1) / 2) - 1 + coord_min[1],
42 (((coord_max[2] - coord_min[2]) + 1) / 2) - 1 + coord_min[2]};
45 inline static auto get_split(std::array<u64, dim> coord_min, std::array<u64, dim> coord_max)
46 -> std::array<PatchCoord, splts_count> {
48 std::array<PatchCoord, splts_count> pret;
50 auto splts = get_split_coord(coord_min, coord_max);
52 u64 split_x = splts[0];
53 u64 split_y = splts[1];
54 u64 split_z = splts[2];
56 pret[0].coord_min[0] = coord_min[0];
57 pret[0].coord_min[1] = coord_min[1];
58 pret[0].coord_min[2] = coord_min[2];
59 pret[0].coord_max[0] = split_x;
60 pret[0].coord_max[1] = split_y;
61 pret[0].coord_max[2] = split_z;
63 pret[1].coord_min[0] = coord_min[0];
64 pret[1].coord_min[1] = coord_min[1];
65 pret[1].coord_min[2] = split_z + 1;
66 pret[1].coord_max[0] = split_x;
67 pret[1].coord_max[1] = split_y;
68 pret[1].coord_max[2] = coord_max[2];
70 pret[2].coord_min[0] = coord_min[0];
71 pret[2].coord_min[1] = split_y + 1;
72 pret[2].coord_min[2] = coord_min[2];
73 pret[2].coord_max[0] = split_x;
74 pret[2].coord_max[1] = coord_max[1];
75 pret[2].coord_max[2] = split_z;
77 pret[3].coord_min[0] = coord_min[0];
78 pret[3].coord_min[1] = split_y + 1;
79 pret[3].coord_min[2] = split_z + 1;
80 pret[3].coord_max[0] = split_x;
81 pret[3].coord_max[1] = coord_max[1];
82 pret[3].coord_max[2] = coord_max[2];
84 pret[4].coord_min[0] = split_x + 1;
85 pret[4].coord_min[1] = coord_min[1];
86 pret[4].coord_min[2] = coord_min[2];
87 pret[4].coord_max[0] = coord_max[0];
88 pret[4].coord_max[1] = split_y;
89 pret[4].coord_max[2] = split_z;
91 pret[5].coord_min[0] = split_x + 1;
92 pret[5].coord_min[1] = coord_min[1];
93 pret[5].coord_min[2] = split_z + 1;
94 pret[5].coord_max[0] = coord_max[0];
95 pret[5].coord_max[1] = split_y;
96 pret[5].coord_max[2] = coord_max[2];
98 pret[6].coord_min[0] = split_x + 1;
99 pret[6].coord_min[1] = split_y + 1;
100 pret[6].coord_min[2] = coord_min[2];
101 pret[6].coord_max[0] = coord_max[0];
102 pret[6].coord_max[1] = coord_max[1];
103 pret[6].coord_max[2] = split_z;
105 pret[7].coord_min[0] = split_x + 1;
106 pret[7].coord_min[1] = split_y + 1;
107 pret[7].coord_min[2] = split_z + 1;
108 pret[7].coord_max[0] = coord_max[0];
109 pret[7].coord_max[1] = coord_max[1];
110 pret[7].coord_max[2] = coord_max[2];
115 inline auto split() -> std::array<PatchCoord, splts_count> {
116 return get_split(coord_min, coord_max);
121 {sycl::min(c1.coord_min[0], c2.coord_min[0]),
122 sycl::min(c1.coord_min[1], c2.coord_min[1]),
123 sycl::min(c1.coord_min[2], c2.coord_min[2])},
124 {sycl::max(c1.coord_max[0], c2.coord_max[0]),
125 sycl::max(c1.coord_max[1], c2.coord_max[1]),
126 sycl::max(c1.coord_max[2], c2.coord_max[2])});
129 inline static PatchCoord merge(std::array<PatchCoord, splts_count> others) {
131 merge(merge(others[0], others[1]), merge(others[2], others[3])),
132 merge(merge(others[4], others[5]), merge(others[6], others[7])));
137 u64_3{coord_min[0], coord_min[1], coord_min[2]},
138 u64_3{coord_max[0], coord_max[1], coord_max[2]} + 1};
142 inline static std::tuple<sycl::vec<T, 3>, sycl::vec<T, 3>> convert_coord(
143 std::array<u64, dim> coord_min,
144 std::array<u64, dim> coord_max,
145 std::array<u64, dim> pcoord_offset,
146 sycl::vec<T, 3> divfact,
147 sycl::vec<T, 3> offset) {
149 using vec = sycl::vec<T, 3>;
151 vec min_bound = vec{coord_min[0] - pcoord_offset[0],
152 coord_min[1] - pcoord_offset[1],
153 coord_min[2] - pcoord_offset[2]}
156 vec max_bound = (vec{coord_max[0] - pcoord_offset[0],
157 coord_max[1] - pcoord_offset[1],
158 coord_max[2] - pcoord_offset[2]}
163 return {min_bound, max_bound};