ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
LBWalberlaImpl.hpp
Go to the documentation of this file.
1/*
2 * Copyright (C) 2019-2023 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 * @ref walberla::LBWalberlaImpl implements the interface of the LB
25 * waLBerla bridge using sweeps generated by lbmpy
26 * (see <tt>maintainer/walberla_kernels</tt>).
27 */
28
29#include <blockforest/Initialization.h>
30#include <blockforest/StructuredBlockForest.h>
31#include <blockforest/communication/UniformBufferedScheme.h>
32#include <domain_decomposition/BlockDataID.h>
33#include <domain_decomposition/IBlock.h>
34#include <field/AddToStorage.h>
35#include <field/GhostLayerField.h>
36#include <field/communication/PackInfo.h>
37#include <field/vtk/FlagFieldCellFilter.h>
38#include <field/vtk/VTKWriter.h>
39#include <stencil/D3Q19.h>
40#include <stencil/D3Q27.h>
41#if defined(__CUDACC__)
42#include <gpu/AddGPUFieldToStorage.h>
43#include <gpu/communication/MemcpyPackInfo.h>
44#include <gpu/communication/UniformGPUScheme.h>
45#endif
46
47#include "../BoundaryHandling.hpp"
48#include "../BoundaryPackInfo.hpp"
49#include "../utils/boundary.hpp"
50#include "../utils/types_conversion.hpp"
52#include "ResetForce.hpp"
53#include "lb_kernels.hpp"
54#if defined(__CUDACC__)
55#include "lb_kernels.cuh"
56#endif
57
63
64#include <utils/Vector.hpp>
67
68#include <array>
69#include <cmath>
70#include <cstddef>
71#include <functional>
72#include <initializer_list>
73#include <limits>
74#include <memory>
75#include <optional>
76#include <stdexcept>
77#include <string>
78#include <tuple>
79#include <type_traits>
80#include <utility>
81#include <variant>
82#include <vector>
83
84namespace walberla {
85
86/** @brief Class that runs and controls the LB on waLBerla. */
87template <typename FloatType, lbmpy::Arch Architecture>
89protected:
91 typename detail::KernelTrait<FloatType,
92 Architecture>::CollisionModelLeesEdwards;
94 typename detail::KernelTrait<FloatType,
95 Architecture>::CollisionModelThermalized;
102 typename detail::BoundaryHandlingTrait<
103 FloatType, Architecture>::Dynamic_UBB>;
105 std::variant<CollisionModelThermalized, CollisionModelLeesEdwards>;
106
107public:
108 /** @brief Stencil for collision and streaming operations. */
109 using Stencil = stencil::D3Q19;
110 /** @brief Stencil for ghost communication (includes domain corners). */
111 using StencilFull = stencil::D3Q27;
112 /** @brief Lattice model (e.g. blockforest). */
114
115protected:
116 template <typename FT, lbmpy::Arch AT = lbmpy::Arch::CPU> struct FieldTrait {
117 using PdfField = field::GhostLayerField<FT, Stencil::Size>;
118 using VectorField = field::GhostLayerField<FT, uint_t{3u}>;
119 template <class Field>
120 using PackInfo = field::communication::PackInfo<Field>;
121 template <class Stencil>
123 blockforest::communication::UniformBufferedScheme<Stencil>;
124 template <class Stencil>
126 blockforest::communication::UniformBufferedScheme<Stencil>;
127 };
128
129#if defined(__CUDACC__)
130 template <typename FT> struct FieldTrait<FT, lbmpy::Arch::GPU> {
131 using PdfField = gpu::GPUField<FT>;
132 using VectorField = gpu::GPUField<FT>;
133 template <class Field>
134 using PackInfo = gpu::communication::MemcpyPackInfo<Field>;
135 template <class Stencil>
136 using RegularCommScheme = gpu::communication::UniformGPUScheme<Stencil>;
137 template <class Stencil>
138 using BoundaryCommScheme =
139 blockforest::communication::UniformBufferedScheme<Stencil>;
140 };
141#endif
142
143 // "underlying" field types (`GPUField` has no f-size info at compile time)
146
147public:
151#if defined(__CUDACC__)
152 using GPUField = gpu::GPUField<FloatType>;
153#endif
154
155public:
156 template <typename T> FloatType FloatType_c(T t) const {
157 return numeric_cast<FloatType>(t);
158 }
159
160 [[nodiscard]] std::size_t stencil_size() const noexcept override {
161 return static_cast<std::size_t>(Stencil::Size);
162 }
163
164 [[nodiscard]] virtual bool is_double_precision() const noexcept override {
165 return std::is_same_v<FloatType, double>;
166 }
167
168private:
169 class CollideSweepVisitor {
170 public:
171 void operator()(CollisionModelThermalized &cm, IBlock *b) { cm(b); }
172
173 void operator()(CollisionModelLeesEdwards &cm, IBlock *b) {
174 cm.v_s_ = static_cast<decltype(cm.v_s_)>(
175 m_lees_edwards_callbacks->get_shear_velocity());
176 cm(b);
177 }
178
179 CollideSweepVisitor() = default;
180 explicit CollideSweepVisitor(std::shared_ptr<LeesEdwardsPack> callbacks) {
181 m_lees_edwards_callbacks = std::move(callbacks);
182 }
183
184 private:
185 std::shared_ptr<LeesEdwardsPack> m_lees_edwards_callbacks{};
186 };
187 CollideSweepVisitor m_run_collide_sweep{};
188
189 FloatType shear_mode_relaxation_rate() const {
190 return FloatType{2} / (FloatType{6} * m_viscosity + FloatType{1});
191 }
192
193 FloatType odd_mode_relaxation_rate(
194 FloatType shear_relaxation,
195 FloatType magic_number = FloatType{3} / FloatType{16}) const {
196 return (FloatType{4} - FloatType{2} * shear_relaxation) /
197 (FloatType{4} * magic_number * shear_relaxation + FloatType{2} -
198 shear_relaxation);
199 }
200
201 void reset_boundary_handling() {
202 auto const &blocks = get_lattice().get_blocks();
203 m_boundary = std::make_shared<BoundaryModel>(blocks, m_pdf_field_id,
205 }
206
207 FloatType pressure_tensor_correction_factor() const {
208 return m_viscosity / (m_viscosity + FloatType{1} / FloatType{6});
209 }
210
211 void pressure_tensor_correction(Matrix3<FloatType> &tensor) const {
212 auto const revert_factor = pressure_tensor_correction_factor();
213 for (auto const i : {1u, 2u, 3u, 5u, 6u, 7u}) {
214 tensor[i] *= revert_factor;
215 }
216 }
217
218 class interpolation_illegal_access : public std::runtime_error {
219 public:
220 explicit interpolation_illegal_access(std::string const &field,
221 Utils::Vector3d const &pos,
222 std::array<int, 3> const &node,
223 double weight)
224 : std::runtime_error("Access to LB " + field + " field failed") {
225 std::cerr << "pos [" << pos << "], node [" << Utils::Vector3i(node)
226 << "], weight " << weight << "\n";
227 }
228 };
229
230 class vtk_runtime_error : public std::runtime_error {
231 public:
232 explicit vtk_runtime_error(std::string const &vtk_uid,
233 std::string const &reason)
234 : std::runtime_error("VTKOutput object '" + vtk_uid + "' " + reason) {}
235 };
236
237protected:
238 // Member variables
239 FloatType m_viscosity; /// kinematic viscosity
240 FloatType m_density;
241 FloatType m_kT;
242
243 // Block data access handles
244 BlockDataID m_pdf_field_id;
246 BlockDataID m_flag_field_id;
247
250
253
254 /** Flag for boundary cells. */
255 FlagUID const Boundary_flag{"boundary"};
256
257 /**
258 * @brief Full communicator.
259 * We use the D3Q27 directions to update cells along the diagonals during
260 * a full ghost communication. This is needed to properly update the corners
261 * of the ghost layer when setting cell velocities or populations.
262 */
264 typename FieldTrait<FloatType, Architecture>::template RegularCommScheme<
265 typename stencil::D3Q27>;
267 typename FieldTrait<FloatType, Architecture>::template BoundaryCommScheme<
268 typename stencil::D3Q27>;
269 /**
270 * @brief Regular communicator.
271 * We use the same directions as the stencil during integration.
272 */
274 typename FieldTrait<FloatType,
275 Architecture>::template RegularCommScheme<Stencil>;
276 template <class Field>
277 using PackInfo =
279
280 // communicators
281 std::shared_ptr<BoundaryFullCommunicator> m_boundary_communicator;
282 std::shared_ptr<RegularFullCommunicator> m_pdf_full_communicator;
283 std::shared_ptr<PDFStreamingCommunicator> m_pdf_streaming_communicator;
284
285 // ResetForce sweep + external force handling
286 std::shared_ptr<ResetForce<PdfField, VectorField>> m_reset_force;
287
288 // Stream sweep
289 std::shared_ptr<StreamSweep> m_stream;
290
291 // Lees Edwards boundary interpolation
292 std::shared_ptr<LeesEdwardsPack> m_lees_edwards_callbacks;
293 std::shared_ptr<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>
295 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
297 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
299
300 // Collision sweep
301 std::shared_ptr<CollisionModel> m_collision_model;
302
303 // boundaries
304 std::shared_ptr<BoundaryModel> m_boundary;
305
306 // lattice
307 std::shared_ptr<LatticeWalberla> m_lattice;
308
309 [[nodiscard]] std::optional<CellInterval>
310 get_interval(Utils::Vector3i const &lower_corner,
311 Utils::Vector3i const &upper_corner) const {
312 auto const &lattice = get_lattice();
313 auto const &cell_min = lower_corner;
314 auto const cell_max = upper_corner - Utils::Vector3i::broadcast(1);
315 auto const lower_bc = get_block_and_cell(lattice, cell_min, true);
316 auto const upper_bc = get_block_and_cell(lattice, cell_max, true);
317 if (not lower_bc or not upper_bc) {
318 return std::nullopt;
319 }
320 assert(&(*(lower_bc->block)) == &(*(upper_bc->block)));
321 return {CellInterval(lower_bc->cell, upper_bc->cell)};
322 }
323
324 /**
325 * @brief Convenience function to add a field with a custom allocator.
326 *
327 * When vectorization is off, let waLBerla decide which memory allocator
328 * to use. When vectorization is on, the aligned memory allocator is
329 * required, otherwise <tt>cpu_vectorize_info["assume_aligned"]</tt> will
330 * trigger assertions. That is because for single-precision kernels the
331 * waLBerla heuristic in <tt>src/field/allocation/FieldAllocator.h</tt>
332 * will fall back to @c StdFieldAlloc, yet @c AllocateAligned is needed
333 * for intrinsics to work.
334 */
335 template <typename Field> auto add_to_storage(std::string const tag) {
336 auto const &blocks = m_lattice->get_blocks();
337 auto const n_ghost_layers = m_lattice->get_ghost_layers();
338 if constexpr (Architecture == lbmpy::Arch::CPU) {
339#ifdef ESPRESSO_BUILD_WITH_AVX_KERNELS
340#if defined(__AVX512F__)
341 constexpr uint_t alignment = 64;
342#elif defined(__AVX__)
343 constexpr uint_t alignment = 32;
344#elif defined(__SSE__)
345 constexpr uint_t alignment = 16;
346#else
347#error "Unsupported arch, check walberla src/field/allocation/FieldAllocator.h"
348#endif
349 using value_type = typename Field::value_type;
350 using Allocator = field::AllocateAligned<value_type, alignment>;
351 auto const allocator = std::make_shared<Allocator>();
352 auto const empty_set = Set<SUID>::emptySet();
353 return field::addToStorage<Field>(
354 blocks, tag, field::internal::defaultSize, FloatType{0}, field::fzyx,
355 n_ghost_layers, false, {}, empty_set, empty_set, allocator);
356#else // ESPRESSO_BUILD_WITH_AVX_KERNELS
357 return field::addToStorage<Field>(blocks, tag, FloatType{0}, field::fzyx,
358 n_ghost_layers);
359#endif // ESPRESSO_BUILD_WITH_AVX_KERNELS
360 }
361#if defined(__CUDACC__)
362 else {
363 auto field_id = gpu::addGPUFieldToStorage<GPUField>(
364 blocks, tag, Field::F_SIZE, field::fzyx, n_ghost_layers);
365 if constexpr (std::is_same_v<Field, _VectorField>) {
366 for (auto block = blocks->begin(); block != blocks->end(); ++block) {
367 auto field = block->template getData<GPUField>(field_id);
368 lbm::accessor::Vector::initialize(field, Vector3<FloatType>{0});
369 }
370 } else if constexpr (std::is_same_v<Field, _PdfField>) {
371 for (auto block = blocks->begin(); block != blocks->end(); ++block) {
372 auto field = block->template getData<GPUField>(field_id);
374 field, std::array<FloatType, Stencil::Size>{});
375 }
376 }
377 return field_id;
378 }
379#endif
380 }
381
382public:
383 LBWalberlaImpl(std::shared_ptr<LatticeWalberla> lattice, double viscosity,
384 double density)
386 m_kT(FloatType{0}), m_lattice(std::move(lattice)) {
387
388 auto const &blocks = m_lattice->get_blocks();
389 auto const n_ghost_layers = m_lattice->get_ghost_layers();
390 if (n_ghost_layers == 0u)
391 throw std::runtime_error("At least one ghost layer must be used");
392
393 // Initialize and register fields (must use the "underlying" types)
394 m_pdf_field_id = add_to_storage<_PdfField>("pdfs");
395 m_pdf_tmp_field_id = add_to_storage<_PdfField>("pdfs_tmp");
396 m_last_applied_force_field_id = add_to_storage<_VectorField>("force last");
397 m_force_to_be_applied_id = add_to_storage<_VectorField>("force next");
398 m_velocity_field_id = add_to_storage<_VectorField>("velocity");
399 m_vec_tmp_field_id = add_to_storage<_VectorField>("velocity_tmp");
400
401 // Initialize and register pdf field
402 auto pdf_setter =
405 for (auto b = blocks->begin(); b != blocks->end(); ++b) {
406 pdf_setter(&*b);
407 }
408
409 // Initialize and register flag field (fluid/boundary)
410 m_flag_field_id = field::addFlagFieldToStorage<FlagField>(
411 blocks, "flag field", n_ghost_layers);
412 // Initialize boundary sweep
413 reset_boundary_handling();
414
415 // Set up the communication and register fields
417 std::make_shared<PDFStreamingCommunicator>(blocks);
418 m_pdf_streaming_communicator->addPackInfo(
419 std::make_shared<PackInfo<PdfField>>(m_pdf_field_id));
420 m_pdf_streaming_communicator->addPackInfo(
422
423 m_pdf_full_communicator = std::make_shared<RegularFullCommunicator>(blocks);
424 m_pdf_full_communicator->addPackInfo(
425 std::make_shared<PackInfo<PdfField>>(m_pdf_field_id));
426 m_pdf_full_communicator->addPackInfo(
428 m_pdf_full_communicator->addPackInfo(
429 std::make_shared<PackInfo<VectorField>>(m_velocity_field_id));
430
432 std::make_shared<BoundaryFullCommunicator>(blocks);
433 m_boundary_communicator->addPackInfo(
434 std::make_shared<field::communication::PackInfo<FlagField>>(
436 auto boundary_packinfo = std::make_shared<
439 boundary_packinfo->setup_boundary_handle(m_lattice, m_boundary);
440 m_boundary_communicator->addPackInfo(boundary_packinfo);
441
442 // Instantiate the sweep responsible for force double buffering and
443 // external forces
444 m_reset_force = std::make_shared<ResetForce<PdfField, VectorField>>(
446
447 // Prepare LB sweeps
448 // Note: For now, combined collide-stream sweeps cannot be used,
449 // because the collide-push variant is not supported by lbmpy.
450 // The following functors are individual in-place collide and stream steps
451 m_stream = std::make_shared<StreamSweep>(
453 }
454
455private:
456 void integrate_stream(std::shared_ptr<Lattice_T> const &blocks) {
457 for (auto b = blocks->begin(); b != blocks->end(); ++b)
458 (*m_stream)(&*b);
459 }
460
461 void integrate_collide(std::shared_ptr<Lattice_T> const &blocks) {
462 auto &cm_variant = *m_collision_model;
463 for (auto b = blocks->begin(); b != blocks->end(); ++b)
464 std::visit(m_run_collide_sweep, cm_variant, std::variant<IBlock *>(&*b));
465 if (auto *cm = std::get_if<CollisionModelThermalized>(&cm_variant)) {
466 cm->time_step_++;
467 }
468 }
469
470 auto has_lees_edwards_bc() const {
471 return std::holds_alternative<CollisionModelLeesEdwards>(
473 }
474
475 void apply_lees_edwards_pdf_interpolation(
476 std::shared_ptr<Lattice_T> const &blocks) {
477 for (auto b = blocks->begin(); b != blocks->end(); ++b)
479 }
480
481 void apply_lees_edwards_vel_interpolation_and_shift(
482 std::shared_ptr<Lattice_T> const &blocks) {
483 for (auto b = blocks->begin(); b != blocks->end(); ++b)
485 }
486
487 void apply_lees_edwards_last_applied_force_interpolation(
488 std::shared_ptr<Lattice_T> const &blocks) {
489 for (auto b = blocks->begin(); b != blocks->end(); ++b)
491 }
492
493 void integrate_reset_force(std::shared_ptr<Lattice_T> const &blocks) {
494 for (auto b = blocks->begin(); b != blocks->end(); ++b)
495 (*m_reset_force)(&*b);
496 }
497
498 void integrate_boundaries(std::shared_ptr<Lattice_T> const &blocks) {
499 for (auto b = blocks->begin(); b != blocks->end(); ++b)
500 (*m_boundary)(&*b);
501 }
502
503 void integrate_push_scheme() {
504 auto const &blocks = get_lattice().get_blocks();
505 // Reset force fields
506 integrate_reset_force(blocks);
507 // LB collide
508 integrate_collide(blocks);
509 m_pdf_streaming_communicator->communicate();
510 // Handle boundaries
511 integrate_boundaries(blocks);
512 // LB stream
513 integrate_stream(blocks);
514 // Refresh ghost layers
516 }
517
518 void integrate_pull_scheme() {
519 auto const &blocks = get_lattice().get_blocks();
520 integrate_reset_force(blocks);
521 // Handle boundaries
522 integrate_boundaries(blocks);
523 // LB stream
524 integrate_stream(blocks);
525 // LB collide
526 integrate_collide(blocks);
527 // Refresh ghost layers
529 }
530
531protected:
532 void integrate_vtk_writers() override {
533 for (auto const &it : m_vtk_auto) {
534 auto &vtk_handle = it.second;
535 if (vtk_handle->enabled) {
536 vtk::writeFiles(vtk_handle->ptr)();
537 vtk_handle->execution_count++;
538 }
539 }
540 }
541
542public:
543 void integrate() override {
544 if (has_lees_edwards_bc()) {
545 integrate_pull_scheme();
546 } else {
547 integrate_push_scheme();
548 }
549 // Handle VTK writers
551 }
552
557
559 m_boundary_communicator->communicate();
560 }
561
563 m_pdf_full_communicator->communicate();
564 if (has_lees_edwards_bc()) {
565 auto const &blocks = get_lattice().get_blocks();
566 apply_lees_edwards_pdf_interpolation(blocks);
567 apply_lees_edwards_vel_interpolation_and_shift(blocks);
568 apply_lees_edwards_last_applied_force_interpolation(blocks);
569 }
570 }
571
572 void set_collision_model(double kT, unsigned int seed) override {
573 auto const omega = shear_mode_relaxation_rate();
574 auto const omega_odd = odd_mode_relaxation_rate(omega);
575 m_kT = FloatType_c(kT);
576 auto obj = CollisionModelThermalized(
578 uint32_t{0u}, uint32_t{0u}, m_kT, omega, omega, omega_odd, omega, seed,
579 uint32_t{0u});
580 obj.block_offset_generator =
581 [this](IBlock *const block, uint32_t &block_offset_0,
582 uint32_t &block_offset_1, uint32_t &block_offset_2) {
583 auto const &blocks = get_lattice().get_blocks();
584 auto const &ci = blocks->getBlockCellBB(*block);
585 block_offset_0 = static_cast<uint32_t>(ci.xMin());
586 block_offset_1 = static_cast<uint32_t>(ci.yMin());
587 block_offset_2 = static_cast<uint32_t>(ci.zMin());
588 };
589 m_collision_model = std::make_shared<CollisionModel>(std::move(obj));
590 }
591
593 std::unique_ptr<LeesEdwardsPack> &&lees_edwards_pack) override {
594 assert(m_kT == 0.);
595#if defined(__CUDACC__)
596 if constexpr (Architecture == lbmpy::Arch::GPU) {
597 throw std::runtime_error("Lees-Edwards LB doesn't support GPU yet");
598 }
599#endif
600 auto const shear_direction = lees_edwards_pack->shear_direction;
601 auto const shear_plane_normal = lees_edwards_pack->shear_plane_normal;
602 auto const shear_vel = FloatType_c(lees_edwards_pack->get_shear_velocity());
603 auto const omega = shear_mode_relaxation_rate();
604 if (shear_plane_normal != 1u) {
605 throw std::domain_error(
606 "Lees-Edwards LB only supports shear_plane_normal=\"y\"");
607 }
608 auto const &lattice = get_lattice();
609 auto const n_ghost_layers = lattice.get_ghost_layers();
610 auto const blocks = lattice.get_blocks();
611 auto const agrid =
612 FloatType_c(lattice.get_grid_dimensions()[shear_plane_normal]);
613 auto obj = CollisionModelLeesEdwards(
614 m_last_applied_force_field_id, m_pdf_field_id, agrid, omega, shear_vel);
615 m_collision_model = std::make_shared<CollisionModel>(std::move(obj));
616 m_lees_edwards_callbacks = std::move(lees_edwards_pack);
617 m_run_collide_sweep = CollideSweepVisitor{m_lees_edwards_callbacks};
619 std::make_shared<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>(
620 blocks, m_pdf_field_id, m_pdf_tmp_field_id, n_ghost_layers,
621 shear_direction, shear_plane_normal,
622 m_lees_edwards_callbacks->get_pos_offset);
623 m_lees_edwards_vel_interpol_sweep = std::make_shared<
625 blocks, m_velocity_field_id, m_vec_tmp_field_id, n_ghost_layers,
626 shear_direction, shear_plane_normal,
627 m_lees_edwards_callbacks->get_pos_offset,
628 m_lees_edwards_callbacks->get_shear_velocity);
632 n_ghost_layers, shear_direction, shear_plane_normal,
633 m_lees_edwards_callbacks->get_pos_offset);
634 }
635
636 void check_lebc(unsigned int shear_direction,
637 unsigned int shear_plane_normal) const override {
639 if (m_lees_edwards_callbacks->shear_direction != shear_direction or
640 m_lees_edwards_callbacks->shear_plane_normal != shear_plane_normal) {
641 throw std::runtime_error(
642 "MD and LB Lees-Edwards boundary conditions disagree");
643 }
644 }
645 }
646
647 void set_viscosity(double viscosity) override {
648 m_viscosity = FloatType_c(viscosity);
649 }
650
651 [[nodiscard]] double get_viscosity() const noexcept override {
652 return numeric_cast<double>(m_viscosity);
653 }
654
655 [[nodiscard]] double get_density() const noexcept override {
656 return numeric_cast<double>(m_density);
657 }
658
659 // Velocity
660 std::optional<Utils::Vector3d>
662 bool consider_ghosts = false) const override {
663 auto const is_boundary = get_node_is_boundary(node, consider_ghosts);
664 if (is_boundary) // is info available locally
665 if (*is_boundary) // is the node a boundary
666 return get_node_velocity_at_boundary(node, consider_ghosts);
667 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
668 if (!bc)
669 return std::nullopt;
670
671 auto field = bc->block->template uncheckedFastGetData<VectorField>(
673 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
674 return to_vector3d(vec);
675 }
676
678 Utils::Vector3d const &v) override {
679 auto bc = get_block_and_cell(get_lattice(), node, false);
680 if (!bc)
681 return false;
682
683 // We have to set both, the pdf and the stored velocity field
684 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
685 auto vel_field =
686 bc->block->template getData<VectorField>(m_velocity_field_id);
687 auto force_field =
688 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
689 auto const vel = to_vector3<FloatType>(v);
690 lbm::accessor::Velocity::set(pdf_field, force_field, vel, bc->cell);
691 lbm::accessor::Vector::set(vel_field, vel, bc->cell);
692
693 return true;
694 }
695
696 std::vector<double>
698 Utils::Vector3i const &upper_corner) const override {
699 std::vector<double> out;
700 if (auto const ci = get_interval(lower_corner, upper_corner)) {
701 auto const &lattice = get_lattice();
702 auto const &block = *(lattice.get_blocks()->begin());
703 auto const field =
704 block.template getData<VectorField>(m_velocity_field_id);
705 auto const local_offset = std::get<0>(lattice.get_local_grid_range());
706 auto const lower_cell = ci->min();
707 auto const upper_cell = ci->max();
708 out.reserve(ci->numCells());
709 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
710 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
711 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
712 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
713 auto const node = local_offset + Utils::Vector3i{{x, y, z}};
714 if (m_boundary->node_is_boundary(node)) {
715 auto const vec = m_boundary->get_node_value_at_boundary(node);
716 for (uint_t f = 0u; f < 3u; ++f) {
717 out.emplace_back(double_c(vec[f]));
718 }
719 } else {
720 auto const vec = lbm::accessor::Vector::get(field, Cell{x, y, z});
721 for (uint_t f = 0u; f < 3u; ++f) {
722 out.emplace_back(double_c(vec[f]));
723 }
724 }
725 }
726 }
727 }
728 }
729 return out;
730 }
731
732 void set_slice_velocity(Utils::Vector3i const &lower_corner,
733 Utils::Vector3i const &upper_corner,
734 std::vector<double> const &velocity) override {
735 if (auto const ci = get_interval(lower_corner, upper_corner)) {
736 auto const &lattice = get_lattice();
737 auto &block = *(lattice.get_blocks()->begin());
738 // We have to set both, the pdf and the stored velocity field
739 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
740 auto vel_field = block.template getData<VectorField>(m_velocity_field_id);
741 auto force_field =
742 block.template getData<VectorField>(m_last_applied_force_field_id);
743 auto const lower_cell = ci->min();
744 auto const upper_cell = ci->max();
745 auto it = velocity.begin();
746 assert(velocity.size() == 3u * ci->numCells());
747 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
748 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
749 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
750 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
751 auto const cell = Cell{x, y, z};
752 Vector3<FloatType> vec;
753 for (uint_t f = 0u; f < 3u; ++f) {
754 vec[f] = FloatType_c(*it);
755 ++it;
756 }
757 lbm::accessor::Velocity::set(pdf_field, force_field, vec, cell);
758 lbm::accessor::Vector::set(vel_field, vec, cell);
759 }
760 }
761 }
762 }
763 }
764
765 [[nodiscard]] bool is_gpu() const noexcept override {
766 return Architecture == lbmpy::Arch::GPU;
767 }
768
769 void add_forces_at_pos(std::vector<Utils::Vector3d> const &pos,
770 std::vector<Utils::Vector3d> const &forces) override {
771 assert(pos.size() == forces.size());
772 if (pos.empty()) {
773 return;
774 }
775 if constexpr (Architecture == lbmpy::Arch::CPU) {
776 for (std::size_t i = 0ul; i < pos.size(); ++i) {
777 add_force_at_pos(pos[i], forces[i]);
778 }
779 }
780#if defined(__CUDACC__)
781 if constexpr (Architecture == lbmpy::Arch::GPU) {
782 auto const &lattice = get_lattice();
783 auto const &block = *(lattice.get_blocks()->begin());
784 auto const origin = block.getAABB().min();
785 std::vector<FloatType> host_pos;
786 std::vector<FloatType> host_force;
787 host_pos.reserve(3ul * pos.size());
788 host_force.reserve(3ul * forces.size());
789 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
790 for (auto const &vec : pos) {
791#pragma unroll
792 for (std::size_t i : {0ul, 1ul, 2ul}) {
793 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
794 }
795 }
796 for (auto const &vec : forces) {
797#pragma unroll
798 for (std::size_t i : {0ul, 1ul, 2ul}) {
799 host_force.emplace_back(static_cast<FloatType>(vec[i]));
800 }
801 }
802 auto const gl = lattice.get_ghost_layers();
803 auto field = block.template uncheckedFastGetData<VectorField>(
805 lbm::accessor::Interpolation::set(field, host_pos, host_force, gl);
806 }
807#endif
808 }
809
810 std::vector<Utils::Vector3d>
811 get_velocities_at_pos(std::vector<Utils::Vector3d> const &pos) override {
812 if (pos.empty()) {
813 return {};
814 }
815 if constexpr (Architecture == lbmpy::Arch::CPU) {
816 std::vector<Utils::Vector3d> vel{};
817 vel.reserve(pos.size());
818 for (auto const &vec : pos) {
819 auto res = get_velocity_at_pos(vec, true);
820 assert(res.has_value());
821 vel.emplace_back(*res);
822 }
823 return vel;
824 }
825#if defined(__CUDACC__)
826 if constexpr (Architecture == lbmpy::Arch::GPU) {
827 auto const &lattice = get_lattice();
828 auto const &block = *(lattice.get_blocks()->begin());
829 auto const origin = block.getAABB().min();
830 std::vector<FloatType> host_pos;
831 host_pos.reserve(3ul * pos.size());
832 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
833 for (auto const &vec : pos) {
834#pragma unroll
835 for (std::size_t i : {0ul, 1ul, 2ul}) {
836 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
837 }
838 }
839 auto const gl = lattice.get_ghost_layers();
840 auto field =
841 block.template uncheckedFastGetData<VectorField>(m_velocity_field_id);
842 auto const res = lbm::accessor::Interpolation::get(field, host_pos, gl);
843 std::vector<Utils::Vector3d> vel{};
844 vel.reserve(res.size() / 3ul);
845 for (auto it = res.begin(); it != res.end(); it += 3) {
846 vel.emplace_back(Utils::Vector3d{static_cast<double>(*(it + 0)),
847 static_cast<double>(*(it + 1)),
848 static_cast<double>(*(it + 2))});
849 }
850 return vel;
851 }
852#endif
853 return {};
854 }
855
856 std::optional<Utils::Vector3d>
858 bool consider_points_in_halo = false) const override {
859 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
860 return std::nullopt;
861 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
862 return std::nullopt;
863 Utils::Vector3d v{0., 0., 0.};
865 pos, [this, &v, &pos](std::array<int, 3> const node, double weight) {
866 // Nodes with zero weight might not be accessible, because they can be
867 // outside ghost layers
868 if (weight != 0.) {
869 auto const res = get_node_velocity(Utils::Vector3i(node), true);
870 if (!res) {
871 throw interpolation_illegal_access("velocity", pos, node, weight);
872 }
873 v += *res * weight;
874 }
875 });
876 return {std::move(v)};
877 }
878
879 std::optional<double>
881 bool consider_points_in_halo = false) const override {
882 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
883 return std::nullopt;
884 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
885 return std::nullopt;
886 double dens = 0.;
888 pos, [this, &dens, &pos](std::array<int, 3> const node, double weight) {
889 // Nodes with zero weight might not be accessible, because they can be
890 // outside ghost layers
891 if (weight != 0.) {
892 auto const res = get_node_density(Utils::Vector3i(node), true);
893 if (!res) {
894 throw interpolation_illegal_access("density", pos, node, weight);
895 }
896 dens += *res * weight;
897 }
898 });
899 return {std::move(dens)};
900 }
901
902 // Local force
904 Utils::Vector3d const &force) override {
905 if (!m_lattice->pos_in_local_halo(pos))
906 return false;
907 auto const force_at_node = [this, &force](std::array<int, 3> const node,
908 double weight) {
909 auto const bc =
911 if (bc) {
912 auto const weighted_force = to_vector3<FloatType>(weight * force);
913 auto force_field =
914 bc->block->template uncheckedFastGetData<VectorField>(
916 lbm::accessor::Vector::add(force_field, weighted_force, bc->cell);
917 }
918 };
919 interpolate_bspline_at_pos(pos, force_at_node);
920 return true;
921 }
922
923 std::optional<Utils::Vector3d>
925 auto const bc = get_block_and_cell(get_lattice(), node, true);
926 if (!bc)
927 return std::nullopt;
928
929 auto field =
930 bc->block->template getData<VectorField>(m_force_to_be_applied_id);
931 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
932 return to_vector3d(vec);
933 }
934
935 std::optional<Utils::Vector3d>
937 bool consider_ghosts = false) const override {
938 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
939 if (!bc)
940 return std::nullopt;
941
942 auto const field =
943 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
944 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
945 return to_vector3d(vec);
946 }
947
949 Utils::Vector3d const &force) override {
950 auto bc = get_block_and_cell(get_lattice(), node, false);
951 if (!bc)
952 return false;
953
954 auto field =
955 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
956 auto const vec = to_vector3<FloatType>(force);
957 lbm::accessor::Vector::set(field, vec, bc->cell);
958
959 return true;
960 }
961
962 std::vector<double> get_slice_last_applied_force(
963 Utils::Vector3i const &lower_corner,
964 Utils::Vector3i const &upper_corner) const override {
965 std::vector<double> out;
966 if (auto const ci = get_interval(lower_corner, upper_corner)) {
967 auto const &lattice = get_lattice();
968 auto const &block = *(lattice.get_blocks()->begin());
969 auto const field =
970 block.template getData<VectorField>(m_last_applied_force_field_id);
971 auto const lower_cell = ci->min();
972 auto const upper_cell = ci->max();
973 auto const n_values = 3u * ci->numCells();
974 out.reserve(n_values);
975 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
976 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
977 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
978 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
979 auto const vec = lbm::accessor::Vector::get(field, Cell{x, y, z});
980 for (uint_t f = 0u; f < 3u; ++f) {
981 out.emplace_back(double_c(vec[f]));
982 }
983 }
984 }
985 }
986 assert(out.size() == n_values);
987 }
988 return out;
989 }
990
992 Utils::Vector3i const &upper_corner,
993 std::vector<double> const &force) override {
994 if (auto const ci = get_interval(lower_corner, upper_corner)) {
995 auto const &lattice = get_lattice();
996 auto &block = *(lattice.get_blocks()->begin());
997 auto field =
998 block.template getData<VectorField>(m_last_applied_force_field_id);
999 auto const lower_cell = ci->min();
1000 auto const upper_cell = ci->max();
1001 auto it = force.begin();
1002 assert(force.size() == 3u * ci->numCells());
1003 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1004 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
1005 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
1006 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
1007 Vector3<FloatType> vec;
1008 for (uint_t f = 0u; f < 3u; ++f) {
1009 vec[f] = FloatType_c(*it);
1010 ++it;
1011 }
1012 lbm::accessor::Vector::set(field, vec, Cell{x, y, z});
1013 }
1014 }
1015 }
1016 }
1017 }
1018
1019 // Population
1020 std::optional<std::vector<double>>
1022 bool consider_ghosts = false) const override {
1023 auto bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1024 if (!bc)
1025 return std::nullopt;
1026
1027 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1028 auto const pop = lbm::accessor::Population::get(pdf_field, bc->cell);
1029 std::vector<double> population(Stencil::Size);
1030 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1031 population[f] = double_c(pop[f]);
1032 }
1033
1034 return {std::move(population)};
1035 }
1036
1038 std::vector<double> const &population) override {
1039 auto bc = get_block_and_cell(get_lattice(), node, false);
1040 if (!bc)
1041 return false;
1042
1043 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1044 std::array<FloatType, Stencil::Size> pop;
1045 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1046 pop[f] = FloatType_c(population[f]);
1047 }
1048 lbm::accessor::Population::set(pdf_field, pop, bc->cell);
1049
1050 return true;
1051 }
1052
1053 std::vector<double>
1055 Utils::Vector3i const &upper_corner) const override {
1056 std::vector<double> out;
1057 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1058 auto const &lattice = get_lattice();
1059 auto const &block = *(lattice.get_blocks()->begin());
1060 auto const pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1061 auto const values = lbm::accessor::Population::get(pdf_field, *ci);
1062 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1063 if constexpr (std::is_same_v<typename decltype(values)::value_type,
1064 double>) {
1065 out = std::move(values);
1066 } else {
1067 out = std::vector<double>(values.begin(), values.end());
1068 }
1069 assert(out.size() == stencil_size() * ci->numCells());
1070 }
1071 return out;
1072 }
1073
1074 void set_slice_population(Utils::Vector3i const &lower_corner,
1075 Utils::Vector3i const &upper_corner,
1076 std::vector<double> const &population) override {
1077 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1078 auto const &lattice = get_lattice();
1079 auto &block = *(lattice.get_blocks()->begin());
1080 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1081 assert(population.size() == stencil_size() * ci->numCells());
1082 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1083 std::vector<FloatType> const values(population.begin(), population.end());
1084 lbm::accessor::Population::set(pdf_field, values, *ci);
1085 }
1086 }
1087
1088 // Density
1089 std::optional<double>
1091 bool consider_ghosts = false) const override {
1092 auto bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1093 if (!bc)
1094 return std::nullopt;
1095
1096 auto pdf_field =
1097 bc->block->template uncheckedFastGetData<PdfField>(m_pdf_field_id);
1098 auto const density = lbm::accessor::Density::get(pdf_field, bc->cell);
1099 return {double_c(density)};
1100 }
1101
1102 bool set_node_density(Utils::Vector3i const &node, double density) override {
1103 auto bc = get_block_and_cell(get_lattice(), node, false);
1104 if (!bc)
1105 return false;
1106
1107 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1108 lbm::accessor::Density::set(pdf_field, FloatType_c(density), bc->cell);
1109
1110 return true;
1111 }
1112
1113 std::vector<double>
1115 Utils::Vector3i const &upper_corner) const override {
1116 std::vector<double> out;
1117 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1118 auto const &lattice = get_lattice();
1119 auto const &block = *(lattice.get_blocks()->begin());
1120 auto const pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1121 auto const values = lbm::accessor::Density::get(pdf_field, *ci);
1122 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1123 if constexpr (std::is_same_v<typename decltype(values)::value_type,
1124 double>) {
1125 out = std::move(values);
1126 } else {
1127 out = std::vector<double>(values.begin(), values.end());
1128 }
1129 assert(out.size() == ci->numCells());
1130 }
1131 return out;
1132 }
1133
1134 void set_slice_density(Utils::Vector3i const &lower_corner,
1135 Utils::Vector3i const &upper_corner,
1136 std::vector<double> const &density) override {
1137 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1138 auto const &lattice = get_lattice();
1139 auto &block = *(lattice.get_blocks()->begin());
1140 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1141 assert(density.size() == ci->numCells());
1142 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1143 std::vector<FloatType> const values(density.begin(), density.end());
1144 lbm::accessor::Density::set(pdf_field, values, *ci);
1145 }
1146 }
1147
1148 std::optional<Utils::Vector3d>
1150 bool consider_ghosts = false) const override {
1151 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1152 if (!bc or !m_boundary->node_is_boundary(node))
1153 return std::nullopt;
1154
1155 return {to_vector3d(m_boundary->get_node_value_at_boundary(node))};
1156 }
1157
1159 Utils::Vector3d const &velocity) override {
1160 auto bc = get_block_and_cell(get_lattice(), node, true);
1161 if (bc) {
1162 m_boundary->set_node_value_at_boundary(
1163 node, to_vector3<FloatType>(velocity), *bc);
1164 }
1165 return bc.has_value();
1166 }
1167
1168 std::vector<std::optional<Utils::Vector3d>> get_slice_velocity_at_boundary(
1169 Utils::Vector3i const &lower_corner,
1170 Utils::Vector3i const &upper_corner) const override {
1171 std::vector<std::optional<Utils::Vector3d>> out;
1172 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1173 auto const &lattice = get_lattice();
1174 auto const local_offset = std::get<0>(lattice.get_local_grid_range());
1175 auto const lower_cell = ci->min();
1176 auto const upper_cell = ci->max();
1177 auto const n_values = ci->numCells();
1178 out.reserve(n_values);
1179 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
1180 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
1181 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
1182 auto const node = local_offset + Utils::Vector3i{{x, y, z}};
1183 if (m_boundary->node_is_boundary(node)) {
1184 out.emplace_back(
1185 to_vector3d(m_boundary->get_node_value_at_boundary(node)));
1186 } else {
1187 out.emplace_back(std::nullopt);
1188 }
1189 }
1190 }
1191 }
1192 assert(out.size() == n_values);
1193 }
1194 return out;
1195 }
1196
1198 Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner,
1199 std::vector<std::optional<Utils::Vector3d>> const &velocity) override {
1200 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1201 auto const &lattice = get_lattice();
1202 auto const local_offset = std::get<0>(lattice.get_local_grid_range());
1203 auto const lower_cell = ci->min();
1204 auto const upper_cell = ci->max();
1205 auto it = velocity.begin();
1206 assert(velocity.size() == ci->numCells());
1207 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
1208 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
1209 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
1210 auto const node = local_offset + Utils::Vector3i{{x, y, z}};
1211 auto const bc = get_block_and_cell(lattice, node, false);
1212 auto const &opt = *it;
1213 if (opt) {
1214 m_boundary->set_node_value_at_boundary(
1215 node, to_vector3<FloatType>(*opt), *bc);
1216 } else {
1217 m_boundary->remove_node_from_boundary(node, *bc);
1218 }
1219 ++it;
1220 }
1221 }
1222 }
1223 }
1224 }
1225
1226 std::optional<Utils::Vector3d>
1228 auto const bc = get_block_and_cell(get_lattice(), node, true);
1229 if (!bc or !m_boundary->node_is_boundary(node))
1230 return std::nullopt;
1231
1232 return get_node_last_applied_force(node, true);
1233 }
1234
1236 auto bc = get_block_and_cell(get_lattice(), node, true);
1237 if (bc) {
1238 m_boundary->remove_node_from_boundary(node, *bc);
1239 }
1240 return bc.has_value();
1241 }
1242
1243 std::optional<bool>
1245 bool consider_ghosts = false) const override {
1246 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1247 if (!bc)
1248 return std::nullopt;
1249
1250 return {m_boundary->node_is_boundary(node)};
1251 }
1252
1253 std::vector<bool>
1255 Utils::Vector3i const &upper_corner) const override {
1256 std::vector<bool> out;
1257 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1258 auto const &lattice = get_lattice();
1259 auto const local_offset = std::get<0>(lattice.get_local_grid_range());
1260 auto const lower_cell = ci->min();
1261 auto const upper_cell = ci->max();
1262 auto const n_values = ci->numCells();
1263 out.reserve(n_values);
1264 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
1265 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
1266 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
1267 auto const node = local_offset + Utils::Vector3i{x, y, z};
1268 out.emplace_back(m_boundary->node_is_boundary(node));
1269 }
1270 }
1271 }
1272 assert(out.size() == n_values);
1273 }
1274 return out;
1275 }
1276
1277 void reallocate_ubb_field() override { m_boundary->boundary_update(); }
1278
1279 void clear_boundaries() override {
1280 reset_boundary_handling();
1282 }
1283
1284 void
1285 update_boundary_from_shape(std::vector<int> const &raster_flat,
1286 std::vector<double> const &data_flat) override {
1287 auto const grid_size = get_lattice().get_grid_dimensions();
1288 auto const data = fill_3D_vector_array(data_flat, grid_size);
1289 set_boundary_from_grid(*m_boundary, get_lattice(), raster_flat, data);
1292 }
1293
1294 // Pressure tensor
1295 std::optional<Utils::VectorXd<9>>
1297 auto bc = get_block_and_cell(get_lattice(), node, false);
1298 if (!bc)
1299 return std::nullopt;
1300
1301 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1302 auto tensor = lbm::accessor::PressureTensor::get(pdf_field, bc->cell);
1303 pressure_tensor_correction(tensor);
1304
1305 return to_vector9d(tensor);
1306 }
1307
1308 std::vector<double> get_slice_pressure_tensor(
1309 Utils::Vector3i const &lower_corner,
1310 Utils::Vector3i const &upper_corner) const override {
1311 std::vector<double> out;
1312 if (auto const ci = get_interval(lower_corner, upper_corner)) {
1313 auto const &lattice = get_lattice();
1314 auto const &block = *(lattice.get_blocks()->begin());
1315 auto const pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1316 auto const lower_cell = ci->min();
1317 auto const upper_cell = ci->max();
1318 auto const n_values = 9u * ci->numCells();
1319 out.reserve(n_values);
1320 assert(++(lattice.get_blocks()->begin()) == lattice.get_blocks()->end());
1321 for (auto x = lower_cell.x(); x <= upper_cell.x(); ++x) {
1322 for (auto y = lower_cell.y(); y <= upper_cell.y(); ++y) {
1323 for (auto z = lower_cell.z(); z <= upper_cell.z(); ++z) {
1324 auto const cell = Cell{x, y, z};
1325 auto tensor = lbm::accessor::PressureTensor::get(pdf_field, cell);
1326 pressure_tensor_correction(tensor);
1327 for (auto i = 0u; i < 9u; ++i) {
1328 out.emplace_back(tensor[i]);
1329 }
1330 }
1331 }
1332 }
1333 assert(out.size() == n_values);
1334 }
1335 return out;
1336 }
1337
1338 // Global pressure tensor
1339 [[nodiscard]] Utils::VectorXd<9> get_pressure_tensor() const override {
1340 auto const &blocks = get_lattice().get_blocks();
1341 Matrix3<FloatType> tensor(FloatType{0});
1342 for (auto block = blocks->begin(); block != blocks->end(); ++block) {
1343 auto pdf_field = block->template getData<PdfField>(m_pdf_field_id);
1344 WALBERLA_FOR_ALL_CELLS_XYZ(pdf_field, {
1345 tensor += lbm::accessor::PressureTensor::get(pdf_field, Cell{x, y, z});
1346 });
1347 }
1348 auto const grid_size = get_lattice().get_grid_dimensions();
1349 auto const number_of_nodes = Utils::product(grid_size);
1350 pressure_tensor_correction(tensor);
1351 return to_vector9d(tensor) * (1. / static_cast<double>(number_of_nodes));
1352 }
1353
1354 // Global momentum
1355 [[nodiscard]] Utils::Vector3d get_momentum() const override {
1356 auto const &blocks = get_lattice().get_blocks();
1357 Vector3<FloatType> mom(FloatType{0});
1358 for (auto block = blocks->begin(); block != blocks->end(); ++block) {
1359 auto pdf_field = block->template getData<PdfField>(m_pdf_field_id);
1360 auto force_field =
1361 block->template getData<VectorField>(m_last_applied_force_field_id);
1362 mom += lbm::accessor::MomentumDensity::reduce(pdf_field, force_field);
1363 }
1364 return to_vector3d(mom);
1365 }
1366
1367 // Global external force
1368 void set_external_force(Utils::Vector3d const &ext_force) override {
1369 m_reset_force->set_ext_force(ext_force);
1370 }
1371
1372 [[nodiscard]] Utils::Vector3d get_external_force() const noexcept override {
1373 return m_reset_force->get_ext_force();
1374 }
1375
1376 [[nodiscard]] double get_kT() const noexcept override {
1377 return numeric_cast<double>(m_kT);
1378 }
1379
1380 [[nodiscard]] std::optional<uint64_t> get_rng_state() const override {
1381 auto const cm = std::get_if<CollisionModelThermalized>(&*m_collision_model);
1382 if (!cm or m_kT == 0.) {
1383 return std::nullopt;
1384 }
1385 return {static_cast<uint64_t>(cm->time_step_)};
1386 }
1387
1388 void set_rng_state(uint64_t counter) override {
1389 auto const cm = std::get_if<CollisionModelThermalized>(&*m_collision_model);
1390 if (!cm or m_kT == 0.) {
1391 throw std::runtime_error("This LB instance is unthermalized");
1392 }
1393 assert(counter <=
1394 static_cast<uint32_t>(std::numeric_limits<uint_t>::max()));
1395 cm->time_step_ = static_cast<uint32_t>(counter);
1396 }
1397
1398 [[nodiscard]] LatticeWalberla const &get_lattice() const noexcept override {
1399 return *m_lattice;
1400 }
1401
1402 [[nodiscard]] std::size_t get_velocity_field_id() const noexcept override {
1403 return m_velocity_field_id;
1404 }
1405
1406 [[nodiscard]] std::size_t get_force_field_id() const noexcept override {
1408 }
1409
1410 void register_vtk_field_filters(walberla::vtk::VTKOutput &vtk_obj) override {
1411 if constexpr (Architecture == lbmpy::Arch::GPU) {
1412 throw std::runtime_error("VTK output not supported for GPU");
1413 } else {
1414 field::FlagFieldCellFilter<FlagField> fluid_filter(m_flag_field_id);
1415 fluid_filter.addFlag(Boundary_flag);
1416 vtk_obj.addCellExclusionFilter(fluid_filter);
1417 }
1418 }
1419
1420protected:
1421 template <typename Field_T, uint_t F_SIZE_ARG, typename OutputType>
1422 class VTKWriter : public vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG> {
1423 public:
1424 VTKWriter(ConstBlockDataID const &block_id, std::string const &id,
1425 FloatType unit_conversion)
1426 : vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG>(id),
1427 m_block_id(block_id), m_field(nullptr),
1428 m_conversion(unit_conversion) {}
1429
1430 protected:
1431 void configure() override {
1432 WALBERLA_ASSERT_NOT_NULLPTR(this->block_);
1433 m_field = this->block_->template getData<Field_T>(m_block_id);
1434 }
1435
1436 ConstBlockDataID const m_block_id;
1437 Field_T const *m_field;
1438 FloatType const m_conversion;
1439 };
1440
1441 template <typename OutputType = float>
1442 class DensityVTKWriter : public VTKWriter<PdfField, 1u, OutputType> {
1443 public:
1444 using VTKWriter<PdfField, 1u, OutputType>::VTKWriter;
1445
1446 protected:
1447 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1448 cell_idx_t const z, cell_idx_t const) override {
1449 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1450 auto const density =
1451 lbm::accessor::Density::get(this->m_field, {x, y, z});
1452 return numeric_cast<OutputType>(this->m_conversion * density);
1453 }
1454 };
1455
1456 template <typename OutputType = float>
1457 class VelocityVTKWriter : public VTKWriter<VectorField, 3u, OutputType> {
1458 public:
1459 using VTKWriter<VectorField, 3u, OutputType>::VTKWriter;
1460
1461 protected:
1462 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1463 cell_idx_t const z, cell_idx_t const f) override {
1464 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1465 auto const velocity =
1466 lbm::accessor::Vector::get(this->m_field, {x, y, z});
1467 return numeric_cast<OutputType>(this->m_conversion * velocity[uint_c(f)]);
1468 }
1469 };
1470
1471 template <typename OutputType = float>
1472 class PressureTensorVTKWriter : public VTKWriter<PdfField, 9u, OutputType> {
1473 public:
1474 PressureTensorVTKWriter(ConstBlockDataID const &block_id,
1475 std::string const &id, FloatType unit_conversion,
1476 FloatType off_diag_factor)
1477 : VTKWriter<PdfField, 9u, OutputType>::VTKWriter(block_id, id,
1478 unit_conversion),
1479 m_off_diag_factor(off_diag_factor) {}
1480
1481 protected:
1482 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1483 cell_idx_t const z, cell_idx_t const f) override {
1484 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1485 auto const pressure =
1487 auto const revert_factor =
1488 (f == 0 or f == 4 or f == 8) ? FloatType{1} : m_off_diag_factor;
1489 return numeric_cast<OutputType>(this->m_conversion * revert_factor *
1490 pressure[uint_c(f)]);
1491 }
1492 FloatType const m_off_diag_factor;
1493 };
1494
1495public:
1496 void register_vtk_field_writers(walberla::vtk::VTKOutput &vtk_obj,
1497 LatticeModel::units_map const &units,
1498 int flag_observables) override {
1499 if constexpr (Architecture == lbmpy::Arch::GPU) {
1500 throw std::runtime_error("VTK output not supported for GPU");
1501 } else {
1502 if (flag_observables & static_cast<int>(OutputVTK::density)) {
1503 auto const unit_conversion = FloatType_c(units.at("density"));
1504 vtk_obj.addCellDataWriter(make_shared<DensityVTKWriter<float>>(
1505 m_pdf_field_id, "density", unit_conversion));
1506 }
1507 if (flag_observables & static_cast<int>(OutputVTK::velocity_vector)) {
1508 auto const unit_conversion = FloatType_c(units.at("velocity"));
1509 vtk_obj.addCellDataWriter(make_shared<VelocityVTKWriter<float>>(
1510 m_velocity_field_id, "velocity_vector", unit_conversion));
1511 }
1512 if (flag_observables & static_cast<int>(OutputVTK::pressure_tensor)) {
1513 auto const unit_conversion = FloatType_c(units.at("pressure"));
1514 vtk_obj.addCellDataWriter(make_shared<PressureTensorVTKWriter<float>>(
1515 m_pdf_field_id, "pressure_tensor", unit_conversion,
1516 pressure_tensor_correction_factor()));
1517 }
1518 }
1519 }
1520
1521 ~LBWalberlaImpl() override = default;
1522};
1523
1524} // namespace walberla
LBWalberlaBase provides the public interface of the LB waLBerla bridge.
Vector implementation and trait types for boost qvm interoperability.
float f[3]
__global__ float * force
__shared__ int pos[MAXDEPTH *THREADS5/WARPSIZE]
__shared__ int node[MAXDEPTH *THREADS5/WARPSIZE]
__shared__ float res[]
float u[3]
Definition Cell.hpp:98
Interface of a lattice-based fluid model.
std::map< std::string, std::shared_ptr< VTKHandle > > m_vtk_auto
VTK writers that are executed automatically.
std::unordered_map< std::string, double > units_map
Class that runs and controls the BlockForest in waLBerla.
walberla::blockforest::StructuredBlockForest Lattice_T
static Vector< T, N > broadcast(T const &s)
Create a vector that has all entries set to one value.
Definition Vector.hpp:110
Boundary class optimized for sparse data.
field::FlagField< uint8_t > FlagField
OutputType evaluate(cell_idx_t const x, cell_idx_t const y, cell_idx_t const z, cell_idx_t const) override
OutputType evaluate(cell_idx_t const x, cell_idx_t const y, cell_idx_t const z, cell_idx_t const f) override
PressureTensorVTKWriter(ConstBlockDataID const &block_id, std::string const &id, FloatType unit_conversion, FloatType off_diag_factor)
VTKWriter(ConstBlockDataID const &block_id, std::string const &id, FloatType unit_conversion)
OutputType evaluate(cell_idx_t const x, cell_idx_t const y, cell_idx_t const z, cell_idx_t const f) override
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
typename FieldTrait< FloatType, Architecture >::template RegularCommScheme< typename stencil::D3Q27 > RegularFullCommunicator
Full communicator.
std::vector< double > get_slice_last_applied_force(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
LatticeWalberla::Lattice_T Lattice_T
Lattice model (e.g.
std::shared_ptr< RegularFullCommunicator > m_pdf_full_communicator
std::vector< std::optional< Utils::Vector3d > > get_slice_velocity_at_boundary(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
std::optional< Utils::Vector3d > get_node_last_applied_force(Utils::Vector3i const &node, bool consider_ghosts=false) const override
stencil::D3Q19 Stencil
Stencil for collision and streaming operations.
std::optional< Utils::Vector3d > get_node_velocity_at_boundary(Utils::Vector3i const &node, bool consider_ghosts=false) const override
void ghost_communication() override
std::optional< Utils::Vector3d > get_node_velocity(Utils::Vector3i const &node, bool consider_ghosts=false) const override
void reallocate_ubb_field() override
typename detail::KernelTrait< FloatType, Architecture >::CollisionModelLeesEdwards CollisionModelLeesEdwards
std::size_t get_force_field_id() const noexcept override
std::shared_ptr< CollisionModel > m_collision_model
typename FieldTrait< FloatType, Architecture >::VectorField VectorField
void set_slice_velocity(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &velocity) override
void integrate_vtk_writers() override
bool remove_node_from_boundary(Utils::Vector3i const &node) override
std::optional< Utils::Vector3d > get_node_boundary_force(Utils::Vector3i const &node) const override
std::optional< double > get_density_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
void set_slice_velocity_at_boundary(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< std::optional< Utils::Vector3d > > const &velocity) override
std::vector< double > get_slice_population(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
void set_rng_state(uint64_t counter) override
std::vector< double > get_slice_velocity(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
std::shared_ptr< LatticeWalberla > m_lattice
std::shared_ptr< LeesEdwardsPack > m_lees_edwards_callbacks
LBWalberlaImpl(std::shared_ptr< LatticeWalberla > lattice, double viscosity, double density)
bool set_node_last_applied_force(Utils::Vector3i const &node, Utils::Vector3d const &force) override
std::shared_ptr< InterpolateAndShiftAtBoundary< _VectorField, FloatType > > m_lees_edwards_vel_interpol_sweep
std::vector< bool > get_slice_is_boundary(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
std::vector< double > get_slice_density(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
void update_boundary_from_shape(std::vector< int > const &raster_flat, std::vector< double > const &data_flat) override
std::shared_ptr< BoundaryModel > m_boundary
std::vector< double > get_slice_pressure_tensor(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
bool set_node_density(Utils::Vector3i const &node, double density) override
void set_collision_model(double kT, unsigned int seed) override
std::variant< CollisionModelThermalized, CollisionModelLeesEdwards > CollisionModel
std::vector< Utils::Vector3d > get_velocities_at_pos(std::vector< Utils::Vector3d > const &pos) override
std::optional< bool > get_node_is_boundary(Utils::Vector3i const &node, bool consider_ghosts=false) const override
std::shared_ptr< ResetForce< PdfField, VectorField > > m_reset_force
virtual bool is_double_precision() const noexcept override
std::shared_ptr< StreamSweep > m_stream
FlagUID const Boundary_flag
Flag for boundary cells.
~LBWalberlaImpl() override=default
std::optional< uint64_t > get_rng_state() const override
std::optional< std::vector< double > > get_node_population(Utils::Vector3i const &node, bool consider_ghosts=false) const override
FloatType m_density
kinematic viscosity
double get_viscosity() const noexcept override
double get_kT() const noexcept override
Utils::Vector3d get_momentum() const override
void set_viscosity(double viscosity) override
std::shared_ptr< BoundaryFullCommunicator > m_boundary_communicator
std::size_t stencil_size() const noexcept override
auto add_to_storage(std::string const tag)
Convenience function to add a field with a custom allocator.
typename detail::KernelTrait< FloatType, Architecture >::InitialPDFsSetter InitialPDFsSetter
bool set_node_velocity_at_boundary(Utils::Vector3i const &node, Utils::Vector3d const &velocity) override
stencil::D3Q27 StencilFull
Stencil for ghost communication (includes domain corners).
void register_vtk_field_writers(walberla::vtk::VTKOutput &vtk_obj, LatticeModel::units_map const &units, int flag_observables) override
typename BoundaryModel::FlagField FlagField
std::optional< double > get_node_density(Utils::Vector3i const &node, bool consider_ghosts=false) const override
void check_lebc(unsigned int shear_direction, unsigned int shear_plane_normal) const override
void set_slice_density(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &density) override
typename FieldTrait< FloatType >::PdfField _PdfField
Utils::Vector3d get_external_force() const noexcept override
FloatType FloatType_c(T t) const
std::size_t get_velocity_field_id() const noexcept override
Utils::VectorXd< 9 > get_pressure_tensor() const override
std::shared_ptr< InterpolateAndShiftAtBoundary< _PdfField, FloatType > > m_lees_edwards_pdf_interpol_sweep
void set_external_force(Utils::Vector3d const &ext_force) override
std::optional< Utils::Vector3d > get_velocity_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
void set_slice_last_applied_force(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &force) override
void set_collision_model(std::unique_ptr< LeesEdwardsPack > &&lees_edwards_pack) override
typename FieldTrait< FloatType, Architecture >::template BoundaryCommScheme< typename stencil::D3Q27 > BoundaryFullCommunicator
typename FieldTrait< FloatType, Architecture >::template RegularCommScheme< Stencil > PDFStreamingCommunicator
Regular communicator.
LatticeWalberla const & get_lattice() const noexcept override
void set_slice_population(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &population) override
typename FieldTrait< FloatType, Architecture >::template PackInfo< Field > PackInfo
double get_density() const noexcept override
typename detail::KernelTrait< FloatType, Architecture >::CollisionModelThermalized CollisionModelThermalized
typename FieldTrait< FloatType, Architecture >::PdfField PdfField
void register_vtk_field_filters(walberla::vtk::VTKOutput &vtk_obj) override
std::optional< Utils::Vector3d > get_node_force_to_be_applied(Utils::Vector3i const &node) const override
typename FieldTrait< FloatType >::VectorField _VectorField
bool set_node_velocity(Utils::Vector3i const &node, Utils::Vector3d const &v) override
std::optional< CellInterval > get_interval(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const
bool set_node_population(Utils::Vector3i const &node, std::vector< double > const &population) override
bool is_gpu() const noexcept override
std::shared_ptr< PDFStreamingCommunicator > m_pdf_streaming_communicator
BlockDataID m_last_applied_force_field_id
std::optional< Utils::VectorXd< 9 > > get_node_pressure_tensor(Utils::Vector3i const &node) const override
std::shared_ptr< InterpolateAndShiftAtBoundary< _VectorField, FloatType > > m_lees_edwards_last_applied_force_interpol_sweep
typename detail::KernelTrait< FloatType, Architecture >::StreamSweep StreamSweep
bool add_force_at_pos(Utils::Vector3d const &pos, Utils::Vector3d const &force) override
void setup_boundary_handle(std::shared_ptr< LatticeWalberla > lattice, std::shared_ptr< Boundary_T > boundary)
static double weight(int type, double r_cut, double k, double r)
Definition dpd.cpp:79
static double * block(double *p, std::size_t index, std::size_t size)
Definition elc.cpp:174
T product(Vector< T, N > const &v)
Definition Vector.hpp:359
VectorXi< 3 > Vector3i
Definition Vector.hpp:166
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, double const rho_in, Cell const &cell)
double get(GhostLayerField< double, uint_t{19u}> const *pdf_field, Cell const &cell)
std::vector< double > get(gpu::GPUField< double > const *vec_field, std::vector< double > const &pos, uint gl)
void set(gpu::GPUField< double > const *vec_field, std::vector< double > const &pos, std::vector< double > const &forces, uint gl)
Vector3< double > reduce(GhostLayerField< double, uint_t{19u}> const *pdf_field, GhostLayerField< double, uint_t{3u}> const *force_field)
std::array< double, 19u > get(GhostLayerField< double, uint_t{19u}> const *pdf_field, Cell const &cell)
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, std::array< double, 19u > const &pop, Cell const &cell)
void initialize(GhostLayerField< double, uint_t{19u}> *pdf_field, std::array< double, 19u > const &pop)
Matrix3< double > get(GhostLayerField< double, uint_t{19u}> const *pdf_field, Cell const &cell)
void initialize(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec)
void add(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec, Cell const &cell)
void set(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec, Cell const &cell)
Vector3< double > get(GhostLayerField< double, uint_t{3u}> const *vec_field, Cell const &cell)
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, GhostLayerField< double, uint_t{3u}> const *force_field, Vector3< double > const &u, Cell const &cell)
std::optional< BlockAndCell > get_block_and_cell(::LatticeWalberla const &lattice, Utils::Vector3i const &node, bool consider_ghost_layers)
void interpolate_bspline_at_pos(Utils::Vector3d const &pos, Function const &f)
Utils::VectorXd< 9 > to_vector9d(Matrix3< double > const &m)
void set_boundary_from_grid(BoundaryModel &boundary, LatticeWalberla const &lattice, std::vector< int > const &raster_flat, std::vector< DataType > const &data_flat)
Definition boundary.hpp:81
std::vector< Utils::Vector3d > fill_3D_vector_array(std::vector< double > const &vec_flat, Utils::Vector3i const &grid_size)
Definition boundary.hpp:36
Utils::Vector3d to_vector3d(Vector3< float > const &v)
static Utils::Vector3d velocity(Particle const &p_ref, Particle const &p_vs)
Velocity of the virtual site.
Definition relative.cpp:64
DEVICE_QUALIFIER constexpr iterator begin() noexcept
Definition Array.hpp:128
DEVICE_QUALIFIER constexpr size_type size() const noexcept
Definition Array.hpp:154
field::GhostLayerField< FT, uint_t{3u}> VectorField
field::communication::PackInfo< Field > PackInfo
field::GhostLayerField< FT, Stencil::Size > PdfField
blockforest::communication::UniformBufferedScheme< Stencil > RegularCommScheme
blockforest::communication::UniformBufferedScheme< Stencil > BoundaryCommScheme