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