diff options
| -rw-r--r-- | examples/one_particle_sedimentation_2d/sim.cpp | 3 | ||||
| -rw-r--r-- | examples/schaefer_turek_durst_krause_rannbacher/common.hpp | 13 | ||||
| -rw-r--r-- | examples/schaefer_turek_durst_krause_rannbacher/init.hpp | 6 | ||||
| -rw-r--r-- | examples/schaefer_turek_durst_krause_rannbacher/sim.cpp | 16 | ||||
| -rw-r--r-- | examples/schaefer_turek_durst_krause_rannbacher/step.hpp | 12 | ||||
| -rw-r--r-- | examples/throat_clog_2d/init.hpp | 34 | ||||
| -rw-r--r-- | modules/core/c++/fplbm.hpp | 157 |
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; |
