summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--examples/one_particle_sedimentation_2d/sim.cpp3
-rw-r--r--examples/schaefer_turek_durst_krause_rannbacher/common.hpp13
-rw-r--r--examples/schaefer_turek_durst_krause_rannbacher/init.hpp6
-rw-r--r--examples/schaefer_turek_durst_krause_rannbacher/sim.cpp16
-rw-r--r--examples/schaefer_turek_durst_krause_rannbacher/step.hpp12
-rw-r--r--examples/throat_clog_2d/init.hpp34
-rw-r--r--modules/core/c++/fplbm.hpp157
7 files changed, 66 insertions, 175 deletions
diff --git a/examples/one_particle_sedimentation_2d/sim.cpp b/examples/one_particle_sedimentation_2d/sim.cpp
index 86a7b1c..db730d4 100644
--- a/examples/one_particle_sedimentation_2d/sim.cpp
+++ b/examples/one_particle_sedimentation_2d/sim.cpp
@@ -126,7 +126,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
auto lsdm_view = make_view(lbm_sycl_macro_data);
auto lsdp_view = make_view(lbm_sycl_particle_data);
- saw::data<sch::UInt64> time_steps{1024ul * 1024ul};
+ saw::data<sch::UInt64> time_steps{1024ul * 1024ul * 64ul};
auto& info_f = lsd_view.template get<"info">();
@@ -220,6 +220,7 @@ saw::error_or<void> k_main(int argc, char** argv){
}
}
auto& an = args.template get<"args">();
+
auto& coupling = an.template get<"coupling">();
if(coupling.stl_view() == "hlbm"){
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/common.hpp b/examples/schaefer_turek_durst_krause_rannbacher/common.hpp
index 59d2f9c..ac5c94e 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/common.hpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/common.hpp
@@ -25,6 +25,19 @@ namespace lbm {
constexpr uint64_t dim_y = 41u;
constexpr uint64_t dim_x = 220u;
+template<uint64_t N>
+struct dimensions {
+ static constexpr std::array<uint64_t,2u> calculate(){
+ return {dim_x * N + 1u, dim_y * N + 1u};
+ }
+};
+
+template<uint64_t N>
+constexpr std::array<uint64_t,2u> dims{
+ dim_x * N + 1u,
+ dim_y * N + 1u
+};
+
constexpr uint64_t particle_amount = 1ul;
namespace sch {
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/init.hpp b/examples/schaefer_turek_durst_krause_rannbacher/init.hpp
index 96e2917..45cbaa2 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/init.hpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/init.hpp
@@ -41,7 +41,7 @@ saw::error_or<void> init(
info_f.at(index).set(3u);
},
{{0u,0u}},
- {{1u,dim_y}},
+ {{1u,dim_y*N}},
{{0u,1u}}
);
@@ -50,8 +50,8 @@ saw::error_or<void> init(
[&](auto& index){
info_f.at(index).set(4u);
},
- {{dim_x-1u,0u}},
- {{dim_x, dim_y}},
+ {{dim_x*N-1u,0u}},
+ {{dim_x*N, dim_y*N}},
{{0u,1u}}
);
//
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
index 35c04a4..71410b2 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
@@ -34,7 +34,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
}
auto& lbm_dir = eo_lbm_dir.get_value();
- auto out_dir = lbm_dir / "stdkr_2d_gpu" / an.template get<"coupling">().stl_view();
+ auto out_dir = lbm_dir / "stdkr_2d" / std::to_string(N) / an.template get<"coupling">().stl_view();
{
std::error_code ec;
@@ -118,7 +118,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
auto lsdm_view = make_view(lbm_sycl_macro_data);
auto lsdp_view = make_view(lbm_sycl_particle_data);
- saw::data<sch::UInt64> time_steps{64u*1024ul};
+ saw::data<sch::UInt64> time_steps{64u*1024ul * N * N};
auto& info_f = lsd_view.template get<"info">();
@@ -140,7 +140,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
}
}
{
- auto eov = write_vtk_file(out_dir,std::string{"ops_2d_"}+an.template get<"coupling">().stl_string(),i.get(), *lbm_macro_data_ptr);
+ auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+std::to_string(N)+std::string{"_"}+an.template get<"coupling">().stl_string(),i.get(), *lbm_macro_data_ptr);
if(eov.is_error()){
return eov;
}
@@ -150,7 +150,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
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){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -180,7 +180,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
}
}
{
- auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+an.template get<"coupling">().stl_string(),time_steps.get(), *lbm_macro_data_ptr);
+ auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+std::to_string(N)+std::string{"_"}+an.template get<"coupling">().stl_string(),time_steps.get(), *lbm_macro_data_ptr);
if(eov.is_error()){
return eov;
}
@@ -212,11 +212,11 @@ saw::error_or<void> k_main(int argc, char** argv){
auto& coupling = an.template get<"coupling">();
if(coupling.stl_view() == "hlbm"){
- return lbm_main<FloatT,DescT,method::Hlbm,2u>(args);
+ return lbm_main<FloatT,DescT,method::Hlbm,1u>(args);
} else if(coupling.stl_view() == "fplbm"){
- return lbm_main<FloatT,DescT,method::FpLbm,2u>(args);
+ return lbm_main<FloatT,DescT,method::FpLbm,1u>(args);
} else if(coupling.stl_view() == "psm"){
- return lbm_main<FloatT,DescT,method::Psm,2u>(args);
+ return lbm_main<FloatT,DescT,method::Psm,1u>(args);
}
return saw::make_error<saw::err::critical>("Invalid coupling");
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
index e88f262..3e5f9ba 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
@@ -23,7 +23,7 @@ saw::error_or<void> step(
if constexpr ( std::is_same_v<Coll, method::Hlbm> ){
q.submit([&](acpp::sycl::handler& h){
component<T,Desc,cmpt::HlbmReset,encode::Sycl<saw::encode::Native>> hlbm_reset;
- h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -50,7 +50,7 @@ saw::error_or<void> step(
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
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){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -96,7 +96,7 @@ saw::error_or<void> step(
}else if constexpr ( std::is_same_v<Coll, method::FpLbm> ){
q.submit([&](acpp::sycl::handler& h){
component<T,Desc,cmpt::FpLbmReset,encode::Sycl<saw::encode::Native>> fplbm_reset;
- h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -126,7 +126,7 @@ saw::error_or<void> step(
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
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){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -172,7 +172,7 @@ saw::error_or<void> step(
}else if constexpr ( std::is_same_v<Coll, method::Psm> ){
q.submit([&](acpp::sycl::handler& h){
component<T,Desc,cmpt::PsmReset,encode::Sycl<saw::encode::Native>> fplbm_reset;
- h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
@@ -202,7 +202,7 @@ saw::error_or<void> step(
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
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){
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x*N,dim_y*N}, [=](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]);
diff --git a/examples/throat_clog_2d/init.hpp b/examples/throat_clog_2d/init.hpp
index 69225be..06ff179 100644
--- a/examples/throat_clog_2d/init.hpp
+++ b/examples/throat_clog_2d/init.hpp
@@ -35,19 +35,33 @@ saw::error_or<void> init(
{{1u,1u}}
);
+ auto dims = info_f.get_dims();
+ auto width = dims.at({{0u}});
+ auto height = dims.at({{1u}});
+ //
+ auto w = width.template cast_to<T>();
+ auto h = height.template cast_to<T>();
+ auto w2 = w * 0.5;
+ auto h4 = h * 0.25;
+
+ saw::data<T> e{0.1};
+ saw::data<T> c = saw::data<T>{1.0} - e;
+
+ saw::data<T> d{0.2};
+ auto b = (saw::data<T>{1.0} - d) * 0.5;
+
+ auto p_r = (c*c + b*b) / (b*2);
+ auto p_i = p_r - b;
+
+ // Circle Top
+ saw::data<sch::FixedArray<sch::Int64,2u>> top_pos{w2.get(),h + p_i};
+
+ // Circle Bot
+ saw::data<sch::FixedArray<sch::Int64,2u>> bot_pos{w2.get(),-p_i.get()};
+
iterator<Desc::D>::apply(
[&](auto& index){
- auto dims = info_f.get_dims();
- auto width = dims.at({{0u}});
- auto height = dims.at({{1u}});
- //
- auto w2 = width / 2;
- auto h4 = height / 4;
- // Circle Top
- saw::data<sch::FixedArray<sch::Int64,2u>> top_pos{{w2.get(),(height+h4).get()}};
- // Circle Bot
- saw::data<sch::FixedArray<sch::Int64,2u>> bot_pos{{w2.get(),-h4.get()}};
},
{},
diff --git a/modules/core/c++/fplbm.hpp b/modules/core/c++/fplbm.hpp
index c48e127..aefc568 100644
--- a/modules/core/c++/fplbm.hpp
+++ b/modules/core/c++/fplbm.hpp
@@ -94,13 +94,14 @@ public:
auto& force_f = macros.template get<"force">();
auto& force = force_f.at(index);
- auto& vel_f = macros.template get<"momentum">();
- saw::data<sch::Vector<T,Descriptor::D>> vel = vel_f.at(index);
+ auto& mom_f = macros.template get<"momentum">();
+ auto& mom = mom_f.at(index);
// We already computed this in reset
// compute_rho_u<T,Descriptor>(dfs,rho,vel);
- vel = vel + force * half;
- auto eq = equilibrium<T,Descriptor>(rho,vel / rho);
+ mom = mom + force * half;
+ auto vel = mom / rho;
+ auto eq = equilibrium<T,Descriptor>(rho,vel);
using dfi = df_info<T,Descriptor>;
@@ -115,12 +116,12 @@ public:
for(uint64_t d{0u}; d < Descriptor::D; ++d){
ci.at({{d}}).set(static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d]));
}
- auto ci_dot_u = saw::math::dot(ci,vel/rho);
+ auto ci_dot_u = saw::math::dot(ci,vel);
saw::data<sch::Scalar<T>> w;
w.at({}).set(dfi::weights[i]);
- auto term1 = (ci-vel/rho) * dfi_inv_cs2;
+ auto term1 = (ci-vel) * dfi_inv_cs2;
auto term2 = ci * (ci_dot_u * dfi_inv_cs2 * dfi_inv_cs2);
auto force_projection = saw::math::dot(term1 + term2, force);
@@ -137,145 +138,6 @@ public:
};
template<typename T, typename Descriptor, typename Encode>
-class component<T, Descriptor, cmpt::FpLbmMomentumOneParticle, Encode> final {
-public:
- using Component = cmpt::FpLbmOneParticle;
-private:
- saw::data<sch::Vector<T,Descriptor::D>> global_force_;
-public:
- component():
- global_force_{}
- {}
- component(saw::data<sch::Vector<T,Descriptor::D>> global_force__):
- global_force_{global_force__}
- {}
-
- 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>& part_group, saw::data<sch::FixedArray<sch::UInt64,1u>> index, saw::data<sch::UInt64> time_step, saw::data<sch::UInt64> sub_steps) const {
- /// Figure out how to access the particle list
- // auto& p = particles.at(i);
-
- /// Iterate over the grid bounds
- // auto& grid = p.template get<"grid">();
- 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& part_spheroid_group = part_group;
- auto& mvel_f = macros.template get<"momentum">();
- auto& mrho_f = macros.template get<"density">();
- auto& mpor_f = macros.template get<"porosity">();
-
- auto& particle_N_f = field.template get<"particle_N">();
- auto& particle_D_f = field.template get<"particle_D">();
- {
- auto parts = part_spheroid_group.template get<"particles">();
- auto parts_size = parts.meta().at({0u});
-
- auto& p_coll = part_spheroid_group.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">();
-
- // TODO !!!! Divide by actual time step - for now it's ok
- 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(part_spheroid_group,{{0u}},mvel_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{
- 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& mpor = mpor_f.at(index_f);
- mpor = particle_porosity<T,Descriptor::D,1u,por::ParticleSpheroid<T>>::calculate(rel_dist,p_rad,eps);
-
- if(mpor.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;
- }
-
- // TODO technically needs to adjust for rotation as well
-
- auto& rho = mrho_f.at(index_f);
-
- macros.template get<"force">().at(index_f) = momentum;
- force_p = force_p + momentum;
- },start,stop);
-
- auto& pirb_acc = pirb.template get<"acceleration">();
- pirb_acc = - force_p / part_spheroid_group.template get<"total_mass">().at({});
-
- for(saw::data<sch::UInt64> i{0u}; i < sub_steps; ++i){
- verlet_step_lambda<T,Descriptor::D>(pi,ts);
- }
- auto vel_p = (pirb_pos-pirb_pos_old) / ts;
-
- iterator<Descriptor::D>::apply([&](const auto& index_f){
- auto& N = particle_N_f.at(index_f);
- auto& D = particle_D_f.at(index_f);
- auto& mpor = mpor_f.at(index_f);
- auto flip_porosity = one - mpor;
-
- D = D + flip_porosity;
- N = N + vel_p * flip_porosity;
- },start,stop);
- // Check
- }
- }
-};
-
-template<typename T, typename Descriptor, typename Encode>
class component<T, Descriptor, cmpt::FpLbmOneParticle, Encode> final {
public:
using Component = cmpt::FpLbmOneParticle;
@@ -342,7 +204,7 @@ public:
sss.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{{}};
//= global_force_;
iterator<Descriptor::D>::apply([&](const auto& index_f) -> void {
@@ -358,7 +220,8 @@ public:
// vel_s is technically time the density of the particle?
- force = global_force_ * por + ( vel_s - mom) * two * flip_por;
+ // force = global_force_ * por + ( vel_s - mom) * two * flip_por;
+ force = ( vel_s - mom) * two * flip_por;
// force = (vel_s - mom) * flip_por / (flip_por / two + one);
force_p = force_p - force;