ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
LBInterpolation.impl.hpp
Go to the documentation of this file.
1/*
2 * Copyright (C) 2019-2026 The ESPResSo project
3 *
4 * This file is part of ESPResSo.
5 *
6 * ESPResSo is free software: you can redistribute it and/or modify
7 * it under the terms of the GNU General Public License as published by
8 * the Free Software Foundation, either version 3 of the License, or
9 * (at your option) any later version.
10 *
11 * ESPResSo is distributed in the hope that it will be useful,
12 * but WITHOUT ANY WARRANTY; without even the implied warranty of
13 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
14 * GNU General Public License for more details.
15 *
16 * You should have received a copy of the GNU General Public License
17 * along with this program. If not, see <http://www.gnu.org/licenses/>.
18 */
19
20#pragma once
21
22/**
23 * @file
24 * Out-of-class position-based interpolation definitions for
25 * @ref walberla::LBWalberlaImpl.
26 */
27
28#include <utils/Vector.hpp>
30
31#include <algorithm>
32#include <array>
33#include <cstddef>
34#include <iostream>
35#include <iterator>
36#include <memory>
37#include <optional>
38#include <stdexcept>
39#include <string>
40#include <type_traits>
41#include <utility>
42#include <vector>
43
44namespace walberla {
45
46/**
47 * @brief Exception for accessing a lattice node outside the local domain
48 * and ghost layers during B-spline interpolation.
49 */
50class interpolation_illegal_access : public std::runtime_error {
51public:
52 interpolation_illegal_access(std::string const &field,
53 Utils::Vector3d const &pos,
54 std::array<int, 3> const &node, double weight)
55 : std::runtime_error("Access to LB " + field + " field failed") {
56 std::cerr << "pos [" << pos << "], node [" << Utils::Vector3i(node)
57 << "], weight " << weight << "\n";
58 }
59};
60
61void interpolate_bspline_at_pos(Utils::Vector3d const &pos, auto const &&f) {
62 Utils::Interpolation::bspline_3d<2>(
63 pos, f, Utils::Vector3d::broadcast(1.), // grid spacing
64 Utils::Vector3d::broadcast(.5)); // offset
65}
66
67template <typename FloatType, lbmpy::Arch Architecture>
68std::function<bool(Utils::Vector3d const &)>
70 bool consider_points_in_halo) const {
71 auto const &lat = *m_lattice;
73 return [&](Utils::Vector3d const &p) { return lat.pos_in_local_halo(p); };
74 }
75 return [&](Utils::Vector3d const &p) { return lat.pos_in_local_domain(p); };
76}
77
78/**
79 * @brief Distribute forces to the lattice at given positions.
80 * Uses B-spline interpolation to spread each force over the surrounding
81 * lattice nodes. On GPU, positions are transformed to block-local
82 * coordinates and the operation is performed in a single kernel launch.
83 */
84template <typename FloatType, lbmpy::Arch Architecture>
86 std::vector<Utils::Vector3d> const &pos,
87 std::vector<Utils::Vector3d> const &forces) {
88 assert(pos.size() == forces.size());
89 if (pos.empty()) {
90 return;
91 }
92 if constexpr (Architecture == lbmpy::Arch::CPU) {
93 auto const kernel = make_force_interpolation_kernel();
94 for (std::size_t i = 0ul; i < pos.size(); ++i) {
95 kernel(pos[i], forces[i]);
96 }
97 }
98#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
99 if constexpr (Architecture == lbmpy::Arch::GPU) {
100 auto const &lattice = get_lattice();
101 auto const &block = *(lattice.get_blocks()->begin());
102 auto const origin = block.getAABB().min();
103 std::vector<FloatType> host_pos;
104 std::vector<FloatType> host_force;
105 host_pos.reserve(3ul * pos.size());
106 host_force.reserve(3ul * forces.size());
107 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
108 for (auto const &vec : pos) {
109#pragma unroll
110 for (std::size_t i : {0ul, 1ul, 2ul}) {
111 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
112 }
113 }
114 for (auto const &vec : forces) {
115#pragma unroll
116 for (std::size_t i : {0ul, 1ul, 2ul}) {
117 host_force.emplace_back(static_cast<FloatType>(vec[i]));
118 }
119 }
120 zero_centered_to_lb_in_place(host_force);
121 auto const gl = lattice.get_ghost_layers();
122 auto field = block.template uncheckedFastGetData<VectorField>(
123 m_force_to_be_applied_id);
125 }
126#endif
127}
128
129template <typename FloatType, lbmpy::Arch Architecture>
131 const {
132 auto const &lattice = *m_lattice;
133 auto const &blocks = *lattice.get_blocks();
134 assert(lattice.get_ghost_layers() == 1u);
135 return [&](Utils::Vector3d const &pos, Utils::Vector3d const &force) {
136 if (not get_block_extended(lattice, pos, 1u)) {
137 return;
138 }
140 pos, [&, conv = m_zc_to_lb, field_id = m_force_to_be_applied_id](
141 std::array<int, 3> const node, double weight) {
142 auto block = get_block_extended(lattice, node, 0u);
143 if (!block)
144 block = get_block_extended(lattice, node, 1u);
145 if (block) {
146 auto cell = to_cell(node);
147 blocks.transformGlobalToBlockLocalCell(cell, *block);
148 weight *= conv;
149 auto const weighted_force = to_vector3<FloatType>(weight * force);
150 auto field =
153 }
154 });
155 };
156}
157
158template <typename FloatType, lbmpy::Arch Architecture>
159auto LBWalberlaImpl<FloatType,
160 Architecture>::make_velocity_interpolation_kernel() const {
161 auto const &lattice = *m_lattice;
162 auto const &blocks = *lattice.get_blocks();
163 assert(lattice.get_ghost_layers() == 1u);
164 return [&](Utils::Vector3d const &pos) {
165 Utils::Vector3d acc{0., 0., 0.};
167 pos, [&, field_id = m_velocity_field_id](std::array<int, 3> const node,
168 double weight) {
169 // Nodes with zero weight might not be accessible, because they can be
170 // outside ghost layers
171 if (weight != 0.) {
172 auto block = get_block_extended(lattice, node, 1u);
173 if (!block)
174 throw interpolation_illegal_access("velocity", pos, node, weight);
176 if (m_has_boundaries and m_boundary->node_is_boundary(node)) {
177 vel = m_boundary->get_node_value_at_boundary(node);
178 } else {
179 auto cell = to_cell(node);
180 blocks.transformGlobalToBlockLocalCell(cell, *block);
181 auto field =
183 vel = lbm::accessor::Vector::get(field, cell);
184 }
185 acc += to_vector3d(vel) * weight;
186 }
187 });
188 return acc;
189 };
190}
191
192template <typename FloatType, lbmpy::Arch Architecture>
193auto LBWalberlaImpl<FloatType,
194 Architecture>::make_density_interpolation_kernel() const {
195 auto const &lattice = *m_lattice;
196 auto const &blocks = *lattice.get_blocks();
197 assert(lattice.get_ghost_layers() == 1u);
198 return [&](Utils::Vector3d const &pos) {
199 double acc = 0.;
201 pos, [&, density = m_density, field_id = m_pdf_field_id](
202 std::array<int, 3> const node, double weight) {
203 // Nodes with zero weight might not be accessible, because they can be
204 // outside ghost layers
205 if (weight != 0.) {
206 auto block = get_block_extended(lattice, node, 1u);
207 if (!block)
208 throw interpolation_illegal_access("density", pos, node, weight);
209 auto cell = to_cell(node);
210 blocks.transformGlobalToBlockLocalCell(cell, *block);
211 auto field =
213 auto const rho = lbm::accessor::Density::get(field, density, cell);
214 acc += rho * weight;
215 }
216 });
217 return acc;
218 };
219}
220
221/**
222 * @brief Interpolate velocities at given positions (batch version).
223 * On GPU, boundary slip velocities are written into the velocity field
224 * before interpolation, since the field has indeterminate values inside
225 * boundary regions.
226 */
227template <typename FloatType, lbmpy::Arch Architecture>
228std::vector<Utils::Vector3d>
230 std::vector<Utils::Vector3d> const &pos) {
231 if (pos.empty()) {
232 return {};
233 }
234 std::vector<Utils::Vector3d> vel{};
235 vel.reserve(pos.size());
236 if constexpr (Architecture == lbmpy::Arch::CPU) {
237 auto const kernel = make_velocity_interpolation_kernel();
238 std::ranges::transform(pos, std::back_inserter(vel), kernel);
239 }
240#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
241 if constexpr (Architecture == lbmpy::Arch::GPU) {
242 auto const &lattice = get_lattice();
243 auto const &block = *(lattice.get_blocks()->begin());
244 auto const origin = block.getAABB().min();
245 std::vector<FloatType> host_pos;
246 host_pos.reserve(3ul * pos.size());
247 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
248 for (auto const &vec : pos) {
249#pragma unroll
250 for (std::size_t i : {0ul, 1ul, 2ul}) {
251 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
252 }
253 }
254 auto const gl = lattice.get_ghost_layers();
255 auto field =
256 block.template uncheckedFastGetData<VectorField>(m_velocity_field_id);
257 // the velocity field has indeterminate values inside boundary regions;
258 // we overwrite them with boundary slip velocities before interpolation
259 auto const [dev_idx, dev_vel] = m_boundary->get_flattened_map_device();
260 if (not dev_idx->empty()) {
262 }
264 for (auto it = res.begin(); it != res.end(); it += 3) {
265 vel.emplace_back(Utils::Vector3d{static_cast<double>(*(it + 0)),
266 static_cast<double>(*(it + 1)),
267 static_cast<double>(*(it + 2))});
268 }
269 }
270#endif
271 return vel;
272}
273
274template <typename FloatType, lbmpy::Arch Architecture>
275std::vector<double>
277 std::vector<Utils::Vector3d> const &pos) {
278 if (pos.empty()) {
279 return {};
280 }
281 std::vector<double> rho{};
282 rho.reserve(pos.size());
283 if constexpr (Architecture == lbmpy::Arch::CPU) {
284 auto const kernel = make_density_interpolation_kernel();
285 std::ranges::transform(pos, std::back_inserter(rho), kernel);
286 }
287#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
288 if constexpr (Architecture == lbmpy::Arch::GPU) {
289 auto const &lattice = get_lattice();
290 auto const &block = *(lattice.get_blocks()->begin());
291 auto const origin = block.getAABB().min();
292 std::vector<FloatType> host_pos;
293 host_pos.reserve(3ul * pos.size());
294 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
295 for (auto const &vec : pos) {
296#pragma unroll
297 for (std::size_t i : {0ul, 1ul, 2ul}) {
298 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
299 }
300 }
301 auto const gl = lattice.get_ghost_layers();
302 auto field = block.template uncheckedFastGetData<PdfField>(m_pdf_field_id);
303 auto res =
305 if constexpr (std::is_same_v<FloatType, double>) {
306 std::swap(rho, res);
307 } else {
308 for (auto const &v : res) {
309 rho.emplace_back(static_cast<double>(v));
310 }
311 }
312 }
313#endif
314 return rho;
315}
316
317template <typename FloatType, lbmpy::Arch Architecture>
318std::optional<Utils::Vector3d>
320 Utils::Vector3d const &pos, bool consider_points_in_halo) const {
321 assert(not m_pending_ghost_comm.test(GhostComm::VEL));
322 assert(not m_pending_ghost_comm.test(GhostComm::UBB));
323 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
324 return std::nullopt;
325 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
326 return std::nullopt;
327 auto const kernel = make_velocity_interpolation_kernel();
328 return {kernel(pos)};
329}
330
331template <typename FloatType, lbmpy::Arch Architecture>
332std::optional<double>
334 Utils::Vector3d const &pos, bool consider_points_in_halo) const {
335 assert(not m_pending_ghost_comm.test(GhostComm::PDF));
336 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
337 return std::nullopt;
338 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
339 return std::nullopt;
340 auto const kernel = make_density_interpolation_kernel();
341 return {kernel(pos)};
342}
343
344template <typename FloatType, lbmpy::Arch Architecture>
346 Utils::Vector3d const &pos, Utils::Vector3d const &force) {
347 if (!m_lattice->pos_in_local_halo(pos))
348 return false;
349 auto const kernel = make_force_interpolation_kernel();
350 kernel(pos, force);
351 return true;
352}
353
354} // namespace walberla
Vector implementation and trait types for boost qvm interoperability.
static DEVICE_QUALIFIER constexpr Vector< T, N > broadcast(typename Base::value_type const &value) noexcept
Create a vector that has all entries set to the same value.
Definition Vector.hpp:131
Class that runs and controls the LB on waLBerla.
void add_forces_at_pos(std::vector< Utils::Vector3d > const &pos, std::vector< Utils::Vector3d > const &forces) override
Distribute forces to the lattice at given positions.
std::optional< double > get_density_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
std::function< bool(Utils::Vector3d const &)> make_lattice_position_checker(bool consider_points_in_halo) const override
std::vector< Utils::Vector3d > get_velocities_at_pos(std::vector< Utils::Vector3d > const &pos) override
Interpolate velocities at given positions (batch version).
std::optional< Utils::Vector3d > get_velocity_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
std::vector< double > get_densities_at_pos(std::vector< Utils::Vector3d > const &pos) override
bool add_force_at_pos(Utils::Vector3d const &pos, Utils::Vector3d const &force) override
Exception for accessing a lattice node outside the local domain and ghost layers during B-spline inte...
interpolation_illegal_access(std::string const &field, Utils::Vector3d const &pos, std::array< int, 3 > const &node, double weight)
cudaStream_t stream[1]
CUDA streams for parallel computing on CPU and GPU.
static double * block(double *p, std::size_t index, std::size_t size)
Definition elc.cpp:177
VectorXi< 3 > Vector3i
Definition Vector.hpp:193
STL namespace.
double get(GhostLayerField< double, uint_t{19u}> const *pdf_field, double const density, Cell const &cell)
void add_force(gpu::GPUField< double > const *field, std::vector< double > const &pos, std::vector< double > const &forces, uint gl)
std::vector< double > get_rho(gpu::GPUField< double > const *field, std::vector< double > const &pos, double const density, uint gl)
std::vector< double > get_vel(gpu::GPUField< double > const *field, std::vector< double > const &pos, uint gl)
void set_from_list(gpu::GPUField< double > const *field, thrust::device_vector< int > const &indices, thrust::device_vector< double > const &values, uint gl)
void add(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec, Cell const &cell)
auto get(GhostLayerField< double, uint_t{3u}> const *vec_field, Cell const &cell)
\file PackInfoPdfDoublePrecision.cpp \author pystencils
IBlock * get_block_extended(LatticeWalberla const &lattice, auto const &pos, unsigned int n_ghost_layers)
auto to_vector3d(Vector3< T > const &v) noexcept
void interpolate_bspline_at_pos(Utils::Vector3d const &pos, auto const &&f)
Cell to_cell(signed_integral_vector auto const &xyz)