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>
48#include "../BoundaryHandling.hpp"
49#include "../BoundaryPackInfo.hpp"
50#include "../utils/boundary.hpp"
51#include "../utils/types_conversion.hpp"
55#if defined(__CUDACC__)
77#include <initializer_list>
92template <
typename FloatType, lbmpy::Arch Architecture>
96 detail::KernelTrait<FloatType,
99 detail::KernelTrait<FloatType,
121 template <
typename FT, lbmpy::Arch AT = lbmpy::Arch::CPU>
struct FieldTrait {
122 using PdfField = field::GhostLayerField<FT, Stencil::Size>;
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>;
136#if defined(__CUDACC__)
140 template <
class Field>
141 using MemcpyPackInfo = gpu::communication::MemcpyPackInfo<Field>;
144 template <
typename Stencil>
145 class UniformGPUScheme
146 :
public gpu::communication::UniformGPUScheme<Stencil> {
148 explicit UniformGPUScheme(
auto const &bf)
149 : gpu::communication::UniformGPUScheme<Stencil>(
155 template <
class Field>
using PackInfo = MemcpyPackInfo<Field>;
158 template <
class Stencil>
160 template <
class Stencil>
162 blockforest::communication::UniformBufferedScheme<Stencil>;
174#if defined(__CUDACC__)
175 using GPUField = gpu::GPUField<FloatType>;
193 return numeric_cast<FloatType>(t);
197 return static_cast<std::size_t
>(Stencil::Size);
201 return std::is_same_v<FloatType, double>;
205 class StreamCollideSweepVisitor {
210 cm.configure(m_storage, b);
215 cm.setV_s(
static_cast<decltype(cm.getV_s())
>(
216 m_lees_edwards_callbacks->get_shear_velocity()));
220 StreamCollideSweepVisitor() =
default;
221 StreamCollideSweepVisitor(std::shared_ptr<StructuredBlockStorage> storage) {
222 m_storage = std::move(storage);
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);
231 std::shared_ptr<StructuredBlockStorage> m_storage{};
232 std::shared_ptr<LeesEdwardsPack> m_lees_edwards_callbacks{};
234 StreamCollideSweepVisitor m_run_stream_collide_sweep{};
236 FloatType shear_mode_relaxation_rate()
const {
237 return FloatType{2} / (FloatType{6} *
m_viscosity + FloatType{1});
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} -
248 void reset_boundary_handling(std::shared_ptr<BlockStorage>
const &blocks) {
249 auto const [lc, uc] =
m_lattice->get_local_grid_range(
true);
255 FloatType pressure_tensor_correction_factor()
const {
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;
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;
273 class interpolation_illegal_access :
public std::runtime_error {
275 interpolation_illegal_access(std::string
const &field,
277 std::array<int, 3>
const &node,
double weight)
278 :
std::runtime_error(
"Access to LB " + field +
" field failed") {
280 <<
"], weight " << weight <<
"\n";
287 :
std::runtime_error(
"VTKOutput object '" + vtk_uid +
"' " + reason) {}
310#if defined(__CUDACC__)
311 std::optional<BlockDataID> m_pdf_cpu_field_id;
312 std::optional<BlockDataID> m_vel_cpu_field_id;
333 Architecture>::template RegularCommScheme<stencil::D3Q27>;
336 Architecture>::template BoundaryCommScheme<stencil::D3Q27>;
343 template <
class Field>
368 std::shared_ptr<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>
370 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
372 std::shared_ptr<InterpolateAndShiftAtBoundary<_VectorField, FloatType>>
375#if defined(__CUDACC__)
376 std::shared_ptr<gpu::HostFieldAllocator<FloatType>> m_host_field_allocator;
391 auto const &blocks =
m_lattice->get_blocks();
392 auto const n_ghost_layers =
m_lattice->get_ghost_layers();
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);
404 return field::addToStorage<Field>(blocks, tag, FloatType{0}, field::fzyx,
408#if defined(__CUDACC__)
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);
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>{});
430 auto const setup = [
this]<
typename PackInfoPdf,
typename PackInfoVec>() {
431 auto const &blocks =
m_lattice->get_blocks();
433 std::make_shared<PDFStreamingCommunicator>(blocks);
445 setup.template operator()<PackInfoPdf, PackInfoVec>();
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");
469#if defined(__CUDACC__)
470 m_host_field_allocator =
471 std::make_shared<gpu::HostFieldAllocator<FloatType>>();
477 for (
auto &
block : *blocks) {
483 blocks,
"flag field", n_ghost_layers);
485 reset_boundary_handling(
m_lattice->get_blocks());
509 std::make_shared<BoundaryFullCommunicator>(blocks);
513 auto boundary_packinfo = std::make_shared<
523 m_reset_force = std::make_shared<ResetForce<PdfField, VectorField>>(
532 void integrate_stream_collide(std::shared_ptr<BlockStorage>
const &blocks) {
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);
538 if (
auto *cm = std::get_if<StreamCollisionModelThermalized>(&cm_variant)) {
539 cm->setTime_step(cm->getTime_step() + 1u);
543 auto has_lees_edwards_bc()
const {
544 return std::holds_alternative<StreamCollisionModelLeesEdwards>(
548 void apply_lees_edwards_pdf_interpolation(
549 std::shared_ptr<BlockStorage>
const &blocks) {
550 for (
auto &
block : *blocks)
554 void apply_lees_edwards_vel_interpolation_and_shift(
555 std::shared_ptr<BlockStorage>
const &blocks) {
556 for (
auto &
block : *blocks)
560 void apply_lees_edwards_last_applied_force_interpolation(
561 std::shared_ptr<BlockStorage>
const &blocks) {
562 for (
auto &
block : *blocks)
566 void integrate_reset_force(std::shared_ptr<BlockStorage>
const &blocks) {
567 for (
auto &
block : *blocks)
571 void integrate_boundaries(std::shared_ptr<BlockStorage>
const &blocks) {
572 for (
auto &
block : *blocks)
576 void integrate_update_velocities_from_pdf(
577 std::shared_ptr<BlockStorage>
const &blocks) {
578 for (
auto &
block : *blocks)
582 void integrate_pull_scheme() {
586 integrate_reset_force(blocks);
588 integrate_stream_collide(blocks);
594 if (has_lees_edwards_bc()) {
595 apply_lees_edwards_pdf_interpolation(blocks);
596 apply_lees_edwards_last_applied_force_interpolation(blocks);
600 integrate_boundaries(blocks);
603 integrate_update_velocities_from_pdf(blocks);
605 if (has_lees_edwards_bc()) {
606 apply_lees_edwards_vel_interpolation_and_shift(blocks);
613 auto &vtk_handle = it.second;
614 if (vtk_handle->enabled) {
615 vtk::writeFiles(vtk_handle->ptr)();
616 vtk_handle->execution_count++;
623 integrate_pull_scheme();
641 if (has_lees_edwards_bc()) {
643 apply_lees_edwards_pdf_interpolation(blocks);
653 if (has_lees_edwards_bc()) {
655 apply_lees_edwards_vel_interpolation_and_shift(blocks);
665 if (has_lees_edwards_bc()) {
667 apply_lees_edwards_last_applied_force_interpolation(blocks);
684 if (has_lees_edwards_bc()) {
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);
700 auto const omega = shear_mode_relaxation_rate();
701 auto const omega_odd = odd_mode_relaxation_rate(omega);
710 m_run_stream_collide_sweep = StreamCollideSweepVisitor(blocks);
715 std::unique_ptr<LeesEdwardsPack> &&lees_edwards_pack)
override {
717#if defined(__CUDACC__)
719 throw std::runtime_error(
"Lees-Edwards LB doesn't support GPU yet");
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\"");
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");
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;
751 lebc_top_index, omega, omega, omega_odd, omega, shear_vel));
753 m_run_stream_collide_sweep =
756 std::make_shared<InterpolateAndShiftAtBoundary<_PdfField, FloatType>>(
758 shear_direction, shear_plane_normal,
763 shear_direction, shear_plane_normal,
769 n_ghost_layers, shear_direction, shear_plane_normal,
775 unsigned int shear_plane_normal)
const override {
779 throw std::runtime_error(
780 "MD and LB Lees-Edwards boundary conditions disagree");
797 template <
typename T>
799 if constexpr (std::is_arithmetic_v<T>) {
800 static_assert(std::is_floating_point_v<T>);
801 data *=
static_cast<T
>(factor);
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; });
818 auto transformed_data = data;
820 return transformed_data;
824 auto transformed_data = data;
826 return transformed_data;
830 std::optional<Utils::Vector3d>
832 bool consider_ghosts =
false)
const override {
837 if (is_boundary and *is_boundary) {
845 auto field = bc->block->template uncheckedFastGetData<VectorField>(
860 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
865 auto vel = to_vector3<FloatType>(v);
875 std::vector<double> out;
876 uint_t values_size = 0;
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);
883 lattice, lower_corner, upper_corner, block_offset,
block)) {
887 assert(values.size() == 3u * bci->numCells());
888 values_size += 3u * bci->numCells();
890 auto kernel = [&values, &out,
this](
unsigned const block_index,
891 unsigned const local_index,
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];
899 for (uint_t f = 0u; f < 3u; ++f) {
900 out[3u * local_index + f] =
901 double_c(values[3u * block_index + f]);
909 assert(values_size == 3u * ci->numCells());
916 std::vector<double>
const &
velocity)
override {
920 if (
auto const ci =
get_interval(lattice, lower_corner, upper_corner)) {
922 for (
auto &
block : *lattice.get_blocks()) {
923 auto const block_offset = lattice.get_block_corner(
block,
true);
925 lattice, lower_corner, upper_corner, block_offset,
block)) {
927 auto force_field =
block.template getData<VectorField>(
931 std::vector<FloatType> values(3u * bci->numCells());
933 auto kernel = [&values, &
velocity](
unsigned const block_index,
934 unsigned const local_index,
936 for (uint_t f = 0u; f < 3u; ++f) {
937 values[3u * block_index + f] =
938 numeric_cast<FloatType>(
velocity[3u * local_index + f]);
950 [[nodiscard]]
bool is_gpu() const noexcept
override {
957 if (consider_points_in_halo) {
958 return [&](
Utils::Vector3d const &p) {
return lat.pos_in_local_halo(p); };
960 return [&](
Utils::Vector3d const &p) {
return lat.pos_in_local_domain(p); };
964 std::vector<Utils::Vector3d>
const &forces)
override {
965 assert(pos.size() == forces.size());
971 for (std::size_t i = 0ul; i < pos.size(); ++i) {
972 kernel(pos[i], forces[i]);
975#if defined(__CUDACC__)
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) {
987 for (std::size_t i : {0ul, 1ul, 2ul}) {
988 host_pos.emplace_back(
static_cast<FloatType
>(vec[i] - origin[i]));
991 for (
auto const &vec : forces) {
993 for (std::size_t i : {0ul, 1ul, 2ul}) {
994 host_force.emplace_back(
static_cast<FloatType
>(vec[i]));
998 auto const gl = lattice.get_ghost_layers();
999 auto field =
block.template uncheckedFastGetData<VectorField>(
1008 auto const &blocks = *lattice.get_blocks();
1009 assert(lattice.get_ghost_layers() == 1u);
1016 std::array<int, 3>
const node,
double weight) {
1022 blocks.transformGlobalToBlockLocalCell(cell, *
block);
1024 auto const weighted_force = to_vector3<FloatType>(weight * force);
1026 block->template uncheckedFastGetData<VectorField>(field_id);
1035 auto const &blocks = *lattice.get_blocks();
1036 assert(lattice.get_ghost_layers() == 1u);
1040 std::array<int, 3>
const node,
1047 throw interpolation_illegal_access(
"velocity", pos, node, weight);
1048 Vector3<FloatType> vel;
1050 vel =
m_boundary->get_node_value_at_boundary(node);
1053 blocks.transformGlobalToBlockLocalCell(cell, *
block);
1055 block->template uncheckedFastGetData<VectorField>(field_id);
1067 auto const &blocks = *lattice.get_blocks();
1068 assert(lattice.get_ghost_layers() == 1u);
1073 std::array<int, 3>
const node,
1080 throw interpolation_illegal_access(
"density", pos, node, weight);
1082 blocks.transformGlobalToBlockLocalCell(cell, *
block);
1083 auto field =
block->template uncheckedFastGetData<PdfField>(field_id);
1085 acc += rho * weight;
1092 std::vector<Utils::Vector3d>
1097 std::vector<Utils::Vector3d> vel{};
1098 vel.reserve(pos.size());
1101 std::ranges::transform(pos, std::back_inserter(vel), kernel);
1103#if defined(__CUDACC__)
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) {
1113 for (std::size_t i : {0ul, 1ul, 2ul}) {
1114 host_pos.emplace_back(
static_cast<FloatType
>(vec[i] - origin[i]));
1117 auto const gl = lattice.get_ghost_layers();
1122 auto const [dev_idx, dev_vel] =
m_boundary->get_flattened_map_device();
1123 if (not dev_idx->empty()) {
1128 for (
auto it = res.begin(); it != res.end(); it += 3) {
1130 static_cast<double>(*(it + 1)),
1131 static_cast<double>(*(it + 2))});
1143 std::vector<double> rho{};
1144 rho.reserve(pos.size());
1147 std::ranges::transform(pos, std::back_inserter(rho), kernel);
1149#if defined(__CUDACC__)
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) {
1159 for (std::size_t i : {0ul, 1ul, 2ul}) {
1160 host_pos.emplace_back(
static_cast<FloatType
>(vec[i] - origin[i]));
1163 auto const gl = lattice.get_ghost_layers();
1168 if constexpr (std::is_same_v<FloatType, double>) {
1169 std::swap(rho, res);
1171 for (
auto const &v : res) {
1172 rho.emplace_back(
static_cast<double>(v));
1180 std::optional<Utils::Vector3d>
1182 bool consider_points_in_halo =
false)
const override {
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;
1190 return {kernel(pos)};
1193 std::optional<double>
1195 bool consider_points_in_halo =
false)
const override {
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;
1202 return {kernel(pos)};
1214 std::optional<Utils::Vector3d>
1218 return std::nullopt;
1226 std::optional<Utils::Vector3d>
1228 bool consider_ghosts =
false)
const override {
1232 return std::nullopt;
1248 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
1253 auto const vec = to_vector3<FloatType>(force);
1263 std::vector<double> out;
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);
1270 lattice, lower_corner, upper_corner, block_offset,
block)) {
1271 auto const field =
block.template getData<VectorField>(
1274 assert(values.size() == 3u * bci->numCells());
1276 auto kernel = [&values, &out](
unsigned const block_index,
1277 unsigned const local_index,
1279 for (uint_t f = 0u; f < 3u; ++f) {
1280 out[3u * local_index + f] = values[3u * block_index + f];
1294 std::vector<double>
const &force)
override {
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);
1303 lattice, lower_corner, upper_corner, block_offset,
block)) {
1305 auto force_field =
block.template getData<VectorField>(
1309 std::vector<FloatType> values(3u * bci->numCells());
1311 auto kernel = [&values, &force](
unsigned const block_index,
1312 unsigned const local_index,
1314 for (uint_t f = 0u; f < 3u; ++f) {
1315 values[3u * block_index + f] =
1316 numeric_cast<FloatType>(force[3u * local_index + f]);
1329 std::optional<std::vector<double>>
1331 bool consider_ghosts =
false)
const override {
1335 return std::nullopt;
1337 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
1339 std::vector<double> population(Stencil::Size);
1340 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1341 population[f] = double_c(pop[f]);
1344 return {std::move(population)};
1348 std::vector<double>
const &population)
override {
1355 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
1360 std::array<FloatType, Stencil::Size> pop;
1361 for (uint_t f = 0u; f < Stencil::Size; ++f) {
1373 std::vector<double> out;
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);
1380 lattice, lower_corner, upper_corner, block_offset,
block)) {
1381 auto const pdf_field =
1384 assert(values.size() ==
stencil_size() * bci->numCells());
1386 auto kernel = [&values, &out,
this](
unsigned const block_index,
1387 unsigned const local_index,
1404 std::vector<double>
const &population)
override {
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);
1411 lattice, lower_corner, upper_corner, block_offset,
block)) {
1413 auto force_field =
block.template getData<VectorField>(
1417 std::vector<FloatType> values(
stencil_size() * bci->numCells());
1419 auto kernel = [&values, &population,
this](
unsigned const block_index,
1420 unsigned const local_index,
1424 numeric_cast<FloatType>(
1438 std::optional<double>
1440 bool consider_ghosts =
false)
const override {
1444 return std::nullopt;
1447 bc->block->template uncheckedFastGetData<PdfField>(
m_pdf_field_id);
1459 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
1469 std::vector<double> out;
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);
1476 lattice, lower_corner, upper_corner, block_offset,
block)) {
1477 auto const pdf_field =
1481 assert(values.size() == bci->numCells());
1483 auto kernel = [&values, &out](
unsigned const block_index,
1484 unsigned const local_index,
1486 out[local_index] = values[block_index];
1498 std::vector<double>
const &
density)
override {
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);
1506 lattice, lower_corner, upper_corner, block_offset,
block)) {
1508 std::vector<FloatType> values(bci->numCells());
1510 auto kernel = [&values, &
density](
unsigned const block_index,
1511 unsigned const local_index,
1513 values[block_index] = numeric_cast<FloatType>(
density[local_index]);
1523 std::optional<Utils::Vector3d>
1525 bool consider_ghosts =
false)
const override {
1528 if (!bc or !
m_boundary->node_is_boundary(node))
1529 return std::nullopt;
1544 node, to_vector3<FloatType>(
velocity), *bc);
1546 return bc.has_value();
1552 std::vector<std::optional<Utils::Vector3d>> out;
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);
1559 lattice, lower_corner, upper_corner, block_offset,
block)) {
1561 auto kernel = [&out,
this](
unsigned const,
unsigned const local_index,
1567 out[local_index] = std::nullopt;
1574 assert(out.size() == ci->numCells());
1581 std::vector<std::optional<Utils::Vector3d>>
const &
velocity)
override {
1585 if (
auto const ci =
get_interval(lattice, lower_corner, upper_corner)) {
1587 for (
auto &
block : *lattice.get_blocks()) {
1588 auto const block_offset = lattice.get_block_corner(
block,
true);
1590 lattice, lower_corner, upper_corner, block_offset,
block)) {
1592 auto kernel = [&,
this](
unsigned const,
unsigned const local_index,
1595 assert(bc->block->getAABB() ==
block.getAABB());
1596 auto const &opt =
velocity[local_index];
1599 node, to_vector3<FloatType>(*opt), *bc);
1601 m_boundary->remove_node_from_boundary(node, *bc);
1611 std::optional<Utils::Vector3d>
1614 if (!bc or !
m_boundary->node_is_boundary(node))
1615 return std::nullopt;
1626 m_boundary->remove_node_from_boundary(node, *bc);
1628 return bc.has_value();
1633 bool consider_ghosts =
false)
const override {
1637 return std::nullopt;
1645 std::vector<bool> out;
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);
1652 lattice, lower_corner, upper_corner, block_offset,
block)) {
1654 auto kernel = [&out,
this](
unsigned const,
unsigned const local_index,
1656 out[local_index] =
m_boundary->node_is_boundary(node);
1662 assert(out.size() == ci->numCells());
1678 reset_boundary_handling(
get_lattice().get_blocks());
1687 std::vector<double>
const &data_flat)
override {
1698 std::optional<Utils::VectorXd<9>>
1702 return std::nullopt;
1704 auto pdf_field = bc->block->template getData<PdfField>(
m_pdf_field_id);
1707 pressure_tensor_correction(tensor);
1714 std::vector<double> out;
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);
1721 lattice, lower_corner, upper_corner, block_offset,
block)) {
1722 auto const pdf_field =
1726 assert(values.size() == 9u * bci->numCells());
1728 auto kernel = [&values, &out,
this](
unsigned const block_index,
1729 unsigned const local_index,
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];
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];
1760 for (
int i = 0; i < neighbor.size(); i++) {
1762 (node[i] - neighbor_offset[i][dir] + grid_size[i]) % grid_size[i];
1768 std::vector<int>
const &raster_flat)
const override {
1775 for (
int i = 0; i < raster_flat.size(); i++) {
1776 if (raster_flat[i] != 0) {
1780 node = (node - offset + grid_size) % grid_size;
1781 for (
int j = 0; j < index_field.size(); j++) {
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;
1799 Vector3<double> force(0.);
1807 Matrix3<FloatType> tensor(FloatType{0});
1814 pressure_tensor_correction(tensor);
1815 return to_vector9d(tensor) * (1. /
static_cast<double>(number_of_nodes));
1820 Vector3<FloatType> mom(FloatType{0});
1840 [[nodiscard]]
double get_kT() const noexcept
override {
1841 return static_cast<double>(
m_kT);
1844 [[nodiscard]]
unsigned int get_seed() const noexcept
override {
1851 if (!cm or
m_kT == 0.) {
1852 return std::nullopt;
1854 return {
static_cast<uint64_t
>(cm->getTime_step())};
1860 if (!cm or
m_kT == 0.) {
1861 throw std::runtime_error(
"This LB instance is unthermalized");
1864 static_cast<uint32_t
>(std::numeric_limits<uint_t>::max()));
1865 cm->setTime_step(
static_cast<uint32_t
>(counter));
1883 vtk_obj.addCellExclusionFilter(fluid_filter);
1887 template <
typename Field_T, u
int_t F_SIZE_ARG,
typename OutputType>
1888 class VTKWriter :
public vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG> {
1890 VTKWriter(ConstBlockDataID
const &block_id, std::string
const &
id,
1891 FloatType unit_conversion)
1892 : vtk::BlockCellDataWriter<OutputType, F_SIZE_ARG>(id),
1898 WALBERLA_ASSERT_NOT_NULLPTR(this->block_);
1907 template <
typename OutputType =
float>
1912 using Base::evaluate;
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);
1924 template <
typename OutputType =
float>
1929 using Base::evaluate;
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);
1941 template <
typename OutputType =
float>
1946 using Base::evaluate;
1949 std::string
const &
id, FloatType unit_conversion,
1950 FloatType off_diag_factor)
1951 :
Base(block_id, id, unit_conversion),
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 =
1960 auto const revert_factor =
1962 return numeric_cast<OutputType>(this->
m_conversion * revert_factor *
1963 pressure[uint_c(f)]);
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);
1984 auto const unit_conversion =
1986#if defined(__CUDACC__)
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>(
1999 auto const unit_conversion =
FloatType_c(units.at(
"velocity"));
2000#if defined(__CUDACC__)
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>(
2014 auto const unit_conversion =
2016#if defined(__CUDACC__)
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>(
2025 vtk_obj.addCellDataWriter(
2028 pressure_tensor_correction_factor()));
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.
auto const & get_grid_dimensions() const
walberla::blockforest::StructuredBlockForest Lattice_T
Utils::Vector3i get_block_corner(IBlock const &block, bool lower) const
DEVICE_QUALIFIER constexpr size_type size() const noexcept
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
FloatType const m_off_diag_factor
PressureTensorVTKWriter(ConstBlockDataID const &block_id, std::string const &id, FloatType unit_conversion, FloatType off_diag_factor)
void configure() override
VTKWriter(ConstBlockDataID const &block_id, std::string const &id, FloatType unit_conversion)
FloatType const m_conversion
ConstBlockDataID const m_block_id
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
BlockDataID m_pdf_tmp_field_id
std::size_t get_force_field_id() const noexcept override
BlockDataID m_pdf_field_id
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
BlockDataID m_vel_tmp_field_id
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
BlockDataID m_velocity_field_id
FieldTrait< FloatType, Architecture >::template RegularCommScheme< Stencil > PDFStreamingCommunicator
Regular communicator.
std::shared_ptr< ResetForce< PdfField, VectorField > > m_reset_force
void setup_streaming_communicator()
BlockDataID m_flag_field_id
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
void ghost_communication_full()
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
BlockDataID m_force_to_be_applied_id
FloatType FloatType_c(T t) const
std::shared_ptr< RegularFullCommunicator > m_laf_communicator
void apply_lees_edwards_interpolation()
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
void clear_boundaries() 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
void integrate() override
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
void ghost_communication_boundary()
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)
T product(Vector< T, N > const &v)
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)
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)
static Utils::Vector3d velocity(Particle const &p_ref, Particle const &p_vs)
Velocity of the virtual site.
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.
@ VEL
velocities communication
@ LAF
last applied forces communication
@ UBB
boundaries communication