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-2024 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/HostFieldAllocator.h>
44#include <gpu/communication/MemcpyPackInfo.h>
45#include <gpu/communication/UniformGPUScheme.h>
46#endif
47
48#include "../BoundaryHandling.hpp"
49#include "../BoundaryPackInfo.hpp"
50#include "../utils/boundary.hpp"
51#include "../utils/types_conversion.hpp"
53#include "ResetForce.hpp"
54#include "lb_kernels.hpp"
55#if defined(__CUDACC__)
56#include "lb_kernels.cuh"
57#endif
58
66
67#include <utils/Vector.hpp>
68#include <utils/index.hpp>
71
72#include <array>
73#include <bitset>
74#include <cmath>
75#include <cstddef>
76#include <functional>
77#include <initializer_list>
78#include <limits>
79#include <memory>
80#include <optional>
81#include <stdexcept>
82#include <string>
83#include <tuple>
84#include <type_traits>
85#include <utility>
86#include <variant>
87#include <vector>
88
89namespace walberla {
90
91/** @brief Class that runs and controls the LB on waLBerla. */
92template <typename FloatType, lbmpy::Arch Architecture>
94protected:
96 detail::KernelTrait<FloatType,
99 detail::KernelTrait<FloatType,
111
112public:
113 /** @brief Stencil for collision and streaming operations. */
114 using Stencil = stencil::D3Q19;
115 /** @brief Stencil for ghost communication (includes domain corners). */
116 using StencilFull = stencil::D3Q27;
117 /** @brief Lattice model (e.g. blockforest). */
119
120protected:
121 template <typename FT, lbmpy::Arch AT = lbmpy::Arch::CPU> struct FieldTrait {
122 using PdfField = field::GhostLayerField<FT, Stencil::Size>;
123 using VectorField = field::GhostLayerField<FT, uint_t{3u}>;
124 template <class Field>
125 using PackInfo = field::communication::PackInfo<Field>;
128 template <class Stencil>
130 blockforest::communication::UniformBufferedScheme<Stencil>;
131 template <class Stencil>
133 blockforest::communication::UniformBufferedScheme<Stencil>;
134 };
135
136#if defined(__CUDACC__)
137 template <typename FT> struct FieldTrait<FT, lbmpy::Arch::GPU> {
138 private:
139 static auto constexpr AT = lbmpy::Arch::GPU;
140 template <class Field>
141 using MemcpyPackInfo = gpu::communication::MemcpyPackInfo<Field>;
142
143 public:
144 template <typename Stencil>
145 class UniformGPUScheme
146 : public gpu::communication::UniformGPUScheme<Stencil> {
147 public:
148 explicit UniformGPUScheme(auto const &bf)
149 : gpu::communication::UniformGPUScheme<Stencil>(
150 bf, /* sendDirectlyFromGPU */ false,
151 /* useLocalCommunication */ false) {}
152 };
153 using PdfField = gpu::GPUField<FT>;
154 using VectorField = gpu::GPUField<FT>;
155 template <class Field> using PackInfo = MemcpyPackInfo<Field>;
158 template <class Stencil>
159 using RegularCommScheme = UniformGPUScheme<Stencil>;
160 template <class Stencil>
161 using BoundaryCommScheme =
162 blockforest::communication::UniformBufferedScheme<Stencil>;
163 };
164#endif
165
166 // "underlying" field types (`GPUField` has no f-size info at compile time)
169
170public:
174#if defined(__CUDACC__)
175 using GPUField = gpu::GPUField<FloatType>;
178#endif
179
180 struct GhostComm {
181 /** @brief Ghost communication operations. */
182 enum GhostCommFlags : unsigned {
183 PDF, ///< PDFs communication
184 VEL, ///< velocities communication
185 LAF, ///< last applied forces communication
186 UBB, ///< boundaries communication
187 SIZE
188 };
189 };
190
191public:
192 template <typename T> FloatType FloatType_c(T t) const {
193 return numeric_cast<FloatType>(t);
194 }
195
196 [[nodiscard]] std::size_t stencil_size() const noexcept override {
197 return static_cast<std::size_t>(Stencil::Size);
198 }
199
200 [[nodiscard]] bool is_double_precision() const noexcept override {
201 return std::is_same_v<FloatType, double>;
202 }
203
204private:
205 class StreamCollideSweepVisitor {
206 public:
207 using StructuredBlockStorage = LatticeWalberla::Lattice_T;
208
209 void operator()(StreamCollisionModelThermalized &cm, IBlock *b) {
210 cm.configure(m_storage, b);
211 cm(b);
212 }
213
214 void operator()(StreamCollisionModelLeesEdwards &cm, IBlock *b) {
215 cm.setV_s(static_cast<decltype(cm.getV_s())>(
216 m_lees_edwards_callbacks->get_shear_velocity()));
217 cm(b);
218 }
219
220 StreamCollideSweepVisitor() = default;
221 StreamCollideSweepVisitor(std::shared_ptr<StructuredBlockStorage> storage) {
222 m_storage = std::move(storage);
223 }
224 StreamCollideSweepVisitor(std::shared_ptr<StructuredBlockStorage> storage,
225 std::shared_ptr<LeesEdwardsPack> callbacks) {
226 m_storage = std::move(storage);
227 m_lees_edwards_callbacks = std::move(callbacks);
228 }
229
230 private:
231 std::shared_ptr<StructuredBlockStorage> m_storage{};
232 std::shared_ptr<LeesEdwardsPack> m_lees_edwards_callbacks{};
233 };
234 StreamCollideSweepVisitor m_run_stream_collide_sweep{};
235
236 FloatType shear_mode_relaxation_rate() const {
237 return FloatType{2} / (FloatType{6} * m_viscosity + FloatType{1});
238 }
239
240 FloatType odd_mode_relaxation_rate(
241 FloatType shear_relaxation,
242 FloatType magic_number = FloatType{3} / FloatType{16}) const {
243 return (FloatType{4} - FloatType{2} * shear_relaxation) /
244 (FloatType{4} * magic_number * shear_relaxation + FloatType{2} -
245 shear_relaxation);
246 }
247
248 void reset_boundary_handling(std::shared_ptr<BlockStorage> const &blocks) {
249 auto const [lc, uc] = m_lattice->get_local_grid_range(true);
250 m_boundary =
251 std::make_shared<BoundaryModel>(blocks, m_pdf_field_id, m_flag_field_id,
252 CellInterval{to_cell(lc), to_cell(uc)});
253 }
254
255 FloatType pressure_tensor_correction_factor() const {
256 return m_viscosity / (m_viscosity + FloatType{1} / FloatType{6});
257 }
258
259 void pressure_tensor_correction(Matrix3<FloatType> &tensor) const {
260 auto const revert_factor = pressure_tensor_correction_factor();
261 for (auto const i : {1u, 2u, 3u, 5u, 6u, 7u}) {
262 tensor[i] *= revert_factor;
263 }
264 }
265
266 void pressure_tensor_correction(std::span<FloatType, 9ul> tensor) const {
267 auto const revert_factor = pressure_tensor_correction_factor();
268 for (auto const i : {1u, 2u, 3u, 5u, 6u, 7u}) {
269 tensor[i] *= revert_factor;
270 }
271 }
272
273 class interpolation_illegal_access : public std::runtime_error {
274 public:
275 interpolation_illegal_access(std::string const &field,
276 Utils::Vector3d const &pos,
277 std::array<int, 3> const &node, double weight)
278 : std::runtime_error("Access to LB " + field + " field failed") {
279 std::cerr << "pos [" << pos << "], node [" << Utils::Vector3i(node)
280 << "], weight " << weight << "\n";
281 }
282 };
283
284 class vtk_runtime_error : public std::runtime_error {
285 public:
286 vtk_runtime_error(std::string const &vtk_uid, std::string const &reason)
287 : std::runtime_error("VTKOutput object '" + vtk_uid + "' " + reason) {}
288 };
289
290protected:
291 // Member variables
292 FloatType m_viscosity; /// kinematic viscosity
293 FloatType m_density;
294 FloatType m_kT;
295 unsigned int m_seed;
296 double m_zc_to_md; // zero-centered conversion factor to MD units
297 double m_zc_to_lb; // zero-centered conversion factor to LB units
298
299 // Block data access handles
300 BlockDataID m_pdf_field_id;
302 BlockDataID m_flag_field_id;
303
306
309
310#if defined(__CUDACC__)
311 std::optional<BlockDataID> m_pdf_cpu_field_id;
312 std::optional<BlockDataID> m_vel_cpu_field_id;
313#endif
314
315 /** Flag for boundary cells. */
316 FlagUID const Boundary_flag{"boundary"};
317 bool m_has_boundaries{false};
318
319 // lattice
320 std::shared_ptr<LatticeWalberla> m_lattice;
321
322 // boundaries
323 std::shared_ptr<BoundaryModel> m_boundary;
324
325 /**
326 * @brief Full communicator.
327 * We use the D3Q27 directions to update cells along the diagonals during
328 * a full ghost communication. This is needed to properly update the corners
329 * of the ghost layer when setting cell velocities or populations.
330 */
332 FieldTrait<FloatType,
333 Architecture>::template RegularCommScheme<stencil::D3Q27>;
335 FieldTrait<FloatType,
336 Architecture>::template BoundaryCommScheme<stencil::D3Q27>;
337 /**
338 * @brief Regular communicator.
339 * We use the same directions as the stencil during integration.
340 */
342 FieldTrait<FloatType, Architecture>::template RegularCommScheme<Stencil>;
343 template <class Field>
344 using PackInfo =
346
347 // communicators
348 std::shared_ptr<BoundaryFullCommunicator> m_boundary_communicator;
349 std::shared_ptr<RegularFullCommunicator> m_full_communicator;
350 std::shared_ptr<RegularFullCommunicator> m_pdf_communicator;
351 std::shared_ptr<RegularFullCommunicator> m_vel_communicator;
352 std::shared_ptr<RegularFullCommunicator> m_laf_communicator;
353 std::shared_ptr<PDFStreamingCommunicator> m_pdf_streaming_communicator;
354 std::bitset<GhostComm::SIZE> m_pending_ghost_comm;
356
357 // collision sweep
358 std::shared_ptr<CollisionModel> m_collision_model;
359
360 // force reset sweep + external force handling
361 std::shared_ptr<ResetForce<PdfField, VectorField>> m_reset_force;
362
363 // velocity update sweep
364 std::shared_ptr<UpdateVelFromPDF> m_update_velocities_from_pdf;
365
366 // Lees-Edwards boundary interpolation
367 std::shared_ptr<LeesEdwardsPack> m_lees_edwards_callbacks;
368 std::shared_ptr<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>
370 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
372 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
374
375#if defined(__CUDACC__)
376 std::shared_ptr<gpu::HostFieldAllocator<FloatType>> m_host_field_allocator;
377#endif
378
379 /**
380 * @brief Convenience function to add a field with a custom allocator.
381 *
382 * When vectorization is off, let waLBerla decide which memory allocator
383 * to use. When vectorization is on, the aligned memory allocator is
384 * required, otherwise <tt>cpu_vectorize_info["assume_aligned"]</tt> will
385 * trigger assertions. That is because for single-precision kernels the
386 * waLBerla heuristic in <tt>src/field/allocation/FieldAllocator.h</tt>
387 * will fall back to @c StdFieldAlloc, yet @c AllocateAligned is needed
388 * for intrinsics to work.
389 */
390 template <typename Field> auto add_to_storage(std::string const tag) {
391 auto const &blocks = m_lattice->get_blocks();
392 auto const n_ghost_layers = m_lattice->get_ghost_layers();
393 if constexpr (Architecture == lbmpy::Arch::CPU) {
394#ifdef ESPRESSO_BUILD_WITH_AVX_KERNELS
395 constexpr auto alignment = field::SIMDAlignment();
396 using value_type = Field::value_type;
397 using Allocator = field::AllocateAligned<value_type, alignment>;
398 auto const allocator = std::make_shared<Allocator>();
399 auto const empty_set = Set<SUID>::emptySet();
400 return field::addToStorage<Field>(
401 blocks, tag, field::internal::defaultSize, FloatType{0}, field::fzyx,
402 n_ghost_layers, false, {}, empty_set, empty_set, allocator);
403#else // ESPRESSO_BUILD_WITH_AVX_KERNELS
404 return field::addToStorage<Field>(blocks, tag, FloatType{0}, field::fzyx,
405 n_ghost_layers);
406#endif // ESPRESSO_BUILD_WITH_AVX_KERNELS
407 }
408#if defined(__CUDACC__)
409 else {
410 auto field_id = gpu::addGPUFieldToStorage<GPUField>(
411 blocks, tag, Field::F_SIZE, field::fzyx, n_ghost_layers);
412 if constexpr (std::is_same_v<Field, _VectorField>) {
413 for (auto &block : *blocks) {
414 auto field = block.template getData<GPUField>(field_id);
415 lbm::accessor::Vector::initialize(field, Vector3<FloatType>{0});
416 }
417 } else if constexpr (std::is_same_v<Field, _PdfField>) {
418 for (auto &block : *blocks) {
419 auto field = block.template getData<GPUField>(field_id);
421 field, std::array<FloatType, Stencil::Size>{});
422 }
423 }
424 return field_id;
425 }
426#endif
427 }
428
430 auto const setup = [this]<typename PackInfoPdf, typename PackInfoVec>() {
431 auto const &blocks = m_lattice->get_blocks();
433 std::make_shared<PDFStreamingCommunicator>(blocks);
434 m_pdf_streaming_communicator->addPackInfo(
435 std::make_shared<PackInfoPdf>(m_pdf_field_id));
436 m_pdf_streaming_communicator->addPackInfo(
437 std::make_shared<PackInfoVec>(m_last_applied_force_field_id));
438 };
440 using PackInfoPdf = FieldTrait::PackInfoStreamingPdf;
441 using PackInfoVec = FieldTrait::PackInfoStreamingVec;
442 if (m_has_boundaries or (m_collision_model and has_lees_edwards_bc())) {
443 setup.template operator()<PackInfo<PdfField>, PackInfoVec>();
444 } else {
445 setup.template operator()<PackInfoPdf, PackInfoVec>();
446 }
447 }
448
449public:
450 LBWalberlaImpl(std::shared_ptr<LatticeWalberla> lattice, double viscosity,
451 double density)
453 m_kT(FloatType{0}), m_seed(0u), m_zc_to_md(density),
454 m_zc_to_lb(1. / density), m_lattice(std::move(lattice)),
456
457 auto const &blocks = m_lattice->get_blocks();
458 auto const n_ghost_layers = m_lattice->get_ghost_layers();
459 if (n_ghost_layers == 0u)
460 throw std::runtime_error("At least one ghost layer must be used");
461
462 // Initialize and register fields (must use the "underlying" types)
463 m_pdf_field_id = add_to_storage<_PdfField>("pdfs");
464 m_pdf_tmp_field_id = add_to_storage<_PdfField>("pdfs_tmp");
465 m_last_applied_force_field_id = add_to_storage<_VectorField>("force last");
466 m_force_to_be_applied_id = add_to_storage<_VectorField>("force next");
467 m_velocity_field_id = add_to_storage<_VectorField>("velocity");
468 m_vel_tmp_field_id = add_to_storage<_VectorField>("velocity_tmp");
469#if defined(__CUDACC__)
470 m_host_field_allocator =
471 std::make_shared<gpu::HostFieldAllocator<FloatType>>();
472#endif
473
474 // Initialize and register pdf field with zero centered density
475 auto pdf_setter = InitialPDFsSetter(
477 for (auto &block : *blocks) {
478 pdf_setter(&block);
479 }
480
481 // Initialize and register flag field (fluid/boundary)
482 m_flag_field_id = field::addFlagFieldToStorage<FlagField>(
483 blocks, "flag field", n_ghost_layers);
484 // Initialize boundary sweep
485 reset_boundary_handling(m_lattice->get_blocks());
486
487 // Set up the communication and register fields
489
490 m_full_communicator = std::make_shared<RegularFullCommunicator>(blocks);
491 m_full_communicator->addPackInfo(
492 std::make_shared<PackInfo<PdfField>>(m_pdf_field_id));
493 m_full_communicator->addPackInfo(
495 m_full_communicator->addPackInfo(
496 std::make_shared<PackInfo<VectorField>>(m_velocity_field_id));
497
498 m_pdf_communicator = std::make_shared<RegularFullCommunicator>(blocks);
499 m_vel_communicator = std::make_shared<RegularFullCommunicator>(blocks);
500 m_laf_communicator = std::make_shared<RegularFullCommunicator>(blocks);
501 m_pdf_communicator->addPackInfo(
502 std::make_shared<PackInfo<PdfField>>(m_pdf_field_id));
503 m_vel_communicator->addPackInfo(
504 std::make_shared<PackInfo<VectorField>>(m_velocity_field_id));
505 m_laf_communicator->addPackInfo(
507
509 std::make_shared<BoundaryFullCommunicator>(blocks);
510 m_boundary_communicator->addPackInfo(
513 auto boundary_packinfo = std::make_shared<
516 boundary_packinfo->setup_boundary_handle(m_lattice, m_boundary);
517 m_boundary_communicator->addPackInfo(boundary_packinfo);
518
520
521 // Instantiate the sweep responsible for force double buffering and
522 // external forces
523 m_reset_force = std::make_shared<ResetForce<PdfField, VectorField>>(
525
526 // Instantiate velocity update sweep
527 m_update_velocities_from_pdf = std::make_shared<UpdateVelFromPDF>(
529 }
530
531private:
532 void integrate_stream_collide(std::shared_ptr<BlockStorage> const &blocks) {
533 auto &cm_variant = *m_collision_model;
534 for (auto &block : *blocks) {
535 auto const block_variant = std::variant<IBlock *>(&block);
536 std::visit(m_run_stream_collide_sweep, cm_variant, block_variant);
537 }
538 if (auto *cm = std::get_if<StreamCollisionModelThermalized>(&cm_variant)) {
539 cm->setTime_step(cm->getTime_step() + 1u);
540 }
541 }
542
543 auto has_lees_edwards_bc() const {
544 return std::holds_alternative<StreamCollisionModelLeesEdwards>(
546 }
547
548 void apply_lees_edwards_pdf_interpolation(
549 std::shared_ptr<BlockStorage> const &blocks) {
550 for (auto &block : *blocks)
552 }
553
554 void apply_lees_edwards_vel_interpolation_and_shift(
555 std::shared_ptr<BlockStorage> const &blocks) {
556 for (auto &block : *blocks)
558 }
559
560 void apply_lees_edwards_last_applied_force_interpolation(
561 std::shared_ptr<BlockStorage> const &blocks) {
562 for (auto &block : *blocks)
564 }
565
566 void integrate_reset_force(std::shared_ptr<BlockStorage> const &blocks) {
567 for (auto &block : *blocks)
568 (*m_reset_force)(&block);
569 }
570
571 void integrate_boundaries(std::shared_ptr<BlockStorage> const &blocks) {
572 for (auto &block : *blocks)
573 (*m_boundary)(&block);
574 }
575
576 void integrate_update_velocities_from_pdf(
577 std::shared_ptr<BlockStorage> const &blocks) {
578 for (auto &block : *blocks)
580 }
581
582 void integrate_pull_scheme() {
584 auto const &blocks = get_lattice().get_blocks();
585 // Reset force fields
586 integrate_reset_force(blocks);
587 // LB stream collide
588 integrate_stream_collide(blocks);
589 // Mark pending ghost layer updates
590 // As pdf and laf are communicated direcly afterwards they are not set.
593 m_pdf_streaming_communicator->communicate();
594 if (has_lees_edwards_bc()) {
595 apply_lees_edwards_pdf_interpolation(blocks);
596 apply_lees_edwards_last_applied_force_interpolation(blocks);
597 }
598 // Handle boundaries
599 if (m_has_boundaries) {
600 integrate_boundaries(blocks);
601 }
602 // Update velocities from pdfs
603 integrate_update_velocities_from_pdf(blocks);
604
605 if (has_lees_edwards_bc()) {
606 apply_lees_edwards_vel_interpolation_and_shift(blocks);
607 }
608 }
609
610protected:
611 void integrate_vtk_writers() override {
612 for (auto const &it : m_vtk_auto) {
613 auto &vtk_handle = it.second;
614 if (vtk_handle->enabled) {
615 vtk::writeFiles(vtk_handle->ptr)();
616 vtk_handle->execution_count++;
617 }
618 }
619 }
620
621public:
622 void integrate() override {
623 integrate_pull_scheme();
625 }
626
636
637 void ghost_communication_pdf() override {
640 m_pdf_communicator->communicate();
641 if (has_lees_edwards_bc()) {
642 auto const &blocks = get_lattice().get_blocks();
643 apply_lees_edwards_pdf_interpolation(blocks);
644 }
646 }
647 }
648
649 void ghost_communication_vel() override {
652 m_vel_communicator->communicate();
653 if (has_lees_edwards_bc()) {
654 auto const &blocks = get_lattice().get_blocks();
655 apply_lees_edwards_vel_interpolation_and_shift(blocks);
656 }
658 }
659 }
660
661 void ghost_communication_laf() override {
664 m_laf_communicator->communicate();
665 if (has_lees_edwards_bc()) {
666 auto const &blocks = get_lattice().get_blocks();
667 apply_lees_edwards_last_applied_force_interpolation(blocks);
668 }
670 }
671 }
672
680
683 m_full_communicator->communicate();
684 if (has_lees_edwards_bc()) {
686 }
690 }
691
693 auto const &blocks = get_lattice().get_blocks();
694 apply_lees_edwards_pdf_interpolation(blocks);
695 apply_lees_edwards_vel_interpolation_and_shift(blocks);
696 apply_lees_edwards_last_applied_force_interpolation(blocks);
697 }
698
699 void set_collision_model(double kT, unsigned int seed) override {
700 auto const omega = shear_mode_relaxation_rate();
701 auto const omega_odd = odd_mode_relaxation_rate(omega);
702 auto const blocks = get_lattice().get_blocks();
703 m_kT = FloatType_c(kT);
704 m_seed = seed;
707 zero_centered_to_lb(m_kT), omega, omega, omega_odd, omega, seed,
708 uint32_t{0u});
709 m_collision_model = std::make_shared<CollisionModel>(std::move(obj));
710 m_run_stream_collide_sweep = StreamCollideSweepVisitor(blocks);
712 }
713
715 std::unique_ptr<LeesEdwardsPack> &&lees_edwards_pack) override {
716 assert(m_kT == 0.);
717#if defined(__CUDACC__)
718 if constexpr (Architecture == lbmpy::Arch::GPU) {
719 throw std::runtime_error("Lees-Edwards LB doesn't support GPU yet");
720 }
721#endif
722 auto const shear_direction = lees_edwards_pack->shear_direction;
723 auto const shear_plane_normal = lees_edwards_pack->shear_plane_normal;
724 auto const shear_vel = FloatType_c(lees_edwards_pack->get_shear_velocity());
725 auto const omega = shear_mode_relaxation_rate();
726 auto const omega_odd = odd_mode_relaxation_rate(omega);
727 if (shear_plane_normal != 1u) {
728 throw std::domain_error(
729 "Lees-Edwards LB only supports shear_plane_normal=\"y\"");
730 }
731 auto const &lattice = get_lattice();
732 auto const n_ghost_layers = lattice.get_ghost_layers();
733 auto const blocks = lattice.get_blocks();
734 if (lattice.get_node_grid()[shear_direction] != 1 or
735 blocks->getSize(shear_direction) != 1ul or
736 blocks->getSize(shear_plane_normal) !=
737 lattice.get_node_grid()[shear_plane_normal]) {
738 throw std::domain_error("LB LEbc doesn't support domain decomposition "
739 "along the shear direction, nor multiple blocks "
740 "along the normal direction");
741 }
742 auto const &grid_dimensions = lattice.get_grid_dimensions();
743 auto const block_origin = lattice.get_local_grid_range(false).first;
744 auto const lebc_slab_origin = block_origin[shear_plane_normal];
745 auto const lebc_slab_total_thickness = grid_dimensions[shear_plane_normal];
746 auto const lebc_bot_index = 0 - lebc_slab_origin;
747 auto const lebc_top_index = lebc_slab_total_thickness - lebc_slab_origin;
749 std::make_shared<CollisionModel>(StreamCollisionModelLeesEdwards(
751 lebc_top_index, omega, omega, omega_odd, omega, shear_vel));
752 m_lees_edwards_callbacks = std::move(lees_edwards_pack);
753 m_run_stream_collide_sweep =
754 StreamCollideSweepVisitor(blocks, m_lees_edwards_callbacks);
756 std::make_shared<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>(
757 blocks, m_pdf_field_id, m_pdf_tmp_field_id, n_ghost_layers,
758 shear_direction, shear_plane_normal,
759 m_lees_edwards_callbacks->get_pos_offset);
760 m_lees_edwards_vel_interpol_sweep = std::make_shared<
762 blocks, m_velocity_field_id, m_vel_tmp_field_id, n_ghost_layers,
763 shear_direction, shear_plane_normal,
764 m_lees_edwards_callbacks->get_pos_offset,
765 m_lees_edwards_callbacks->get_shear_velocity);
769 n_ghost_layers, shear_direction, shear_plane_normal,
770 m_lees_edwards_callbacks->get_pos_offset);
772 }
773
774 void check_lebc(unsigned int shear_direction,
775 unsigned int shear_plane_normal) const override {
777 if (m_lees_edwards_callbacks->shear_direction != shear_direction or
778 m_lees_edwards_callbacks->shear_plane_normal != shear_plane_normal) {
779 throw std::runtime_error(
780 "MD and LB Lees-Edwards boundary conditions disagree");
781 }
782 }
783 }
784
785 void set_viscosity(double viscosity) override {
786 m_viscosity = FloatType_c(viscosity);
787 }
788
789 [[nodiscard]] double get_viscosity() const noexcept override {
790 return static_cast<double>(m_viscosity);
791 }
792
793 [[nodiscard]] double get_density() const noexcept override {
794 return static_cast<double>(m_density);
795 }
796
797 template <typename T>
798 void zero_centered_transform_impl(T &data, auto const factor) const {
799 if constexpr (std::is_arithmetic_v<T>) {
800 static_assert(std::is_floating_point_v<T>);
801 data *= static_cast<T>(factor);
802 } else {
803 auto const coef = static_cast<typename T::value_type>(factor);
804 std::transform(std::begin(data), std::end(data), std::begin(data),
805 [coef](auto value) { return value * coef; });
806 }
807 }
808
809 void zero_centered_to_lb_in_place(auto &data) const {
811 }
812
813 void zero_centered_to_md_in_place(auto &data) const {
815 }
816
817 auto zero_centered_to_lb(auto const &data) const {
818 auto transformed_data = data;
819 zero_centered_to_lb_in_place(transformed_data);
820 return transformed_data;
821 }
822
823 auto zero_centered_to_md(auto const &data) const {
824 auto transformed_data = data;
825 zero_centered_to_md_in_place(transformed_data);
826 return transformed_data;
827 }
828
829 // Velocity
830 std::optional<Utils::Vector3d>
832 bool consider_ghosts = false) const override {
833 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::VEL)));
834 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::UBB)));
835 if (m_has_boundaries) {
836 auto const is_boundary = get_node_is_boundary(node, consider_ghosts);
837 if (is_boundary and *is_boundary) {
838 return get_node_velocity_at_boundary(node, consider_ghosts);
839 }
840 }
841 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
842 if (!bc)
843 return std::nullopt;
844
845 auto field = bc->block->template uncheckedFastGetData<VectorField>(
847 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
848 return to_vector3d(vec);
849 }
850
852 Utils::Vector3d const &v) override {
855 auto bc = get_block_and_cell(get_lattice(), node, false);
856 if (!bc)
857 return false;
858
859 // We have to set both, the pdf and the stored velocity field
860 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
861 auto vel_field =
862 bc->block->template getData<VectorField>(m_velocity_field_id);
863 auto force_field =
864 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
865 auto vel = to_vector3<FloatType>(v);
866 lbm::accessor::Velocity::set(pdf_field, vel_field, force_field, vel,
867 bc->cell);
868
869 return true;
870 }
871
872 std::vector<double>
874 Utils::Vector3i const &upper_corner) const override {
875 std::vector<double> out;
876 uint_t values_size = 0;
877 auto const &lattice = get_lattice();
878 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
879 out = std::vector<double>(3u * ci->numCells());
880 for (auto &block : *lattice.get_blocks()) {
881 auto const block_offset = lattice.get_block_corner(block, true);
882 if (auto const bci = get_block_interval(
883 lattice, lower_corner, upper_corner, block_offset, block)) {
884 auto const field =
885 block.template getData<VectorField>(m_velocity_field_id);
886 auto values = lbm::accessor::Vector::get(field, *bci);
887 assert(values.size() == 3u * bci->numCells());
888 values_size += 3u * bci->numCells();
889
890 auto kernel = [&values, &out, this](unsigned const block_index,
891 unsigned const local_index,
892 Utils::Vector3i const &node) {
893 if (m_boundary->node_is_boundary(node)) {
894 auto const &vec = m_boundary->get_node_value_at_boundary(node);
895 for (uint_t f = 0u; f < 3u; ++f) {
896 out[3u * local_index + f] = vec[f];
897 }
898 } else {
899 for (uint_t f = 0u; f < 3u; ++f) {
900 out[3u * local_index + f] =
901 double_c(values[3u * block_index + f]);
902 }
903 }
904 };
905
906 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
907 }
908 }
909 assert(values_size == 3u * ci->numCells());
910 }
911 return out;
912 }
913
914 void set_slice_velocity(Utils::Vector3i const &lower_corner,
915 Utils::Vector3i const &upper_corner,
916 std::vector<double> const &velocity) override {
919 auto const &lattice = get_lattice();
920 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
921 assert(velocity.size() == 3u * ci->numCells());
922 for (auto &block : *lattice.get_blocks()) {
923 auto const block_offset = lattice.get_block_corner(block, true);
924 if (auto const bci = get_block_interval(
925 lattice, lower_corner, upper_corner, block_offset, block)) {
926 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
927 auto force_field = block.template getData<VectorField>(
929 auto vel_field =
930 block.template getData<VectorField>(m_velocity_field_id);
931 std::vector<FloatType> values(3u * bci->numCells());
932
933 auto kernel = [&values, &velocity](unsigned const block_index,
934 unsigned const local_index,
935 Utils::Vector3i const &) {
936 for (uint_t f = 0u; f < 3u; ++f) {
937 values[3u * block_index + f] =
938 numeric_cast<FloatType>(velocity[3u * local_index + f]);
939 }
940 };
941
942 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
943 lbm::accessor::Velocity::set(pdf_field, vel_field, force_field,
944 values, *bci);
945 }
946 }
947 }
948 }
949
950 [[nodiscard]] bool is_gpu() const noexcept override {
951 return Architecture == lbmpy::Arch::GPU;
952 }
953
954 std::function<bool(Utils::Vector3d const &)>
955 make_lattice_position_checker(bool consider_points_in_halo) const override {
956 auto const &lat = *m_lattice;
957 if (consider_points_in_halo) {
958 return [&](Utils::Vector3d const &p) { return lat.pos_in_local_halo(p); };
959 }
960 return [&](Utils::Vector3d const &p) { return lat.pos_in_local_domain(p); };
961 }
962
963 void add_forces_at_pos(std::vector<Utils::Vector3d> const &pos,
964 std::vector<Utils::Vector3d> const &forces) override {
965 assert(pos.size() == forces.size());
966 if (pos.empty()) {
967 return;
968 }
969 if constexpr (Architecture == lbmpy::Arch::CPU) {
970 auto const kernel = make_force_interpolation_kernel();
971 for (std::size_t i = 0ul; i < pos.size(); ++i) {
972 kernel(pos[i], forces[i]);
973 }
974 }
975#if defined(__CUDACC__)
976 if constexpr (Architecture == lbmpy::Arch::GPU) {
977 auto const &lattice = get_lattice();
978 auto const &block = *(lattice.get_blocks()->begin());
979 auto const origin = block.getAABB().min();
980 std::vector<FloatType> host_pos;
981 std::vector<FloatType> host_force;
982 host_pos.reserve(3ul * pos.size());
983 host_force.reserve(3ul * forces.size());
984 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
985 for (auto const &vec : pos) {
986#pragma unroll
987 for (std::size_t i : {0ul, 1ul, 2ul}) {
988 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
989 }
990 }
991 for (auto const &vec : forces) {
992#pragma unroll
993 for (std::size_t i : {0ul, 1ul, 2ul}) {
994 host_force.emplace_back(static_cast<FloatType>(vec[i]));
995 }
996 }
998 auto const gl = lattice.get_ghost_layers();
999 auto field = block.template uncheckedFastGetData<VectorField>(
1001 lbm::accessor::Interpolation::add_force(field, host_pos, host_force, gl);
1002 }
1003#endif
1004 }
1005
1007 auto const &lattice = *m_lattice;
1008 auto const &blocks = *lattice.get_blocks();
1009 assert(lattice.get_ghost_layers() == 1u);
1010 return [&](Utils::Vector3d const &pos, Utils::Vector3d const &force) {
1011 if (not get_block_extended(lattice, pos, 1u)) {
1012 return;
1013 }
1015 pos, [&, conv = m_zc_to_lb, field_id = m_force_to_be_applied_id](
1016 std::array<int, 3> const node, double weight) {
1017 auto block = get_block_extended(lattice, node, 0u);
1018 if (!block)
1019 block = get_block_extended(lattice, node, 1u);
1020 if (block) {
1021 auto cell = to_cell(node);
1022 blocks.transformGlobalToBlockLocalCell(cell, *block);
1023 weight *= conv;
1024 auto const weighted_force = to_vector3<FloatType>(weight * force);
1025 auto field =
1026 block->template uncheckedFastGetData<VectorField>(field_id);
1027 lbm::accessor::Vector::add(field, weighted_force, cell);
1028 }
1029 });
1030 };
1031 }
1032
1034 auto const &lattice = *m_lattice;
1035 auto const &blocks = *lattice.get_blocks();
1036 assert(lattice.get_ghost_layers() == 1u);
1037 return [&](Utils::Vector3d const &pos) {
1038 Utils::Vector3d acc{0., 0., 0.};
1040 std::array<int, 3> const node,
1041 double weight) {
1042 // Nodes with zero weight might not be accessible, because they can be
1043 // outside ghost layers
1044 if (weight != 0.) {
1045 auto block = get_block_extended(lattice, node, 1u);
1046 if (!block)
1047 throw interpolation_illegal_access("velocity", pos, node, weight);
1048 Vector3<FloatType> vel;
1049 if (m_has_boundaries and m_boundary->node_is_boundary(node)) {
1050 vel = m_boundary->get_node_value_at_boundary(node);
1051 } else {
1052 auto cell = to_cell(node);
1053 blocks.transformGlobalToBlockLocalCell(cell, *block);
1054 auto field =
1055 block->template uncheckedFastGetData<VectorField>(field_id);
1056 vel = lbm::accessor::Vector::get(field, cell);
1057 }
1058 acc += to_vector3d(vel) * weight;
1059 }
1060 });
1061 return acc;
1062 };
1063 }
1064
1066 auto const &lattice = *m_lattice;
1067 auto const &blocks = *lattice.get_blocks();
1068 assert(lattice.get_ghost_layers() == 1u);
1069 return [&](Utils::Vector3d const &pos) {
1070 double acc = 0.;
1072 field_id = m_pdf_field_id](
1073 std::array<int, 3> const node,
1074 double weight) {
1075 // Nodes with zero weight might not be accessible, because they can be
1076 // outside ghost layers
1077 if (weight != 0.) {
1078 auto block = get_block_extended(lattice, node, 1u);
1079 if (!block)
1080 throw interpolation_illegal_access("density", pos, node, weight);
1081 auto cell = to_cell(node);
1082 blocks.transformGlobalToBlockLocalCell(cell, *block);
1083 auto field = block->template uncheckedFastGetData<PdfField>(field_id);
1084 auto const rho = lbm::accessor::Density::get(field, density, cell);
1085 acc += rho * weight;
1086 }
1087 });
1088 return acc;
1089 };
1090 }
1091
1092 std::vector<Utils::Vector3d>
1093 get_velocities_at_pos(std::vector<Utils::Vector3d> const &pos) override {
1094 if (pos.empty()) {
1095 return {};
1096 }
1097 std::vector<Utils::Vector3d> vel{};
1098 vel.reserve(pos.size());
1099 if constexpr (Architecture == lbmpy::Arch::CPU) {
1100 auto const kernel = make_velocity_interpolation_kernel();
1101 std::ranges::transform(pos, std::back_inserter(vel), kernel);
1102 }
1103#if defined(__CUDACC__)
1104 if constexpr (Architecture == lbmpy::Arch::GPU) {
1105 auto const &lattice = get_lattice();
1106 auto const &block = *(lattice.get_blocks()->begin());
1107 auto const origin = block.getAABB().min();
1108 std::vector<FloatType> host_pos;
1109 host_pos.reserve(3ul * pos.size());
1110 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
1111 for (auto const &vec : pos) {
1112#pragma unroll
1113 for (std::size_t i : {0ul, 1ul, 2ul}) {
1114 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
1115 }
1116 }
1117 auto const gl = lattice.get_ghost_layers();
1118 auto field =
1119 block.template uncheckedFastGetData<VectorField>(m_velocity_field_id);
1120 // the velocity field has indeterminate values inside boundary regions;
1121 // we overwrite them with boundary slip velocities before interpolation
1122 auto const [dev_idx, dev_vel] = m_boundary->get_flattened_map_device();
1123 if (not dev_idx->empty()) {
1124 lbm::accessor::Vector::set_from_list(field, *dev_idx, *dev_vel, gl);
1125 }
1126 auto const res =
1127 lbm::accessor::Interpolation::get_vel(field, host_pos, gl);
1128 for (auto it = res.begin(); it != res.end(); it += 3) {
1129 vel.emplace_back(Utils::Vector3d{static_cast<double>(*(it + 0)),
1130 static_cast<double>(*(it + 1)),
1131 static_cast<double>(*(it + 2))});
1132 }
1133 }
1134#endif
1135 return vel;
1136 }
1137
1138 std::vector<double>
1139 get_densities_at_pos(std::vector<Utils::Vector3d> const &pos) override {
1140 if (pos.empty()) {
1141 return {};
1142 }
1143 std::vector<double> rho{};
1144 rho.reserve(pos.size());
1145 if constexpr (Architecture == lbmpy::Arch::CPU) {
1146 auto const kernel = make_density_interpolation_kernel();
1147 std::ranges::transform(pos, std::back_inserter(rho), kernel);
1148 }
1149#if defined(__CUDACC__)
1150 if constexpr (Architecture == lbmpy::Arch::GPU) {
1151 auto const &lattice = get_lattice();
1152 auto const &block = *(lattice.get_blocks()->begin());
1153 auto const origin = block.getAABB().min();
1154 std::vector<FloatType> host_pos;
1155 host_pos.reserve(3ul * pos.size());
1156 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
1157 for (auto const &vec : pos) {
1158#pragma unroll
1159 for (std::size_t i : {0ul, 1ul, 2ul}) {
1160 host_pos.emplace_back(static_cast<FloatType>(vec[i] - origin[i]));
1161 }
1162 }
1163 auto const gl = lattice.get_ghost_layers();
1164 auto field =
1165 block.template uncheckedFastGetData<PdfField>(m_pdf_field_id);
1166 auto res =
1168 if constexpr (std::is_same_v<FloatType, double>) {
1169 std::swap(rho, res);
1170 } else {
1171 for (auto const &v : res) {
1172 rho.emplace_back(static_cast<double>(v));
1173 }
1174 }
1175 }
1176#endif
1177 return rho;
1178 }
1179
1180 std::optional<Utils::Vector3d>
1182 bool consider_points_in_halo = false) const override {
1183 assert(not m_pending_ghost_comm.test(GhostComm::VEL));
1184 assert(not m_pending_ghost_comm.test(GhostComm::UBB));
1185 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
1186 return std::nullopt;
1187 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
1188 return std::nullopt;
1189 auto const kernel = make_velocity_interpolation_kernel();
1190 return {kernel(pos)};
1191 }
1192
1193 std::optional<double>
1195 bool consider_points_in_halo = false) const override {
1196 assert(not m_pending_ghost_comm.test(GhostComm::PDF));
1197 if (!consider_points_in_halo and !m_lattice->pos_in_local_domain(pos))
1198 return std::nullopt;
1199 if (consider_points_in_halo and !m_lattice->pos_in_local_halo(pos))
1200 return std::nullopt;
1201 auto const kernel = make_density_interpolation_kernel();
1202 return {kernel(pos)};
1203 }
1204
1206 Utils::Vector3d const &force) override {
1207 if (!m_lattice->pos_in_local_halo(pos))
1208 return false;
1209 auto const kernel = make_force_interpolation_kernel();
1210 kernel(pos, force);
1211 return true;
1212 }
1213
1214 std::optional<Utils::Vector3d>
1215 get_node_force_to_be_applied(Utils::Vector3i const &node) const override {
1216 auto const bc = get_block_and_cell(get_lattice(), node, true);
1217 if (!bc)
1218 return std::nullopt;
1219
1220 auto field =
1221 bc->block->template getData<VectorField>(m_force_to_be_applied_id);
1222 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
1223 return zero_centered_to_md(to_vector3d(vec));
1224 }
1225
1226 std::optional<Utils::Vector3d>
1228 bool consider_ghosts = false) const override {
1229 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::LAF)));
1230 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1231 if (!bc)
1232 return std::nullopt;
1233
1234 auto const field =
1235 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
1236 auto const vec = lbm::accessor::Vector::get(field, bc->cell);
1237 return zero_centered_to_md(to_vector3d(vec));
1238 }
1239
1241 Utils::Vector3d const &force) override {
1244 auto bc = get_block_and_cell(get_lattice(), node, false);
1245 if (!bc)
1246 return false;
1247
1248 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1249 auto force_field =
1250 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
1251 auto vel_field =
1252 bc->block->template getData<VectorField>(m_velocity_field_id);
1253 auto const vec = to_vector3<FloatType>(force);
1254 lbm::accessor::Force::set(pdf_field, vel_field, force_field, vec, m_density,
1255 bc->cell);
1256
1257 return true;
1258 }
1259
1260 std::vector<double> get_slice_last_applied_force(
1261 Utils::Vector3i const &lower_corner,
1262 Utils::Vector3i const &upper_corner) const override {
1263 std::vector<double> out;
1264 auto const &lattice = get_lattice();
1265 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1266 out = std::vector<double>(3u * ci->numCells());
1267 for (auto const &block : *lattice.get_blocks()) {
1268 auto const block_offset = lattice.get_block_corner(block, true);
1269 if (auto const bci = get_block_interval(
1270 lattice, lower_corner, upper_corner, block_offset, block)) {
1271 auto const field = block.template getData<VectorField>(
1273 auto const values = lbm::accessor::Vector::get(field, *bci);
1274 assert(values.size() == 3u * bci->numCells());
1275
1276 auto kernel = [&values, &out](unsigned const block_index,
1277 unsigned const local_index,
1278 Utils::Vector3i const &) {
1279 for (uint_t f = 0u; f < 3u; ++f) {
1280 out[3u * local_index + f] = values[3u * block_index + f];
1281 }
1282 };
1283
1284 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1285 }
1286 }
1287 }
1289 return out;
1290 }
1291
1293 Utils::Vector3i const &upper_corner,
1294 std::vector<double> const &force) override {
1297 auto const &lattice = get_lattice();
1298 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1299 assert(force.size() == 3u * ci->numCells());
1300 for (auto &block : *lattice.get_blocks()) {
1301 auto const block_offset = lattice.get_block_corner(block, true);
1302 if (auto const bci = get_block_interval(
1303 lattice, lower_corner, upper_corner, block_offset, block)) {
1304 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1305 auto force_field = block.template getData<VectorField>(
1307 auto vel_field =
1308 block.template getData<VectorField>(m_velocity_field_id);
1309 std::vector<FloatType> values(3u * bci->numCells());
1310
1311 auto kernel = [&values, &force](unsigned const block_index,
1312 unsigned const local_index,
1313 Utils::Vector3i const &) {
1314 for (uint_t f = 0u; f < 3u; ++f) {
1315 values[3u * block_index + f] =
1316 numeric_cast<FloatType>(force[3u * local_index + f]);
1317 }
1318 };
1319
1320 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1321 lbm::accessor::Force::set(pdf_field, vel_field, force_field, values,
1322 m_density, *bci);
1323 }
1324 }
1325 }
1326 }
1327
1328 // Population
1329 std::optional<std::vector<double>>
1331 bool consider_ghosts = false) const override {
1332 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::PDF)));
1333 auto bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1334 if (!bc)
1335 return std::nullopt;
1336
1337 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1338 auto const pop = lbm::accessor::Population::get(pdf_field, bc->cell);
1339 std::vector<double> population(Stencil::Size);
1340 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1341 population[f] = double_c(pop[f]);
1342 }
1343
1344 return {std::move(population)};
1345 }
1346
1348 std::vector<double> const &population) override {
1351 auto bc = get_block_and_cell(get_lattice(), node, false);
1352 if (!bc)
1353 return false;
1354
1355 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1356 auto force_field =
1357 bc->block->template getData<VectorField>(m_last_applied_force_field_id);
1358 auto vel_field =
1359 bc->block->template getData<VectorField>(m_velocity_field_id);
1360 std::array<FloatType, Stencil::Size> pop;
1361 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1362 pop[f] = FloatType_c(population[f]);
1363 }
1364 lbm::accessor::Population::set(pdf_field, vel_field, force_field, pop,
1365 bc->cell);
1366
1367 return true;
1368 }
1369
1370 std::vector<double>
1372 Utils::Vector3i const &upper_corner) const override {
1373 std::vector<double> out;
1374 auto const &lattice = get_lattice();
1375 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1376 out = std::vector<double>(stencil_size() * ci->numCells());
1377 for (auto const &block : *lattice.get_blocks()) {
1378 auto const block_offset = lattice.get_block_corner(block, true);
1379 if (auto const bci = get_block_interval(
1380 lattice, lower_corner, upper_corner, block_offset, block)) {
1381 auto const pdf_field =
1382 block.template getData<PdfField>(m_pdf_field_id);
1383 auto const values = lbm::accessor::Population::get(pdf_field, *bci);
1384 assert(values.size() == stencil_size() * bci->numCells());
1385
1386 auto kernel = [&values, &out, this](unsigned const block_index,
1387 unsigned const local_index,
1388 Utils::Vector3i const &) {
1389 for (uint_t f = 0u; f < stencil_size(); ++f) {
1390 out[stencil_size() * local_index + f] =
1391 values[stencil_size() * block_index + f];
1392 }
1393 };
1394
1395 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1396 }
1397 }
1398 }
1399 return out;
1400 }
1401
1402 void set_slice_population(Utils::Vector3i const &lower_corner,
1403 Utils::Vector3i const &upper_corner,
1404 std::vector<double> const &population) override {
1405 auto const &lattice = get_lattice();
1406 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1407 assert(population.size() == stencil_size() * ci->numCells());
1408 for (auto &block : *lattice.get_blocks()) {
1409 auto const block_offset = lattice.get_block_corner(block, true);
1410 if (auto const bci = get_block_interval(
1411 lattice, lower_corner, upper_corner, block_offset, block)) {
1412 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1413 auto force_field = block.template getData<VectorField>(
1415 auto vel_field =
1416 block.template getData<VectorField>(m_velocity_field_id);
1417 std::vector<FloatType> values(stencil_size() * bci->numCells());
1418
1419 auto kernel = [&values, &population, this](unsigned const block_index,
1420 unsigned const local_index,
1421 Utils::Vector3i const &) {
1422 for (uint_t f = 0u; f < stencil_size(); ++f) {
1423 values[stencil_size() * block_index + f] =
1424 numeric_cast<FloatType>(
1425 population[stencil_size() * local_index + f]);
1426 }
1427 };
1428
1429 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1430 lbm::accessor::Population::set(pdf_field, vel_field, force_field,
1431 values, *bci);
1432 }
1433 }
1434 }
1435 }
1436
1437 // Density
1438 std::optional<double>
1440 bool consider_ghosts = false) const override {
1441 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::PDF)));
1442 auto bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1443 if (!bc)
1444 return std::nullopt;
1445
1446 auto pdf_field =
1447 bc->block->template uncheckedFastGetData<PdfField>(m_pdf_field_id);
1448 auto const density =
1449 lbm::accessor::Density::get(pdf_field, m_density, bc->cell);
1450 return {double_c(density)};
1451 }
1452
1453 bool set_node_density(Utils::Vector3i const &node, double density) override {
1455 auto bc = get_block_and_cell(get_lattice(), node, false);
1456 if (!bc)
1457 return false;
1458
1459 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1461 bc->cell);
1462
1463 return true;
1464 }
1465
1466 std::vector<double>
1468 Utils::Vector3i const &upper_corner) const override {
1469 std::vector<double> out;
1470 auto const &lattice = get_lattice();
1471 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1472 out = std::vector<double>(ci->numCells());
1473 for (auto const &block : *lattice.get_blocks()) {
1474 auto const block_offset = lattice.get_block_corner(block, true);
1475 if (auto const bci = get_block_interval(
1476 lattice, lower_corner, upper_corner, block_offset, block)) {
1477 auto const pdf_field =
1478 block.template getData<PdfField>(m_pdf_field_id);
1479 auto const values =
1480 lbm::accessor::Density::get(pdf_field, m_density, *bci);
1481 assert(values.size() == bci->numCells());
1482
1483 auto kernel = [&values, &out](unsigned const block_index,
1484 unsigned const local_index,
1485 Utils::Vector3i const &) {
1486 out[local_index] = values[block_index];
1487 };
1488
1489 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1490 }
1491 }
1492 }
1493 return out;
1494 }
1495
1496 void set_slice_density(Utils::Vector3i const &lower_corner,
1497 Utils::Vector3i const &upper_corner,
1498 std::vector<double> const &density) override {
1500 auto const &lattice = get_lattice();
1501 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1502 assert(density.size() == ci->numCells());
1503 for (auto &block : *lattice.get_blocks()) {
1504 auto const block_offset = lattice.get_block_corner(block, true);
1505 if (auto const bci = get_block_interval(
1506 lattice, lower_corner, upper_corner, block_offset, block)) {
1507 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1508 std::vector<FloatType> values(bci->numCells());
1509
1510 auto kernel = [&values, &density](unsigned const block_index,
1511 unsigned const local_index,
1512 Utils::Vector3i const &) {
1513 values[block_index] = numeric_cast<FloatType>(density[local_index]);
1514 };
1515
1516 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1517 lbm::accessor::Density::set(pdf_field, values, m_density, *bci);
1518 }
1519 }
1520 }
1521 }
1522
1523 std::optional<Utils::Vector3d>
1525 bool consider_ghosts = false) const override {
1526 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::UBB)));
1527 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1528 if (!bc or !m_boundary->node_is_boundary(node))
1529 return std::nullopt;
1530
1531 return {to_vector3d(m_boundary->get_node_value_at_boundary(node))};
1532 }
1533
1535 Utils::Vector3d const &velocity) override {
1538 auto bc = get_block_and_cell(get_lattice(), node, false);
1539 if (!bc) {
1540 bc = get_block_and_cell(get_lattice(), node, true);
1541 }
1542 if (bc) {
1543 m_boundary->set_node_value_at_boundary(
1544 node, to_vector3<FloatType>(velocity), *bc);
1545 }
1546 return bc.has_value();
1547 }
1548
1549 std::vector<std::optional<Utils::Vector3d>> get_slice_velocity_at_boundary(
1550 Utils::Vector3i const &lower_corner,
1551 Utils::Vector3i const &upper_corner) const override {
1552 std::vector<std::optional<Utils::Vector3d>> out;
1553 auto const &lattice = get_lattice();
1554 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1555 out = std::vector<std::optional<Utils::Vector3d>>(ci->numCells());
1556 for (auto const &block : *lattice.get_blocks()) {
1557 auto const block_offset = lattice.get_block_corner(block, true);
1558 if (auto const bci = get_block_interval(
1559 lattice, lower_corner, upper_corner, block_offset, block)) {
1560
1561 auto kernel = [&out, this](unsigned const, unsigned const local_index,
1562 Utils::Vector3i const &node) {
1563 if (m_boundary->node_is_boundary(node)) {
1564 out[local_index] =
1565 to_vector3d(m_boundary->get_node_value_at_boundary(node));
1566 } else {
1567 out[local_index] = std::nullopt;
1568 }
1569 };
1570
1571 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1572 }
1573 }
1574 assert(out.size() == ci->numCells());
1575 }
1576 return out;
1577 }
1578
1580 Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner,
1581 std::vector<std::optional<Utils::Vector3d>> const &velocity) override {
1584 auto const &lattice = get_lattice();
1585 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1586 assert(velocity.size() == ci->numCells());
1587 for (auto &block : *lattice.get_blocks()) {
1588 auto const block_offset = lattice.get_block_corner(block, true);
1589 if (auto const bci = get_block_interval(
1590 lattice, lower_corner, upper_corner, block_offset, block)) {
1591
1592 auto kernel = [&, this](unsigned const, unsigned const local_index,
1593 Utils::Vector3i const &node) {
1594 auto const bc = get_block_and_cell(lattice, node, false);
1595 assert(bc->block->getAABB() == block.getAABB());
1596 auto const &opt = velocity[local_index];
1597 if (opt) {
1598 m_boundary->set_node_value_at_boundary(
1599 node, to_vector3<FloatType>(*opt), *bc);
1600 } else {
1601 m_boundary->remove_node_from_boundary(node, *bc);
1602 }
1603 };
1604
1605 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1606 }
1607 }
1608 }
1609 }
1610
1611 std::optional<Utils::Vector3d>
1612 get_node_boundary_force(Utils::Vector3i const &node) const override {
1613 auto const bc = get_block_and_cell(get_lattice(), node, true);
1614 if (!bc or !m_boundary->node_is_boundary(node))
1615 return std::nullopt;
1616
1617 return get_node_last_applied_force(node, true);
1618 }
1619
1620 bool remove_node_from_boundary(Utils::Vector3i const &node) override {
1621 auto bc = get_block_and_cell(get_lattice(), node, false);
1622 if (!bc) {
1623 bc = get_block_and_cell(get_lattice(), node, true);
1624 }
1625 if (bc) {
1626 m_boundary->remove_node_from_boundary(node, *bc);
1627 }
1628 return bc.has_value();
1629 }
1630
1631 std::optional<bool>
1633 bool consider_ghosts = false) const override {
1634 assert(not(consider_ghosts and m_pending_ghost_comm.test(GhostComm::UBB)));
1635 auto const bc = get_block_and_cell(get_lattice(), node, consider_ghosts);
1636 if (!bc)
1637 return std::nullopt;
1638
1639 return {m_boundary->node_is_boundary(node)};
1640 }
1641
1642 std::vector<bool>
1644 Utils::Vector3i const &upper_corner) const override {
1645 std::vector<bool> out;
1646 auto const &lattice = get_lattice();
1647 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1648 out = std::vector<bool>(ci->numCells());
1649 for (auto const &block : *lattice.get_blocks()) {
1650 auto const block_offset = lattice.get_block_corner(block, true);
1651 if (auto const bci = get_block_interval(
1652 lattice, lower_corner, upper_corner, block_offset, block)) {
1653
1654 auto kernel = [&out, this](unsigned const, unsigned const local_index,
1655 Utils::Vector3i const &node) {
1656 out[local_index] = m_boundary->node_is_boundary(node);
1657 };
1658
1659 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1660 }
1661 }
1662 assert(out.size() == ci->numCells());
1663 }
1664 return out;
1665 }
1666
1667 void reallocate_ubb_field() override { m_boundary->boundary_update(); }
1668
1670 if (not m_has_boundaries) {
1671 m_has_boundaries = true;
1673 }
1674 m_has_boundaries = true;
1675 }
1676
1677 void clear_boundaries() override {
1678 reset_boundary_handling(get_lattice().get_blocks());
1681 m_has_boundaries = false;
1683 }
1684
1685 void
1686 update_boundary_from_shape(std::vector<int> const &raster_flat,
1687 std::vector<double> const &data_flat) override {
1690 auto const &grid_size = get_lattice().get_grid_dimensions();
1691 auto data = fill_3D_vector_array(data_flat, grid_size);
1692 set_boundary_from_grid(*m_boundary, get_lattice(), raster_flat, data);
1695 }
1696
1697 // Pressure tensor
1698 std::optional<Utils::VectorXd<9>>
1699 get_node_pressure_tensor(Utils::Vector3i const &node) const override {
1700 auto bc = get_block_and_cell(get_lattice(), node, false);
1701 if (!bc)
1702 return std::nullopt;
1703
1704 auto pdf_field = bc->block->template getData<PdfField>(m_pdf_field_id);
1705 auto tensor =
1706 lbm::accessor::PressureTensor::get(pdf_field, m_density, bc->cell);
1707 pressure_tensor_correction(tensor);
1708 return to_vector9d(tensor);
1709 }
1710
1711 std::vector<double> get_slice_pressure_tensor(
1712 Utils::Vector3i const &lower_corner,
1713 Utils::Vector3i const &upper_corner) const override {
1714 std::vector<double> out;
1715 auto const &lattice = get_lattice();
1716 if (auto const ci = get_interval(lattice, lower_corner, upper_corner)) {
1717 out = std::vector<double>(9u * ci->numCells());
1718 for (auto const &block : *lattice.get_blocks()) {
1719 auto const block_offset = lattice.get_block_corner(block, true);
1720 if (auto const bci = get_block_interval(
1721 lattice, lower_corner, upper_corner, block_offset, block)) {
1722 auto const pdf_field =
1723 block.template getData<PdfField>(m_pdf_field_id);
1724 auto values =
1726 assert(values.size() == 9u * bci->numCells());
1727
1728 auto kernel = [&values, &out, this](unsigned const block_index,
1729 unsigned const local_index,
1730 Utils::Vector3i const &) {
1731 pressure_tensor_correction(
1732 std::span<FloatType, 9ul>(&values[9u * block_index], 9ul));
1733 for (uint_t f = 0u; f < 9u; ++f) {
1734 out[9u * local_index + f] = values[9u * block_index + f];
1735 }
1736 };
1737
1738 copy_block_buffer(*bci, *ci, block_offset, lower_corner, kernel);
1739 }
1740 }
1741 }
1742 return out;
1743 }
1744
1745 [[nodiscard]] Utils::Vector3i flat_index_to_node(int index) const {
1746 Utils::Vector3i node({0, 0, 0});
1747 auto const &grid_size = get_lattice().get_grid_dimensions();
1748 node[2] = index % grid_size[2];
1749 int tmp = index / grid_size[2];
1750 node[1] = tmp % grid_size[1];
1751 node[0] = tmp / grid_size[1];
1752 return node;
1753 }
1754
1756 int dir) const {
1757 Utils::Vector3i neighbor({0, 0, 0});
1758 auto const &grid_size = get_lattice().get_grid_dimensions();
1759 auto constexpr neighbor_offset = DynamicUBB::neighborOffset;
1760 for (int i = 0; i < neighbor.size(); i++) {
1761 neighbor[i] =
1762 (node[i] - neighbor_offset[i][dir] + grid_size[i]) % grid_size[i];
1763 }
1764 return neighbor;
1765 }
1766
1768 std::vector<int> const &raster_flat) const override {
1769 Utils::Vector3d force({0, 0, 0});
1770 auto const &grid_size = get_lattice().get_grid_dimensions();
1771 for (auto &block : *get_lattice().get_blocks()) {
1772 auto const offset = get_lattice().get_block_corner(block, true);
1773 auto const &force_field = m_boundary->get_force_vector(&block);
1774 auto const &index_field = m_boundary->get_index_vector(&block);
1775 for (int i = 0; i < raster_flat.size(); i++) {
1776 if (raster_flat[i] != 0) {
1777 auto node = flat_index_to_node(i);
1778 if (get_lattice().node_in_local_halo(node)) {
1779 // shift node to local frame
1780 node = (node - offset + grid_size) % grid_size;
1781 for (int j = 0; j < index_field.size(); j++) {
1782 auto neighbor_node = get_neighbor_node(node, index_field[j].dir);
1783 if (index_field[j].x == neighbor_node[0] &&
1784 index_field[j].y == neighbor_node[1] &&
1785 index_field[j].z == neighbor_node[2]) {
1786 force[0] += force_field[j].F_0;
1787 force[1] += force_field[j].F_1;
1788 force[2] += force_field[j].F_2;
1789 }
1790 }
1791 }
1792 }
1793 }
1794 }
1795 return zero_centered_to_md(force);
1796 }
1797 // Global boundary force
1798 [[nodiscard]] Utils::Vector3d get_boundary_force() const override {
1799 Vector3<double> force(0.);
1800 for (auto &block : *get_lattice().get_blocks()) {
1801 force += m_boundary->get_total_force(&block);
1802 }
1803 return zero_centered_to_md(to_vector3d(force));
1804 }
1805 // Global pressure tensor
1806 [[nodiscard]] Utils::VectorXd<9> get_pressure_tensor() const override {
1807 Matrix3<FloatType> tensor(FloatType{0});
1808 for (auto const &block : *get_lattice().get_blocks()) {
1809 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1811 }
1812 auto const &grid_size = get_lattice().get_grid_dimensions();
1813 auto const number_of_nodes = Utils::product(grid_size);
1814 pressure_tensor_correction(tensor);
1815 return to_vector9d(tensor) * (1. / static_cast<double>(number_of_nodes));
1816 }
1817
1818 // Global momentum
1819 [[nodiscard]] Utils::Vector3d get_momentum() const override {
1820 Vector3<FloatType> mom(FloatType{0});
1821 for (auto const &block : *get_lattice().get_blocks()) {
1822 auto pdf_field = block.template getData<PdfField>(m_pdf_field_id);
1823 auto force_field =
1824 block.template getData<VectorField>(m_last_applied_force_field_id);
1825 mom += lbm::accessor::MomentumDensity::reduce(pdf_field, force_field,
1826 m_density);
1827 }
1828 return to_vector3d(mom);
1829 }
1830
1831 // Global external force
1832 void set_external_force(Utils::Vector3d const &ext_force) override {
1833 m_reset_force->set_ext_force(zero_centered_to_lb(ext_force));
1834 }
1835
1836 [[nodiscard]] Utils::Vector3d get_external_force() const noexcept override {
1837 return zero_centered_to_md(m_reset_force->get_ext_force());
1838 }
1839
1840 [[nodiscard]] double get_kT() const noexcept override {
1841 return static_cast<double>(m_kT);
1842 }
1843
1844 [[nodiscard]] unsigned int get_seed() const noexcept override {
1845 return m_seed;
1846 }
1847
1848 [[nodiscard]] std::optional<uint64_t> get_rng_state() const override {
1849 auto const cm =
1850 std::get_if<StreamCollisionModelThermalized>(&*m_collision_model);
1851 if (!cm or m_kT == 0.) {
1852 return std::nullopt;
1853 }
1854 return {static_cast<uint64_t>(cm->getTime_step())};
1855 }
1856
1857 void set_rng_state(uint64_t counter) override {
1858 auto const cm =
1859 std::get_if<StreamCollisionModelThermalized>(&*m_collision_model);
1860 if (!cm or m_kT == 0.) {
1861 throw std::runtime_error("This LB instance is unthermalized");
1862 }
1863 assert(counter <=
1864 static_cast<uint32_t>(std::numeric_limits<uint_t>::max()));
1865 cm->setTime_step(static_cast<uint32_t>(counter));
1866 }
1867
1868 [[nodiscard]] LatticeWalberla const &get_lattice() const noexcept override {
1869 return *m_lattice;
1870 }
1871
1872 [[nodiscard]] std::size_t get_velocity_field_id() const noexcept override {
1873 return m_velocity_field_id;
1874 }
1875
1876 [[nodiscard]] std::size_t get_force_field_id() const noexcept override {
1878 }
1879
1880 void register_vtk_field_filters(walberla::vtk::VTKOutput &vtk_obj) override {
1881 field::FlagFieldCellFilter<FlagField> fluid_filter(m_flag_field_id);
1882 fluid_filter.addFlag(Boundary_flag);
1883 vtk_obj.addCellExclusionFilter(fluid_filter);
1884 }
1885
1886protected:
1887 template <typename Field_T, uint_t F_SIZE_ARG, typename OutputType>
1888 class VTKWriter : public vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG> {
1889 public:
1890 VTKWriter(ConstBlockDataID const &block_id, std::string const &id,
1891 FloatType unit_conversion)
1892 : vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG>(id),
1893 m_block_id(block_id), m_field(nullptr),
1894 m_conversion(unit_conversion) {}
1895
1896 protected:
1897 void configure() override {
1898 WALBERLA_ASSERT_NOT_NULLPTR(this->block_);
1899 m_field = this->block_->template getData<Field_T>(m_block_id);
1900 }
1901
1902 ConstBlockDataID const m_block_id;
1903 Field_T const *m_field;
1904 FloatType const m_conversion;
1905 };
1906
1907 template <typename OutputType = float>
1908 class DensityVTKWriter : public VTKWriter<PdfField, 1u, OutputType> {
1909 public:
1911 using Base::Base;
1912 using Base::evaluate;
1913
1914 protected:
1915 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1916 cell_idx_t const z, cell_idx_t const) override {
1917 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1918 auto const density =
1919 lbm::accessor::Density::get(this->m_field, 1., {x, y, z});
1920 return numeric_cast<OutputType>(this->m_conversion * density);
1921 }
1922 };
1923
1924 template <typename OutputType = float>
1925 class VelocityVTKWriter : public VTKWriter<VectorField, 3u, OutputType> {
1926 public:
1928 using Base::Base;
1929 using Base::evaluate;
1930
1931 protected:
1932 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1933 cell_idx_t const z, cell_idx_t const f) override {
1934 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1935 auto const velocity =
1936 lbm::accessor::Vector::get(this->m_field, {x, y, z});
1937 return numeric_cast<OutputType>(this->m_conversion * velocity[uint_c(f)]);
1938 }
1939 };
1940
1941 template <typename OutputType = float>
1942 class PressureTensorVTKWriter : public VTKWriter<PdfField, 9u, OutputType> {
1943 public:
1945 using Base::Base;
1946 using Base::evaluate;
1947
1948 PressureTensorVTKWriter(ConstBlockDataID const &block_id,
1949 std::string const &id, FloatType unit_conversion,
1950 FloatType off_diag_factor)
1951 : Base(block_id, id, unit_conversion),
1952 m_off_diag_factor(off_diag_factor) {}
1953
1954 protected:
1955 OutputType evaluate(cell_idx_t const x, cell_idx_t const y,
1956 cell_idx_t const z, cell_idx_t const f) override {
1957 WALBERLA_ASSERT_NOT_NULLPTR(this->m_field);
1958 auto const pressure =
1959 lbm::accessor::PressureTensor::get(this->m_field, 1., {x, y, z});
1960 auto const revert_factor =
1961 (f == 0 or f == 4 or f == 8) ? FloatType{1} : m_off_diag_factor;
1962 return numeric_cast<OutputType>(this->m_conversion * revert_factor *
1963 pressure[uint_c(f)]);
1964 }
1965 FloatType const m_off_diag_factor;
1966 };
1967
1968public:
1969 void register_vtk_field_writers(walberla::vtk::VTKOutput &vtk_obj,
1970 LatticeModel::units_map const &units,
1971 int flag_observables) override {
1972#if defined(__CUDACC__)
1973 auto const allocate_cpu_field_if_empty =
1974 [&]<typename Field>(auto const &blocks, std::string name,
1975 std::optional<BlockDataID> &cpu_field) {
1976 if (not cpu_field) {
1977 cpu_field = field::addToStorage<Field>(
1978 blocks, name, FloatType{0}, field::fzyx,
1979 m_lattice->get_ghost_layers(), m_host_field_allocator);
1980 }
1981 };
1982#endif
1983 if (flag_observables & static_cast<int>(OutputVTK::density)) {
1984 auto const unit_conversion =
1985 FloatType_c(zero_centered_to_md(units.at("density")));
1986#if defined(__CUDACC__)
1987 if constexpr (Architecture == lbmpy::Arch::GPU) {
1988 auto const &blocks = m_lattice->get_blocks();
1989 allocate_cpu_field_if_empty.template operator()<PdfFieldCpu>(
1990 blocks, "pdfs_cpu", m_pdf_cpu_field_id);
1991 vtk_obj.addBeforeFunction(gpu::fieldCpyFunctor<PdfFieldCpu, PdfField>(
1992 blocks, *m_pdf_cpu_field_id, m_pdf_field_id));
1993 }
1994#endif
1995 vtk_obj.addCellDataWriter(std::make_shared<DensityVTKWriter<float>>(
1996 m_pdf_field_id, "density", unit_conversion));
1997 }
1998 if (flag_observables & static_cast<int>(OutputVTK::velocity_vector)) {
1999 auto const unit_conversion = FloatType_c(units.at("velocity"));
2000#if defined(__CUDACC__)
2001 if constexpr (Architecture == lbmpy::Arch::GPU) {
2002 auto const &blocks = m_lattice->get_blocks();
2003 allocate_cpu_field_if_empty.template operator()<VectorFieldCpu>(
2004 blocks, "vel_cpu", m_vel_cpu_field_id);
2005 vtk_obj.addBeforeFunction(
2006 gpu::fieldCpyFunctor<VectorFieldCpu, VectorField>(
2007 blocks, *m_vel_cpu_field_id, m_velocity_field_id));
2008 }
2009#endif
2010 vtk_obj.addCellDataWriter(std::make_shared<VelocityVTKWriter<float>>(
2011 m_velocity_field_id, "velocity_vector", unit_conversion));
2012 }
2013 if (flag_observables & static_cast<int>(OutputVTK::pressure_tensor)) {
2014 auto const unit_conversion =
2015 FloatType_c(zero_centered_to_md(units.at("pressure")));
2016#if defined(__CUDACC__)
2017 if constexpr (Architecture == lbmpy::Arch::GPU) {
2018 auto const &blocks = m_lattice->get_blocks();
2019 allocate_cpu_field_if_empty.template operator()<PdfFieldCpu>(
2020 blocks, "pdfs_cpu", m_pdf_cpu_field_id);
2021 vtk_obj.addBeforeFunction(gpu::fieldCpyFunctor<PdfFieldCpu, PdfField>(
2022 blocks, *m_pdf_cpu_field_id, m_pdf_field_id));
2023 }
2024#endif
2025 vtk_obj.addCellDataWriter(
2026 std::make_shared<PressureTensorVTKWriter<float>>(
2027 m_pdf_field_id, "pressure_tensor", unit_conversion,
2028 pressure_tensor_correction_factor()));
2029 }
2030 }
2031
2032 ~LBWalberlaImpl() override = default;
2033};
2034
2035} // namespace walberla
LBWalberlaBase provides the public interface of the LB waLBerla bridge.
Vector implementation and trait types for boost qvm interoperability.
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
Utils::Vector3i get_block_corner(IBlock const &block, bool lower) const
DEVICE_QUALIFIER constexpr size_type size() const noexcept
Definition Array.hpp:166
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
std::vector< double > get_slice_last_applied_force(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner) const override
void zero_centered_transform_impl(T &data, auto const factor) const
std::shared_ptr< RegularFullCommunicator > m_pdf_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
FieldTrait< FloatType, Architecture >::VectorField VectorField
std::optional< Utils::Vector3d > get_node_velocity(Utils::Vector3i const &node, bool consider_ghosts=false) const override
void reallocate_ubb_field() override
Utils::Vector3d get_boundary_force_from_shape(std::vector< int > const &raster_flat) const override
std::shared_ptr< RegularFullCommunicator > m_full_communicator
std::size_t get_force_field_id() const noexcept override
std::shared_ptr< CollisionModel > m_collision_model
void set_slice_velocity(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &velocity) override
unsigned int get_seed() const noexcept override
void integrate_vtk_writers() override
BoundaryModel::FlagField FlagField
FieldTrait< FloatType, Architecture >::template PackInfo< Field > PackInfo
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 zero_centered_to_md_in_place(auto &data) const
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
Utils::Vector3d get_boundary_force() 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
FieldTrait< FloatType >::PdfField _PdfField
std::function< bool(Utils::Vector3d const &)> make_lattice_position_checker(bool consider_points_in_halo) const 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
bool is_double_precision() const noexcept 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
Utils::Vector3i get_neighbor_node(Utils::Vector3i const &node, int dir) const
std::vector< Utils::Vector3d > get_velocities_at_pos(std::vector< Utils::Vector3d > const &pos) override
void ghost_communication_laf() override
std::optional< bool > get_node_is_boundary(Utils::Vector3i const &node, bool consider_ghosts=false) const override
FieldTrait< FloatType, Architecture >::template RegularCommScheme< Stencil > PDFStreamingCommunicator
Regular communicator.
std::shared_ptr< ResetForce< PdfField, VectorField > > m_reset_force
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
FieldTrait< FloatType, Architecture >::template RegularCommScheme< stencil::D3Q27 > RegularFullCommunicator
Full 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.
FieldTrait< FloatType, Architecture >::PdfField PdfField
detail::KernelTrait< FloatType, Architecture >::InitialPDFsSetter InitialPDFsSetter
auto zero_centered_to_lb(auto const &data) const
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
auto make_density_interpolation_kernel() const
auto make_force_interpolation_kernel() const
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
void ghost_communication_vel() override
Utils::Vector3d get_external_force() const noexcept override
FloatType FloatType_c(T t) const
std::shared_ptr< RegularFullCommunicator > m_laf_communicator
detail::KernelTrait< FloatType, Architecture >::StreamCollisionModelLeesEdwards StreamCollisionModelLeesEdwards
std::size_t get_velocity_field_id() const noexcept override
Utils::VectorXd< 9 > get_pressure_tensor() const override
ResourceObserver m_mpi_cart_comm_observer
std::shared_ptr< InterpolateAndShiftAtBoundary< _PdfField, FloatType > > m_lees_edwards_pdf_interpol_sweep
Utils::Vector3i flat_index_to_node(int index) const
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
LatticeWalberla::Lattice_T BlockStorage
Lattice model (e.g.
std::shared_ptr< UpdateVelFromPDF > m_update_velocities_from_pdf
void set_collision_model(std::unique_ptr< LeesEdwardsPack > &&lees_edwards_pack) override
std::shared_ptr< RegularFullCommunicator > m_vel_communicator
std::bitset< GhostComm::SIZE > m_pending_ghost_comm
detail::KernelTrait< FloatType, Architecture >::StreamCollisionModelThermalized StreamCollisionModelThermalized
LatticeWalberla const & get_lattice() const noexcept override
auto zero_centered_to_md(auto const &data) const
void set_slice_population(Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, std::vector< double > const &population) override
FieldTrait< FloatType, Architecture >::template BoundaryCommScheme< stencil::D3Q27 > BoundaryFullCommunicator
double get_density() const noexcept override
void ghost_communication_pdf() override
auto make_velocity_interpolation_kernel() const
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
std::vector< double > get_densities_at_pos(std::vector< Utils::Vector3d > const &pos) override
void zero_centered_to_lb_in_place(auto &data) const
bool set_node_velocity(Utils::Vector3i const &node, Utils::Vector3d const &v) override
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
FieldTrait< FloatType >::VectorField _VectorField
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
bool add_force_at_pos(Utils::Vector3d const &pos, Utils::Vector3d const &force) override
std::variant< StreamCollisionModelThermalized, StreamCollisionModelLeesEdwards > CollisionModel
void setup_boundary_handle(std::shared_ptr< LatticeWalberla > lattice, std::shared_ptr< Boundary_T > boundary)
static constexpr std::array< std::array< int, 19u >, 3u > neighborOffset
static double * block(double *p, std::size_t index, std::size_t size)
Definition elc.cpp:176
T product(Vector< T, N > const &v)
Definition Vector.hpp:373
VectorXi< 3 > Vector3i
Definition Vector.hpp:194
STL namespace.
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, double const rho_in, double const density, Cell const &cell)
double get(GhostLayerField< double, uint_t{19u}> const *pdf_field, double const density, Cell const &cell)
void set(GhostLayerField< double, uint_t{19u}> const *pdf_field, GhostLayerField< double, uint_t{3u}> *velocity_field, GhostLayerField< double, uint_t{3u}> *force_field, Vector3< double > const &force, 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)
auto reduce(GhostLayerField< double, uint_t{19u}> const *pdf_field, GhostLayerField< double, uint_t{3u}> const *force_field, double const density)
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, std::array< double, 19u > const &pop, Cell const &cell)
auto get(GhostLayerField< double, uint_t{19u}> const *pdf_field, Cell const &cell)
void initialize(GhostLayerField< double, uint_t{19u}> *pdf_field, std::array< double, 19u > const &pop)
auto get(GhostLayerField< double, uint_t{19u}> const *pdf_field, double const density, Cell const &cell)
auto reduce(GhostLayerField< double, uint_t{19u}> const *pdf_field, double const density)
void initialize(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec)
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)
void set(GhostLayerField< double, uint_t{19u}> *pdf_field, GhostLayerField< double, uint_t{3u}> *velocity_field, GhostLayerField< double, uint_t{3u}> const *force_field, Vector3< double > const &u, 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, Function const &f)
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::optional< BlockAndCell > get_block_and_cell(::LatticeWalberla const &lattice, signed_integral_vector auto const &node, bool consider_ghost_layers)
Cell to_cell(signed_integral_vector auto const &xyz)
ResourceObserver get_mpi_cart_comm_observer()
Get an observer on waLBerla's MPI Cartesian communicator status.
std::optional< walberla::cell::CellInterval > get_block_interval(::LatticeWalberla const &lattice, Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner, Utils::Vector3i const &block_offset, IBlock const &block)
void copy_block_buffer(CellInterval const &bci, CellInterval const &ci, Utils::Vector3i const &block_offset, Utils::Vector3i const &lower_corner, Kernel &&kernel)
Synchronize data between a sliced block and a container.
auto to_vector9d(Matrix3< T > const &m) noexcept
std::optional< walberla::cell::CellInterval > get_interval(::LatticeWalberla const &lattice, Utils::Vector3i const &lower_corner, Utils::Vector3i const &upper_corner)
std::vector< Utils::Vector3d > fill_3D_vector_array(std::vector< double > const &vec_flat, Utils::Vector3i const &grid_size)
Definition boundary.hpp:36
static Utils::Vector3d velocity(Particle const &p_ref, Particle const &p_vs)
Velocity of the virtual site.
Definition relative.cpp:64
Observer to monitor the lifetime of a shared resource.
detail::KernelTrait< FT, AT >::PackInfoVec PackInfoStreamingVec
field::GhostLayerField< FT, uint_t{3u}> VectorField
detail::KernelTrait< FT, AT >::PackInfoPdf PackInfoStreamingPdf
field::communication::PackInfo< Field > PackInfo
field::GhostLayerField< FT, Stencil::Size > PdfField
blockforest::communication::UniformBufferedScheme< Stencil > RegularCommScheme
blockforest::communication::UniformBufferedScheme< Stencil > BoundaryCommScheme
GhostCommFlags
Ghost communication operations.
@ LAF
last applied forces communication