summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--default.nix5
-rw-r--r--examples/moving_poiseulle_particles_2d_fplbm_gpu/sim.cpp15
-rw-r--r--examples/moving_poiseulle_particles_2d_psm_gpu/.nix/derivation.nix41
-rw-r--r--examples/moving_poiseulle_particles_2d_psm_gpu/SConscript34
-rw-r--r--examples/moving_poiseulle_particles_2d_psm_gpu/SConstruct81
-rw-r--r--examples/moving_poiseulle_particles_2d_psm_gpu/sim.cpp458
-rw-r--r--modules/core/c++/abstract/data.hpp89
-rw-r--r--modules/core/c++/abstract/error.hpp2
-rw-r--r--modules/core/c++/fplbm.hpp125
-rw-r--r--modules/core/c++/particle/aabb.hpp1
-rw-r--r--modules/core/c++/particle/cuboid/aabb.hpp40
-rw-r--r--modules/core/c++/particle/cuboid/common.hpp10
-rw-r--r--modules/core/c++/particle/cuboid/particle.hpp9
-rw-r--r--modules/core/c++/particle/cuboid/porosity.hpp10
-rw-r--r--modules/core/c++/particle/particle.hpp4
-rw-r--r--modules/core/c++/particle/spheroid/particle.hpp9
-rw-r--r--modules/core/c++/psm.hpp123
17 files changed, 1024 insertions, 32 deletions
diff --git a/default.nix b/default.nix
index 241d84a..f560d7e 100644
--- a/default.nix
+++ b/default.nix
@@ -132,6 +132,11 @@ in rec {
inherit kel;
};
+ moving_poiseulle_particles_2d_psm_gpu = pkgs.callPackage ./examples/moving_poiseulle_particles_2d_psm_gpu/.nix/derivation.nix {
+ inherit pname version stdenv forstio adaptive-cpp;
+ inherit kel;
+ };
+
poiseulle_particles_2d_psm_gpu = pkgs.callPackage ./examples/poiseulle_particles_2d_psm_gpu/.nix/derivation.nix {
inherit pname version stdenv forstio adaptive-cpp;
inherit kel;
diff --git a/examples/moving_poiseulle_particles_2d_fplbm_gpu/sim.cpp b/examples/moving_poiseulle_particles_2d_fplbm_gpu/sim.cpp
index f6c09e9..c82f632 100644
--- a/examples/moving_poiseulle_particles_2d_fplbm_gpu/sim.cpp
+++ b/examples/moving_poiseulle_particles_2d_fplbm_gpu/sim.cpp
@@ -141,7 +141,7 @@ saw::error_or<void> init(
rad.at({}) = dim_y * 0.1f;
saw::data<sch::Scalar<T>> dense;
dense.at({}) = 1.0f;
- particles = create_spheroid_particle_group<T,Desc::D,particle_amount>(rad,dense,{16u});
+ particles = create_spheroid_particle_group<T,Desc::D,particle_amount>(rad,dense,{1u});
{
auto& rb = particles.template get<"particles">().at({0u}).template get<"rigid_body">();
auto& rbp = rb.template get<"position">();
@@ -179,13 +179,13 @@ saw::error_or<void> step(
for(uint64_t i = 0u; i < Desc::D; ++i){
index.at({{i}}).set(idx[i]);
}
-
+
fplbm_reset.apply(fields,macros,index,t_i);
});
}).wait();
q.submit([&](acpp::sycl::handler& h){
- component<T,Desc,cmpt::FpLbmOneParticle,encode::Sycl<saw::encode::Native>> fplbm_one_part;
+ component<T,Desc,cmpt::FpLbmOneParticleImplicit,encode::Sycl<saw::encode::Native>> fplbm_one_part;
h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
saw::data<sch::FixedArray<sch::UInt64,1u>> index;
@@ -193,12 +193,13 @@ saw::error_or<void> step(
index.at({{i}}).set(idx[i]);
}
- fplbm_one_part.apply(fields,macros,particles,index,t_i,{1u});
+ fplbm_one_part.apply(fields,macros,particles,index,t_i,{16u});
});
}).wait();
// auto coll_ev =
q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{0.8};
component<T,Desc,cmpt::FpLbm,encode::Sycl<saw::encode::Native>> collision{0.8};
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
@@ -232,11 +233,11 @@ saw::error_or<void> step(
break;
case 3u:
flow_in.apply(fields,index,t_i);
- collision.apply(fields,macros,index,t_i);
+ bgk.apply(fields,macros,index,t_i);
break;
case 4u:
flow_out.apply(fields,index,t_i);
- collision.apply(fields,macros,index,t_i);
+ bgk.apply(fields,macros,index,t_i);
break;
default:
break;
@@ -261,7 +262,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
auto& lbm_dir = eo_lbm_dir.get_value();
- auto out_dir = lbm_dir / "moving_poiseulle_particles_2d_fplbm_gpu";
+ auto out_dir = lbm_dir / "moving_poiseulle_particles_2d_fplbm_imp_gpu";
{
std::error_code ec;
diff --git a/examples/moving_poiseulle_particles_2d_psm_gpu/.nix/derivation.nix b/examples/moving_poiseulle_particles_2d_psm_gpu/.nix/derivation.nix
new file mode 100644
index 0000000..01e1f12
--- /dev/null
+++ b/examples/moving_poiseulle_particles_2d_psm_gpu/.nix/derivation.nix
@@ -0,0 +1,41 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, python3
+, pname
+, version
+, adaptive-cpp
+, kel
+}:
+
+stdenv.mkDerivation {
+ pname = pname + "-examples-" + "moving_poiseulle_particles_2d_psm_gpu";
+ inherit version;
+ src = ./..;
+
+ nativeBuildInputs = [
+ scons
+ clang-tools
+ python3
+ ];
+
+ buildInputs = [
+ forstio.core
+ forstio.async
+ forstio.codec
+ forstio.codec-unit
+ forstio.io
+ forstio.remote
+ forstio.remote-filesystem
+ forstio.codec-json
+ adaptive-cpp
+ kel.lbm.core
+ kel.lbm.sycl
+ ];
+
+ preferLocalBuild = true;
+
+ outputs = [ "out" "dev" ];
+}
diff --git a/examples/moving_poiseulle_particles_2d_psm_gpu/SConscript b/examples/moving_poiseulle_particles_2d_psm_gpu/SConscript
new file mode 100644
index 0000000..e8102aa
--- /dev/null
+++ b/examples/moving_poiseulle_particles_2d_psm_gpu/SConscript
@@ -0,0 +1,34 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+examples_env['CXX'] = 'syclcc-clang';
+examples_env['CXXFLAGS'] += ['-O3'];
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+# Cavity2D
+examples_objects = [];
+examples_env.add_source_files(examples_objects, ['sim.cpp'], shared=False);
+examples_env.poiseulle_2d_gpu = examples_env.Program('#bin/moving_poiseulle_particles_2d_psm_gpu', [examples_objects]);
+
+# Set Alias
+env.examples = [
+ examples_env.poiseulle_2d_gpu
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/moving_poiseulle_particles_2d_psm_gpu/SConstruct b/examples/moving_poiseulle_particles_2d_psm_gpu/SConstruct
new file mode 100644
index 0000000..0611b67
--- /dev/null
+++ b/examples/moving_poiseulle_particles_2d_psm_gpu/SConstruct
@@ -0,0 +1,81 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+ def isbasestring(s):
+ return isinstance(s,basestring)
+else:
+ def isbasestring(s):
+ return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+ if isbasestring(filetype):
+ dir_path = self.Dir('.').abspath
+ filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+ for path in filetype:
+ target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+ if shared:
+ target_name+='.os'
+ sources.append( self.SharedObject( target=target_name, source=path ) )
+ else:
+ target_name+='.o'
+ sources.append( self.StaticObject( target=target_name, source=path ) )
+ pass
+
+def isAbsolutePath(key, dirname, env):
+ assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+ args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+ help='Installation target location of build results and headers',
+ default='/usr/local/',
+ validator=isAbsolutePath
+)
+
+env_vars.Add('build_examples',
+ help='If examples should be built',
+ default="true"
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+ CPPDEFINES=['SAW_UNIX'],
+ CXXFLAGS=[
+ '-std=c++20',
+ '-g',
+ '-Wall',
+ '-Wextra'
+ ],
+ LIBS=[
+ 'forstio-core',
+ 'forstio-async',
+ 'forstio-io'
+ ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/examples/moving_poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/moving_poiseulle_particles_2d_psm_gpu/sim.cpp
new file mode 100644
index 0000000..2c4c9d1
--- /dev/null
+++ b/examples/moving_poiseulle_particles_2d_psm_gpu/sim.cpp
@@ -0,0 +1,458 @@
+#include <kel/lbm/sycl/lbm.hpp>
+#include <kel/lbm/lbm.hpp>
+#include <kel/lbm/particle.hpp>
+
+#include <forstio/io/io.hpp>
+#include <forstio/remote/filesystem/easy.hpp>
+#include <forstio/codec/json/json.hpp>
+#include <forstio/codec/simple.hpp>
+
+namespace kel {
+namespace lbm {
+
+constexpr uint64_t dim_y = 256ul;
+constexpr uint64_t dim_x = dim_y * 20ul;
+
+constexpr uint64_t particle_amount = 1ul;
+
+namespace sch {
+using namespace saw::schema;
+
+using InfoChunk = Chunk<UInt8, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using ScalarChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using VectorChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using ChunkStruct = Struct<
+ Member<InfoChunk, "info">,
+ Member<DfChunk<T,Desc>, "dfs">,
+ Member<DfChunk<T,Desc>, "dfs_old">,
+ Member<VectorChunk<T,Desc>, "particle_N">,
+ Member<ScalarChunk<T,Desc>, "particle_D">
+>;
+
+template<typename T, typename Desc>
+using VelChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T>
+using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using MacroStruct = Struct<
+ Member<VelChunk<T,Desc>, "velocity">,
+ Member<RhoChunk<T>, "density">,
+ Member<ScalarChunk<T,Desc>, "porosity">,
+ Member<VelChunk<T,Desc>, "force">
+>;
+
+//template<typename T, typename Desc>
+//using ParticleArray = Array<
+// Particle<T,Desc::D>
+//>;
+
+template<typename T, typename Desc>
+using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,1u,coll::Spheroid<T>>;
+}
+
+template<typename T, typename Desc>
+saw::error_or<void> init(
+ saw::data<sch::ChunkStruct<T,Desc>>& fields,
+ saw::data<sch::MacroStruct<T,Desc>>& macros,
+ saw::data<sch::ParticleSpheroidGroup<T,Desc>>& particles
+){
+ auto& info_f = fields.template get<"info">();
+ // Set everything as walls
+ iterator<Desc::D>::apply(
+ [&](auto& index){
+ info_f.at(index).set(1u);
+ },
+ {},
+ info_f.get_dims(),
+ {}
+ );
+ // Fluid
+ iterator<Desc::D>::apply(
+ [&](auto& index){
+ info_f.at(index).set(2u);
+ },
+ {},
+ info_f.get_dims(),
+ {{1u,1u}}
+ );
+
+ // Inflow
+ iterator<Desc::D>::apply(
+ [&](auto& index){
+ info_f.at(index).set(3u);
+ },
+ {{0u,0u}},
+ {{1u,dim_y}},
+ {{0u,1u}}
+ );
+
+ // Outflow
+ iterator<Desc::D>::apply(
+ [&](auto& index){
+ info_f.at(index).set(4u);
+ },
+ {{dim_x-1u,0u}},
+ {{dim_x, dim_y}},
+ {{0u,1u}}
+ );
+ //
+ auto& dfs_old_f = fields.template get<"dfs_old">();
+ auto& dfs_f = fields.template get<"dfs">();
+ auto& rho_f = macros.template get<"density">();
+ auto& vel_f = macros.template get<"velocity">();
+ auto& por_f = macros.template get<"porosity">();
+ auto& force_f = macros.template get<"force">();
+
+ iterator<Desc::D>::apply(
+ [&](auto& index){
+ auto& dfs = dfs_f.at(index);
+ auto& rho = rho_f.at(index);
+ por_f.at(index).at({}) = {1};
+ rho.at({}) = {1};
+ auto& vel = vel_f.at(index);
+ for(uint64_t i{0u}; i < Desc::D; ++i){
+ vel.at({{i}}).set(0);
+ }
+ auto eq = equilibrium<T,Desc>(rho,vel);
+
+ dfs = eq;
+ dfs_old_f.at(index) = dfs;
+ auto& force = force_f.at(index);
+ for(uint64_t i{0u}; i < Desc::D; ++i){
+ force.at({{i}}).set(0);
+ }
+ },
+ {},// 0-index
+ dfs_f.get_dims()
+ );
+
+ saw::data<sch::Scalar<T>> rad;
+ rad.at({}) = dim_y * 0.1f;
+ saw::data<sch::Scalar<T>> dense;
+ dense.at({}) = 1.0f;
+ particles = create_spheroid_particle_group<T,Desc::D,particle_amount>(rad,dense,{1u});
+ {
+ auto& rb = particles.template get<"particles">().at({0u}).template get<"rigid_body">();
+ auto& rbp = rb.template get<"position">();
+ rbp.at({{0u}}) = 0.25 * dim_x;
+ rbp.at({{1u}}) = 0.5 * dim_y;
+ rb.template get<"position_old">() = rbp;
+ saw::data<sch::Vector<T,Desc::D>> zero;
+ zero.at({{0u}}) = 0.0f;
+ zero.at({{1u}}) = 0.0f;
+ rb.template get<"acceleration">() = zero;
+ rb.template get<"rotation">() = {};
+ rb.template get<"rotation_old">() = {};
+ rb.template get<"angular_acceleration">() = {};
+ }
+
+ return saw::make_void();
+}
+
+template<typename T, typename Desc>
+saw::error_or<void> step(
+ saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields,
+ saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros,
+ saw::data<sch::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& particles,
+ saw::data<sch::UInt64> t_i,
+ device& dev
+){
+ auto& q = dev.get_handle();
+ auto& info_f = fields.template get<"info">();
+ auto& porous_f = macros.template get<"porosity">();
+
+ q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::PsmReset,encode::Sycl<saw::encode::Native>> psm_reset;
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+ for(uint64_t i = 0u; i < Desc::D; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
+
+ psm_reset.apply(fields,macros,index,t_i);
+ });
+ }).wait();
+
+ q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::PsmOneParticle,encode::Sycl<saw::encode::Native>> psm_one_part;
+
+ h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+ saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+ for(uint64_t i = 0u; i < 1u; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
+
+ psm_one_part.apply(fields,macros,particles,index,t_i,{16u});
+ });
+ }).wait();
+
+ // auto coll_ev =
+ q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{0.8};
+ component<T,Desc,cmpt::Psm,encode::Sycl<saw::encode::Native>> collision{0.8};
+ component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+
+ component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
+ [&](){
+ uint64_t target_t_i = 8u;
+ if(t_i.get() < target_t_i){
+ return 1.0 + (0.01 / target_t_i) * t_i.get();
+ }
+ return 1.01;
+ }()
+ };
+ component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+ for(uint64_t i = 0u; i < Desc::D; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
+
+ auto info = info_f.at(index);
+
+ switch(info.get()){
+ case 0u:
+ break;
+ case 1u:
+ bb.apply(fields,index,t_i);
+ break;
+ case 2u:
+ collision.apply(fields,macros,index,t_i);
+ break;
+ case 3u:
+ flow_in.apply(fields,index,t_i);
+ bgk.apply(fields,macros,index,t_i);
+ break;
+ case 4u:
+ flow_out.apply(fields,index,t_i);
+ bgk.apply(fields,macros,index,t_i);
+ break;
+ default:
+ break;
+ }
+ });
+ }).wait();
+
+ return saw::make_void();
+}
+}
+}
+
+template<typename T, typename Desc>
+saw::error_or<void> lbm_main(int argc, char** argv){
+ using namespace kel::lbm;
+
+ using dfi = df_info<T,Desc>;
+
+ auto eo_lbm_dir = output_directory();
+ if(eo_lbm_dir.is_error()){
+ return std::move(eo_lbm_dir.get_error());
+ }
+ auto& lbm_dir = eo_lbm_dir.get_value();
+
+ auto out_dir = lbm_dir / "moving_poiseulle_particles_2d_psm_gpu";
+
+ {
+ std::error_code ec;
+ std::filesystem::create_directories(out_dir,ec);
+ if(ec != std::errc{}){
+ return saw::make_error<saw::err::critical>("Could not create output directory");
+ }
+ }
+
+ converter<T> conv {
+ // delta_x
+ {{1.0}},
+ // delta_t
+ {{1.0}}
+ };
+
+ print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{0.4 * dim_y});
+
+ // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
+ auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>();
+ auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>();
+ auto lbm_parts_data_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>();
+ //saw::data<sch::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& particles
+
+ std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl;
+
+ auto eo_aio = saw::setup_async_io();
+ if(eo_aio.is_error()){
+ return std::move(eo_aio.get_error());
+ }
+ auto& aio = eo_aio.get_value();
+ saw::wait_scope wait{aio.event_loop};
+
+ bool krun = true;
+ bool print_status = false;
+ aio.event_port.on_signal(saw::Signal::Terminate).then([&](){
+ krun = false;
+ }).detach();
+ aio.event_port.on_signal(saw::Signal::User1).then([&](){
+ print_status = true;
+ }).detach();
+
+ device dev;
+
+ auto& sycl_q = dev.get_handle();
+
+ sycl_q.wait();
+ {
+ auto eov = init<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_parts_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = write_vtk_file(out_dir,"initial_state",0u,*lbm_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+
+ saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q};
+ saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q};
+ saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_parts_data{sycl_q};
+ sycl_q.wait();
+
+ {
+ auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = dev.copy_to_device(*lbm_macro_data_ptr,lbm_sycl_macro_data);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = dev.copy_to_device(*lbm_parts_data_ptr,lbm_sycl_parts_data);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ sycl_q.wait();
+ auto lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
+ auto lsdp_view = make_view(lbm_sycl_parts_data);
+
+ saw::data<sch::UInt64> time_steps{32u*2048ul};
+ auto& info_f = lsd_view.template get<"info">();
+
+ for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
+ // BC + Collision
+ {
+ auto eov = step<T,Desc>(lsd_view,lsdm_view,lsdp_view,i,dev);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ sycl_q.wait();
+
+ if( i.get() % 32u == 0u){
+ {
+ auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = write_vtk_file(out_dir,"m",i.get(), *lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ }
+ {
+ auto eov = dev.copy_to_host(lbm_sycl_parts_data,*lbm_parts_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ auto& p = lbm_parts_data_ptr->template get<"particles">().at({0u}).template get<"rigid_body">();
+ auto& ppos = p.template get<"position">();
+ auto& pacc = p.template get<"acceleration">();
+ std::cout<<"Particle: "<<ppos.at({{0u}}).get()<<", "<<ppos.at({{1u}}).get()<<"; "<<pacc.at({{0u}}).get()<<", "<<pacc.at({{1u}}).get()<<";"<<i.get()<<std::endl;
+ }
+ // Stream
+ sycl_q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream;
+
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+ for(uint64_t i = 0u; i < Desc::D; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
+
+ auto info = info_f.at(index);
+
+ if(info.get() > 0u){
+ stream.apply(lsd_view,index,i);
+ }
+ });
+ }).wait();
+
+ // EPoll
+ wait.poll();
+ if(print_status){
+ std::cout<<"Status: "<<i.get()<<" of "<<time_steps.get()<<" - "<<(i.template cast_to<sch::Float64>().get() * 100 / time_steps.get())<<"%"<<std::endl;
+ print_status = false;
+ }
+ // print_progress_bar(i.get(), time_steps.get()-1u);
+ }
+
+ // After Loop
+ sycl_q.wait();
+ {
+ auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = write_vtk_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+ {
+ auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
+
+ sycl_q.wait();
+ return saw::make_void();
+}
+
+using FloatT = kel::lbm::sch::Float64;
+
+int main(int argc, char** argv){
+ auto eov = lbm_main<FloatT,kel::lbm::sch::D2Q9>(argc, argv);
+ if(eov.is_error()){
+ auto& err = eov.get_error();
+ std::cerr<<"[Error] "<<err.get_category();
+ auto err_msg = err.get_message();
+ if(err_msg.size() > 0u){
+ std::cerr<<" - "<<err_msg;
+ }
+ std::cerr<<std::endl;
+ return err.get_id();
+ }
+ return 0;
+}
diff --git a/modules/core/c++/abstract/data.hpp b/modules/core/c++/abstract/data.hpp
index f1ae5a7..327fb63 100644
--- a/modules/core/c++/abstract/data.hpp
+++ b/modules/core/c++/abstract/data.hpp
@@ -10,8 +10,97 @@ class data final {};
template<typename T, uint64_t N>
class data<sch::Primitive<T,N>,Encode> final {
+public:
+ using Schema = sch::Primitive<T,N>;
+ using Encode = Encode;
private:
+ native_data_type<Schema>::type value_;
public:
+ data():
+ value_{}
+ {}
+
+ data(native_data_type<Schema>::type value__):
+ value_{value__}
+ {}
+
+ constexpr auto get() const {
+ return value_;
+ }
+
+ void set(native_data_type value__){
+ value_ = value__;
+ }
+
+ constexpr bool operator==(const data<Schema,Encode>& rhs) const {
+ return value_ == rhs.value_;
+ }
+
+ constexpr bool operator!=(const data<Schema,Encode>& rhs) const {
+ return value_ != rhs.value_;
+ }
+
+ constexpr bool operator>(const data<Schema,Encode>& rhs) const {
+ return value_ > rhs.value_;
+ }
+
+ constexpr bool operator<(const data<Schema,Encode>& rhs) const {
+ return value_ < rhs.value_;
+ }
+
+ constexpr bool operator>=(const data<Schema,Encode>& rhs) const {
+ return value_ >= rhs.value_;
+ }
+
+ constexpr bool operator<=(const data<Schema,Encode>& rhs) const {
+ return value_ >= rhs.value_;
+ }
+
+ constexpr bool equals(const data<Schema,Encode>& rhs) const {
+ return (*this) == rhs;
+ }
+
+ data<Schema,Encode> operator+(const data<Schema,Encode>& rhs) const {
+ return {value_ + rhs.value_};
+ }
+
+ data<Schema,Encode> operator-(const data<Schema,Encode>& rhs) const {
+ return {value_ - rhs.value_};
+ }
+
+ data<Schema,Encode> operator*(const data<Schema,Encode>& rhs) const {
+ return {value_ * rhs.value_};
+ }
+
+ data<Schema,Encode> operator/(const data<Schema,Encode>& rhs) const {
+ return {value_ / rhs.value_};
+ }
+
+ data<Schema,Encode>& operator+=(const data<Schema,Encode>& rhs) {
+ value_ += rhs.value_;
+ return *this
+ }
+
+ data<Schema,Encode>& operator-=(const data<Schema,Encode>& rhs) {
+ value_ -= rhs.value_;
+ return *this
+ }
+
+ data<Schema,Encode>& operator*=(const data<Schema,Encode>& rhs) {
+ value_ *= rhs.value_;
+ return *this
+ }
+
+ data<Schema,Encode>& operator/=(const data<Schema,Encode>& rhs) {
+ value_ /= rhs.value_;
+ return *this
+ }
+
+ template<typename T>
+ data<T,Encode> cast_to() const {
+ data<T,Encode> val{static_cast<typename native_data_type<T>::type>(value_)};
+ return val;
+ }
};
diff --git a/modules/core/c++/abstract/error.hpp b/modules/core/c++/abstract/error.hpp
index 19917de..94a2c35 100644
--- a/modules/core/c++/abstract/error.hpp
+++ b/modules/core/c++/abstract/error.hpp
@@ -305,7 +305,7 @@ public:
/**
* This tries to catch cases where error starts including itself as a type which can happen in more complicated cases.
- * So this acts as a type safe guard.
+ * So this acts as a type safe guard in case internal handling fails to treat conversion correctly.
*/
template <typename T> class error_or<error_or<T>> {
private:
diff --git a/modules/core/c++/fplbm.hpp b/modules/core/c++/fplbm.hpp
index cfb0062..1f38f51 100644
--- a/modules/core/c++/fplbm.hpp
+++ b/modules/core/c++/fplbm.hpp
@@ -10,6 +10,7 @@ namespace cmpt {
struct FpLbmReset{};
struct FpLbm {};
struct FpLbmOneParticle {};
+struct FpLbmOneParticleImplicit {};
struct FpLbmOneParticleNoVelocity{};
}
namespace method {
@@ -36,18 +37,20 @@ public:
for(uint64_t i{0u}; i < Descriptor::D; ++i){
force.at({{i}}) = 0.0;
}
+ auto info = field.template get<"info">().at(index).get();
+ if(info == 2u){
+ bool is_even = ((time_step.get() % 2) == 0);
- bool is_even = ((time_step.get() % 2) == 0);
-
- auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
- auto& dfs = dfs_old_f.at(index);
+ auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
+ auto& dfs = dfs_old_f.at(index);
- auto& rho_f = macros.template get<"density">();
- auto& rho = rho_f.at(index);
- auto& vel_f = macros.template get<"velocity">();
- auto& vel = vel_f.at(index);
+ auto& rho_f = macros.template get<"density">();
+ auto& rho = rho_f.at(index);
+ auto& vel_f = macros.template get<"velocity">();
+ auto& vel = vel_f.at(index);
- compute_rho_u<T,Descriptor>(dfs,rho,vel);
+ compute_rho_u<T,Descriptor>(dfs,rho,vel);
+ }
}
};
@@ -83,7 +86,7 @@ public:
auto& vel_f = macros.template get<"velocity">();
saw::data<sch::Vector<T,Descriptor::D>>& vel = vel_f.at(index);
- compute_rho_u<T,Descriptor>(dfs,rho,vel);
+ // compute_rho_u<T,Descriptor>(dfs,rho,vel);
auto eq = equilibrium<T,Descriptor>(rho,vel);
using dfi = df_info<T,Descriptor>;
@@ -158,10 +161,94 @@ public:
auto& pirb_pos = pirb.template get<"position">();
auto& pirb_pos_old = pirb.template get<"position_old">();
+
+ auto& p_coll = pg.template get<"collision">().at({});
+ auto& p_rad = p_coll.template get<"radius">();
+
+ auto eo_aabb = particle_aabb<typename ParticleSchema::ValueType>::calculate(pg,{{0u}},vel_f.meta());
+ if(eo_aabb.is_error()){
+ return;
+ }
+ auto& aabb = eo_aabb.get_value();
+
+ saw::data<sch::Scalar<T>> two;
+ two.at({}).set(2.0f);
+
+ saw::data<sch::Scalar<T>> one;
+ one.at({}) = 1.0;
+
+ saw::data<sch::Scalar<T>> eps;
+ eps.at({}) = 1.5f;
+
saw::data<sch::Scalar<T>> sss;
- sss.at({}) = sub_steps.template cast_to<T>();
+ sss.at({}) = one.at({}) / sub_steps.template cast_to<T>();
auto vel_s = (pirb_pos-pirb_pos_old) / sss;
+ saw::data<sch::Vector<T,Descriptor::D>> force_p{};
+
+ iterator<Descriptor::D>::apply([&](const auto& index_f) -> void {
+ auto& force = force_f.at(index_f);
+ auto& vel = vel_f.at(index_f);
+ auto& por = por_f.at(index_f);
+ auto& rho = rho_f.at(index_f);
+
+ saw::data<sch::Vector<T,Descriptor::D>> rel_dist = saw::math::vectorize_data(index_f).template cast_to<T>() - pirb_pos;
+
+ por = particle_porosity<T,Descriptor::D,1u,por::ParticleSpheroid<T>>::calculate(rel_dist,p_rad,eps);
+ auto flip_por = one - por;
+
+ // vel_s is technically time the density of the particle?
+
+ force = ( vel_s * rho - vel * rho ) * two * flip_por;
+
+ force_p = force_p - force;
+ }, aabb.template get<"a">(), aabb.template get<"b">());
+
+ auto& pirb_acc = pirb.template get<"acceleration">();
+ pirb_acc = force_p / pg.template get<"total_mass">().at({});
+
+ for(saw::data<sch::UInt64> i{0u}; i < sub_steps; ++i){
+ verlet_step_lambda<T,Descriptor::D>(pi,sss);
+ }
+ }
+};
+
+template<typename T, typename Descriptor, typename Encode>
+class component<T, Descriptor, cmpt::FpLbmOneParticleImplicit, Encode> final {
+public:
+ using Component = cmpt::FpLbmOneParticle;
+private:
+public:
+ component() = default;
+
+ template<typename CellFieldSchema, typename MacroFieldSchema, typename ParticleSchema>
+ void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, const saw::data<ParticleSchema,Encode>& pg, saw::data<sch::FixedArray<sch::UInt64,1u>> index, saw::data<sch::UInt64> time_step, saw::data<sch::UInt64> sub_steps) const {
+ // void apply(saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64, Descriptor::D>> index, saw::data<sch::UInt64> time_step){
+ bool is_even = ((time_step.get() % 2) == 0);
+
+ //auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
+ // auto& dfs = dfs_old_f.at(index);
+
+ auto& rho_f = macros.template get<"density">();
+ auto& vel_f = macros.template get<"velocity">();
+ auto& por_f = macros.template get<"porosity">();
+ auto& force_f = macros.template get<"force">();
+
+ /**
+ * The other methods all work with a flipped porosity.
+ * For compat reasons this is also flipped
+ */
+ // TODO - Change to tuple later
+ auto parts = pg.template get<"particles">();
+
+ auto parts_size = parts.meta().at({0u});
+ auto& pi = parts.at(index);
+
+ auto& pirb = pi.template get<"rigid_body">();
+ auto& pirb_pos = pirb.template get<"position">();
+ auto& pirb_pos_old = pirb.template get<"position_old">();
+
+
auto& p_coll = pg.template get<"collision">().at({});
auto& p_rad = p_coll.template get<"radius">();
@@ -172,15 +259,19 @@ public:
auto& aabb = eo_aabb.get_value();
saw::data<sch::Scalar<T>> two;
- two.at({}).set(2);
+ two.at({}).set(2.0f);
saw::data<sch::Scalar<T>> one;
one.at({}) = 1.0;
saw::data<sch::Scalar<T>> eps;
eps.at({}) = 1.5f;
+
+ saw::data<sch::Scalar<T>> sss;
+ sss.at({}) = one.at({}) / sub_steps.template cast_to<T>();
+ auto vel_s = (pirb_pos-pirb_pos_old) / sss;
- saw::data<sch::Vector<T,Descriptor::D>> force_p;
+ saw::data<sch::Vector<T,Descriptor::D>> force_p{};
iterator<Descriptor::D>::apply([&](const auto& index_f) -> void {
auto& force = force_f.at(index_f);
@@ -193,13 +284,15 @@ public:
por = particle_porosity<T,Descriptor::D,1u,por::ParticleSpheroid<T>>::calculate(rel_dist,p_rad,eps);
auto flip_por = one - por;
- force = ( vel_s * rho - vel * rho ) * two * flip_por;
+ // vel_s is technically time the density of the particle?
+
+ force = ( vel_s * rho - vel * rho ) * two * flip_por / (one + flip_por);
- force_p = force_p + force;
+ force_p = force_p - force;
}, aabb.template get<"a">(), aabb.template get<"b">());
auto& pirb_acc = pirb.template get<"acceleration">();
- pirb_acc = - force_p / pg.template get<"total_mass">().at({});
+ pirb_acc = force_p / pg.template get<"total_mass">().at({});
for(saw::data<sch::UInt64> i{0u}; i < sub_steps; ++i){
verlet_step_lambda<T,Descriptor::D>(pi,sss);
diff --git a/modules/core/c++/particle/aabb.hpp b/modules/core/c++/particle/aabb.hpp
index da39ec2..e4c930b 100644
--- a/modules/core/c++/particle/aabb.hpp
+++ b/modules/core/c++/particle/aabb.hpp
@@ -26,7 +26,6 @@ public:
static constexpr saw::error_or<saw::data<AABB>> calculate(const saw::data<Sch,Encode>& p_grp, const saw::data<sch::FixedArray<sch::UInt64,1u>>& index, const saw::data<sch::FixedArray<sch::UInt64,D>>& meta){
static_assert(PC > 0u, "Can't calculate from no particles");
if(not (index.at({{0u}}).get() < PC) ){
- std::cerr.flush();
return saw::make_error<saw::err::critical>("Too large i in particle_aabb");
}
diff --git a/modules/core/c++/particle/cuboid/aabb.hpp b/modules/core/c++/particle/cuboid/aabb.hpp
new file mode 100644
index 0000000..502132a
--- /dev/null
+++ b/modules/core/c++/particle/cuboid/aabb.hpp
@@ -0,0 +1,40 @@
+#pragma once
+
+#include "common.hpp"
+#include "../aabb.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename T, uint64_t D, uint64_t PC>
+class particle_aabb<
+ sch::ParticleGroup<T,D,PC,coll::Cuboid<T>>
+> final {
+public:
+ using Schema = sch::ParticleGroup<T,D,PC,coll::Cuboid<T>>;
+
+ using AABB = sch::Struct<
+ sch::Member<sch::FixedArray<sch::UInt64,D>, "a">,
+ sch::Member<sch::FixedArray<sch::UInt64,D>, "b">
+ >;
+
+public:
+ template<typename Sch, typename Encode>
+ static constexpr saw::error_or<saw::data<AABB>> calculate(const saw::data<Sch,Encode>& pg, const saw::data<sch::FixedArray<sch::UInt64,1u>>& index, const saw::data<sch::FixedArray<sch::UInt64,D>>& meta){
+ static_assert(PC > 0u, "Can't calculate from no particle");
+ if(not (index.at({0u}).get() < PC)){
+ return saw::make_error<saw::err::critical>("Too large i in particle_aabb coll::Cuboid<T>");
+ }
+
+ saw::data<AABB> aabb;
+ auto& parts = pg.template get<"particles">();
+ auto& pi = parts.at(index);
+ auto& pirb = pi.template get<"rigid_body">();
+
+ /// TODO
+
+ return aabb;
+ }
+};
+}
+}
diff --git a/modules/core/c++/particle/cuboid/common.hpp b/modules/core/c++/particle/cuboid/common.hpp
new file mode 100644
index 0000000..a859cbe
--- /dev/null
+++ b/modules/core/c++/particle/cuboid/common.hpp
@@ -0,0 +1,10 @@
+#pragma once
+
+namespace kel {
+namespace lbm {
+namespace coll {
+template<typename T>
+struct Cuboid {};
+}
+}
+}
diff --git a/modules/core/c++/particle/cuboid/particle.hpp b/modules/core/c++/particle/cuboid/particle.hpp
new file mode 100644
index 0000000..e19d543
--- /dev/null
+++ b/modules/core/c++/particle/cuboid/particle.hpp
@@ -0,0 +1,9 @@
+#pragma once
+
+#include "common.hpp"
+#include "porosity.hpp"
+
+namespace kel {
+namespace lbm {
+}
+}
diff --git a/modules/core/c++/particle/cuboid/porosity.hpp b/modules/core/c++/particle/cuboid/porosity.hpp
new file mode 100644
index 0000000..bccfebf
--- /dev/null
+++ b/modules/core/c++/particle/cuboid/porosity.hpp
@@ -0,0 +1,10 @@
+#pragma once
+
+#include "common.hpp"
+#include "../porosity.hpp"
+
+namespace kel {
+namespace lbm {
+
+}
+}
diff --git a/modules/core/c++/particle/particle.hpp b/modules/core/c++/particle/particle.hpp
index a3669b4..21b4b34 100644
--- a/modules/core/c++/particle/particle.hpp
+++ b/modules/core/c++/particle/particle.hpp
@@ -82,8 +82,8 @@ saw::data<sch::ParticleGroup<T, D, PartAmount, coll::Spheroid<T>>> create_sphero
total_mass.at({}) = rad_d * rad_d * density.at({}) * 3.141592;
}else if constexpr ( D == 3u ){
- }else if constexpr ( D== 1u ){
-
+ }else if constexpr ( D == 1u ){
+ total_mass.at({}) = rad_d * 2.0;
}
std::cout<<"Total Mass: "<<total_mass.at({}).get()<<std::endl;
diff --git a/modules/core/c++/particle/spheroid/particle.hpp b/modules/core/c++/particle/spheroid/particle.hpp
new file mode 100644
index 0000000..9c0912e
--- /dev/null
+++ b/modules/core/c++/particle/spheroid/particle.hpp
@@ -0,0 +1,9 @@
+#pragma once
+
+namespace kel {
+namespace lbm {
+namespace coll {
+
+}
+}
+}
diff --git a/modules/core/c++/psm.hpp b/modules/core/c++/psm.hpp
index 02db1e1..15ad0f8 100644
--- a/modules/core/c++/psm.hpp
+++ b/modules/core/c++/psm.hpp
@@ -7,15 +7,30 @@
namespace kel {
namespace lbm {
namespace cmpt {
-struct PSM {};
+struct PsmReset {};
struct PsmOneParticle {};
+struct Psm {};
}
+template<typename T, typename Descriptor, typename Encode>
+class component<T, Descriptor, cmpt::PsmReset, Encode> final {
+public:
+ component() = default;
+
+ template<typename CellFieldSchema, typename MacroFieldSchema>
+ void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step) const {
+ auto& porosity_f = macros.template get<"porosity">();
+
+ auto& por = porosity_f.at(index);
+ por.at({}) = 1.0;
+ }
+};
+
/**
* PSM collision operator for LBM
*/
template<typename T, typename Descriptor, typename Encode>
-class component<T, Descriptor, cmpt::PSM, Encode> {
+class component<T, Descriptor, cmpt::Psm, Encode> {
private:
saw::data<T> relaxation_;
saw::data<T> frequency_;
@@ -90,17 +105,115 @@ public:
component() = default;
template<typename CellFieldSchema, typename MacroFieldSchema, typename ParticleSchema>
- void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, const saw::data<ParticleSchema,Encode>& particles, saw::data<sch::FixedArray<sch::UInt64,1u>> index, saw::data<sch::UInt64> time_step) const {
+ void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, const saw::data<ParticleSchema,Encode>& pg, saw::data<sch::FixedArray<sch::UInt64,1u>> index, saw::data<sch::UInt64> time_step, saw::data<sch::UInt64> sub_steps) const {
using dfi = df_info<T,Descriptor>;
bool is_even = ((time_step.get() % 2) == 0);
+
+ saw::data<sch::Scalar<T>> one;
+ one.at({}) = 1.0;
auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
- auto& porous_f = macros.template get<"porosity">();
+ auto& por_f = macros.template get<"porosity">();
auto& rho_f = macros.template get<"density">();
auto& vel_f = macros.template get<"velocity">();
-
+ auto& force_f = macros.template get<"force">();
+
+ {
+ auto parts = pg.template get<"particles">();
+ auto parts_size = parts.meta().at({0u});
+
+ auto& p_coll = pg.template get<"collision">().at({});
+ auto& p_rad = p_coll.template get<"radius">();
+
+ auto& pi = parts.at(index);
+ auto& pirb = pi.template get<"rigid_body">();
+ saw::data<sch::Vector<T,Descriptor::D>>& pirb_pos = pirb.template get<"position">();
+ auto& pirb_pos_old = pirb.template get<"position_old">();
+
+ saw::data<sch::Scalar<T>> ts;
+ ts.at({}) = one.at({}) / sub_steps.template cast_to<T>();
+
+ saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> start;
+ saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> stop;
+
+ auto eo_aabb = particle_aabb<typename ParticleSchema::ValueType>::calculate(pg,{{0u}},vel_f.meta());
+ if(eo_aabb.is_error()){
+ return;
+ }
+ auto& aabb = eo_aabb.get_value();
+
+ /// Ok, I iterate over the space which covers our particle? So lower bounds to upper bounds
+ start = aabb.template get<"a">();
+ stop = aabb.template get<"b">();
+
+ saw::data<sch::Vector<T,Descriptor::D>> force_p;
+ for(uint64_t i{0u}; i < Descriptor::D; ++i){
+ force_p.at({{i}}) = 0.0;
+ }
+ auto vel_p_old = (pirb_pos-pirb_pos_old) / ts;
+
+ iterator<Descriptor::D>::apply([&](const auto& index_f) -> void{
+ // ask for the d_k value here.
+ // For every value im iterating over I need sth
+ // std::cout<<"Pos: "<<index.at({0u}).get()<<" "<<index.at({1u}).get()<<std::endl;
+
+ auto& dfs = dfs_old_f.at(index_f);
+ saw::data<sch::Vector<T,Descriptor::D>> rel_dist = saw::math::vectorize_data(index_f).template cast_to<T>() - pirb_pos;
+ saw::data<sch::Scalar<T>> eps;
+ eps.at({}) = 1.5f;
+
+ auto& por = por_f.at(index_f);
+ por = particle_porosity<T,Descriptor::D,1u,por::ParticleSpheroid<T>>::calculate(rel_dist,p_rad,eps);
+
+ if(por.at({}).get() >= 1.0f){
+ return;
+ }
+
+ saw::data<sch::Vector<T,Descriptor::D>> momentum;
+ for(uint64_t i{0u}; i < Descriptor::D; ++i){
+ momentum.at({{i}}) = 0.0;
+ }
+
+ for(uint64_t i{0u}; i < Descriptor::Q; ++i){
+ saw::data<sch::Vector<T,Descriptor::D>> e_i;
+ saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> n_ind_i;
+ for(uint64_t k{0u}; k < Descriptor::D; ++k){
+ e_i.at({{k}}) = (dfi::directions[i])[k];
+ n_ind_i.at({k}) = index_f.at({k}) + (dfi::directions[i])[k];
+ }
+
+ uint64_t i_opp = dfi::opposite_index[i];
+
+ auto u_p_e = saw::math::dot(e_i,vel_p_old);
+
+ saw::data<T> dfs_added = dfs.at({i})*(saw::data<T>{1}-u_p_e.at({})) + dfs_old_f.at(n_ind_i).at({i_opp})*(saw::data<T>{1}+u_p_e.at({}));
+ saw::data<sch::Scalar<T>> dfs_added_v;
+ dfs_added_v.at({}) = dfs_added;
+ auto ei_dfs = e_i * dfs_added_v;
+
+ momentum = momentum + ei_dfs;
+ }
+ // technically needs to adjust for rotation as well
+
+ auto& force = force_f.at(index_f);
+ auto& rho = rho_f.at(index_f);
+ // To Fluid
+
+ auto flip_por = one - por;
+ force = momentum * flip_por;
+ // To Particle
+ force_p = force_p - force;
+ },start,stop);
+
+ auto& pirb_acc = pirb.template get<"acceleration">();
+ pirb_acc = force_p / pg.template get<"total_mass">().at({});
+
+ for(saw::data<sch::UInt64> i{0u}; i < sub_steps; ++i){
+ verlet_step_lambda<T,Descriptor::D>(pi,ts);
+ }
+ }
}
};