54 std::array<int, 3>
const &node,
double weight)
55 :
std::runtime_error(
"Access to LB " + field +
" field failed") {
57 <<
"], weight " <<
weight <<
"\n";
62 Utils::Interpolation::bspline_3d<2>(
67template <
typename FloatType, lbmpy::Arch Architecture>
71 auto const &
lat = *m_lattice;
84template <
typename FloatType, lbmpy::Arch Architecture>
86 std::vector<Utils::Vector3d>
const &pos,
87 std::vector<Utils::Vector3d>
const &forces) {
88 assert(pos.size() == forces.size());
93 auto const kernel = make_force_interpolation_kernel();
94 for (std::size_t i = 0
ul; i < pos.size(); ++i) {
95 kernel(pos[i], forces[i]);
98#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
100 auto const &lattice = get_lattice();
101 auto const &
block = *(lattice.get_blocks()->begin());
102 auto const origin =
block.getAABB().min();
107 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
108 for (
auto const &
vec : pos) {
110 for (std::size_t i : {0
ul, 1ul, 2ul}) {
111 host_pos.emplace_back(
static_cast<FloatType
>(
vec[i] - origin[i]));
114 for (
auto const &
vec : forces) {
116 for (std::size_t i : {0
ul, 1ul, 2ul}) {
121 auto const gl = lattice.get_ghost_layers();
123 m_force_to_be_applied_id);
129template <
typename FloatType, lbmpy::Arch Architecture>
132 auto const &lattice = *m_lattice;
133 auto const &
blocks = *lattice.get_blocks();
134 assert(lattice.get_ghost_layers() == 1u);
140 pos, [&,
conv = m_zc_to_lb,
field_id = m_force_to_be_applied_id](
141 std::array<int, 3>
const node,
double weight) {
147 blocks.transformGlobalToBlockLocalCell(cell, *
block);
158template <
typename FloatType, lbmpy::Arch Architecture>
159auto LBWalberlaImpl<FloatType,
160 Architecture>::make_velocity_interpolation_kernel()
const {
161 auto const &lattice = *m_lattice;
162 auto const &
blocks = *lattice.get_blocks();
163 assert(lattice.get_ghost_layers() == 1u);
167 pos, [&,
field_id = m_velocity_field_id](std::array<int, 3>
const node,
174 throw interpolation_illegal_access(
"velocity", pos, node,
weight);
176 if (m_has_boundaries
and m_boundary->node_is_boundary(node)) {
177 vel = m_boundary->get_node_value_at_boundary(node);
180 blocks.transformGlobalToBlockLocalCell(cell, *
block);
192template <
typename FloatType, lbmpy::Arch Architecture>
193auto LBWalberlaImpl<FloatType,
194 Architecture>::make_density_interpolation_kernel()
const {
195 auto const &lattice = *m_lattice;
196 auto const &
blocks = *lattice.get_blocks();
197 assert(lattice.get_ghost_layers() == 1u);
202 std::array<int, 3>
const node,
double weight) {
208 throw interpolation_illegal_access(
"density", pos, node,
weight);
210 blocks.transformGlobalToBlockLocalCell(cell, *
block);
227template <
typename FloatType, lbmpy::Arch Architecture>
228std::vector<Utils::Vector3d>
230 std::vector<Utils::Vector3d>
const &pos) {
234 std::vector<Utils::Vector3d> vel{};
235 vel.reserve(pos.size());
237 auto const kernel = make_velocity_interpolation_kernel();
238 std::ranges::transform(pos, std::back_inserter(vel), kernel);
240#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
242 auto const &lattice = get_lattice();
243 auto const &
block = *(lattice.get_blocks()->begin());
244 auto const origin =
block.getAABB().min();
247 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
248 for (
auto const &
vec : pos) {
250 for (std::size_t i : {0
ul, 1ul, 2ul}) {
251 host_pos.emplace_back(
static_cast<FloatType
>(
vec[i] - origin[i]));
254 auto const gl = lattice.get_ghost_layers();
259 auto const [
dev_idx,
dev_vel] = m_boundary->get_flattened_map_device();
266 static_cast<double>(*(
it + 1)),
267 static_cast<double>(*(
it + 2))});
274template <
typename FloatType, lbmpy::Arch Architecture>
277 std::vector<Utils::Vector3d>
const &pos) {
281 std::vector<double>
rho{};
282 rho.reserve(pos.size());
284 auto const kernel = make_density_interpolation_kernel();
285 std::ranges::transform(pos, std::back_inserter(
rho), kernel);
287#if defined(__CUDACC__) and defined(WALBERLA_BUILD_WITH_CUDA)
289 auto const &lattice = get_lattice();
290 auto const &
block = *(lattice.get_blocks()->begin());
291 auto const origin =
block.getAABB().min();
294 assert(lattice.get_blocks()->getNumberOfBlocks() == 1u);
295 for (
auto const &
vec : pos) {
297 for (std::size_t i : {0
ul, 1ul, 2ul}) {
298 host_pos.emplace_back(
static_cast<FloatType
>(
vec[i] - origin[i]));
301 auto const gl = lattice.get_ghost_layers();
305 if constexpr (std::is_same_v<FloatType, double>) {
308 for (
auto const &v :
res) {
309 rho.emplace_back(
static_cast<double>(v));
317template <
typename FloatType, lbmpy::Arch Architecture>
318std::optional<Utils::Vector3d>
321 assert(
not m_pending_ghost_comm.test(GhostComm::VEL));
322 assert(
not m_pending_ghost_comm.test(GhostComm::UBB));
327 auto const kernel = make_velocity_interpolation_kernel();
328 return {kernel(pos)};
331template <
typename FloatType, lbmpy::Arch Architecture>
335 assert(
not m_pending_ghost_comm.test(GhostComm::PDF));
340 auto const kernel = make_density_interpolation_kernel();
341 return {kernel(pos)};
344template <
typename FloatType, lbmpy::Arch Architecture>
347 if (!m_lattice->pos_in_local_halo(pos))
349 auto const kernel = make_force_interpolation_kernel();
Vector implementation and trait types for boost qvm interoperability.
static DEVICE_QUALIFIER constexpr Vector< T, N > broadcast(typename Base::value_type const &value) noexcept
Create a vector that has all entries set to the same value.
Class that runs and controls the LB on waLBerla.
void add_forces_at_pos(std::vector< Utils::Vector3d > const &pos, std::vector< Utils::Vector3d > const &forces) override
Distribute forces to the lattice at given positions.
std::optional< double > get_density_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
std::function< bool(Utils::Vector3d const &)> make_lattice_position_checker(bool consider_points_in_halo) const override
std::vector< Utils::Vector3d > get_velocities_at_pos(std::vector< Utils::Vector3d > const &pos) override
Interpolate velocities at given positions (batch version).
std::optional< Utils::Vector3d > get_velocity_at_pos(Utils::Vector3d const &pos, bool consider_points_in_halo=false) const override
std::vector< double > get_densities_at_pos(std::vector< Utils::Vector3d > const &pos) override
bool add_force_at_pos(Utils::Vector3d const &pos, Utils::Vector3d const &force) override
Exception for accessing a lattice node outside the local domain and ghost layers during B-spline inte...
interpolation_illegal_access(std::string const &field, Utils::Vector3d const &pos, std::array< int, 3 > const &node, double weight)
cudaStream_t stream[1]
CUDA streams for parallel computing on CPU and GPU.
static double * block(double *p, std::size_t index, std::size_t size)
double get(GhostLayerField< double, uint_t{19u}> const *pdf_field, double const density, Cell const &cell)
void add_force(gpu::GPUField< double > const *field, std::vector< double > const &pos, std::vector< double > const &forces, uint gl)
std::vector< double > get_rho(gpu::GPUField< double > const *field, std::vector< double > const &pos, double const density, uint gl)
std::vector< double > get_vel(gpu::GPUField< double > const *field, std::vector< double > const &pos, uint gl)
void set_from_list(gpu::GPUField< double > const *field, thrust::device_vector< int > const &indices, thrust::device_vector< double > const &values, uint gl)
void add(GhostLayerField< double, uint_t{3u}> *vec_field, Vector3< double > const &vec, Cell const &cell)
auto get(GhostLayerField< double, uint_t{3u}> const *vec_field, Cell const &cell)
\file PackInfoPdfDoublePrecision.cpp \author pystencils
IBlock * get_block_extended(LatticeWalberla const &lattice, auto const &pos, unsigned int n_ghost_layers)
auto to_vector3d(Vector3< T > const &v) noexcept
void interpolate_bspline_at_pos(Utils::Vector3d const &pos, auto const &&f)
Cell to_cell(signed_integral_vector auto const &xyz)