diff options
Diffstat (limited to 'examples/poiseulle_particles_2d_gpu')
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/common.hpp | 2 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/init.hpp | 6 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 41 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/step.hpp | 54 |
4 files changed, 55 insertions, 48 deletions
diff --git a/examples/poiseulle_particles_2d_gpu/common.hpp b/examples/poiseulle_particles_2d_gpu/common.hpp index 6c05b64..90b7687 100644 --- a/examples/poiseulle_particles_2d_gpu/common.hpp +++ b/examples/poiseulle_particles_2d_gpu/common.hpp @@ -55,7 +55,7 @@ using MacroStruct = Struct< >; template<typename T, typename Desc> -using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,coll::Spheroid<T>>; +using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,1u,coll::Spheroid<T>>; } } diff --git a/examples/poiseulle_particles_2d_gpu/init.hpp b/examples/poiseulle_particles_2d_gpu/init.hpp index 617b296..cdce25e 100644 --- a/examples/poiseulle_particles_2d_gpu/init.hpp +++ b/examples/poiseulle_particles_2d_gpu/init.hpp @@ -5,8 +5,8 @@ namespace kel { namespace lbm { -template<typename T, typename Desc> -saw::error_or<void> setup_initial_conditions( +template<typename T, typename Desc, typename Coll> +saw::error_or<void> init( const converter<T>& conv, saw::data<sch::ChunkStruct<T,Desc>>& fields, saw::data<sch::MacroStruct<T,Desc>>& macros, @@ -117,7 +117,7 @@ saw::error_or<void> setup_initial_conditions( radius_p.at({}).set(2); saw::data<sch::Scalar<T>> dense_p; dense_p.at({}).set(1); - particles = create_spheroid_particle_group<T,Desc::D>(radius_p, dense_p, {{16u}}); + particles = create_spheroid_particle_group<T,Desc::D,1u>(radius_p, dense_p, {{16u}}); } return saw::make_void(); diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 865babf..fd6cdca 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -22,18 +22,11 @@ using LbmArgs = Args< } template<typename T, typename Desc, typename Coll> -saw::error_or<void> lbm_main(int argc, char** argv){ +saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){ using namespace kel::lbm; using dfi = df_info<T,Desc>; - auto eo_args = saw::parse_args<args::LbmArgs>(argc,argv); - if(eo_args.is_error()){ - return std::move(eo_args.get_error()); - } - auto& args = eo_args.get_value(); - (void)args; - auto eo_lbm_dir = output_directory(); if(eo_lbm_dir.is_error()){ return std::move(eo_lbm_dir.get_error()); @@ -88,7 +81,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ sycl_q.wait(); { - auto eov = setup_initial_conditions<T,Desc>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); + auto eov = init<T,Desc,Coll>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); if(eov.is_error()){ return eov; } @@ -118,26 +111,25 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } 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_particle_data); - saw::data<sch::UInt64> time_steps{16u*4096ul}; + saw::data<sch::UInt64> time_steps{64u*1024ul}; 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>(conv,lsd_view,lsdm_view,lsdp_view,i,dev); + auto eov = step<T,Desc,Coll>(conv,lsd_view,lsdm_view,lsdp_view,i,dev); if(eov.is_error()){ return eov; } } sycl_q.wait(); - if(i.get() % 32u == 0u){ + if(i.get() % 64u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); if(eov.is_error()){ @@ -195,17 +187,32 @@ saw::error_or<void> lbm_main(int argc, char** argv){ return saw::make_void(); } - using FloatT = kel::lbm::sch::Float32; +using DescT = kel::lbm::sch::D2Q9; -saw::error_or<void> kmain(int argc, char** argv){ +saw::error_or<void> k_main(int argc, char** argv){ using namespace kel::lbm; + + saw::data<args::LbmArgs> args; + { + auto eov = saw::parse_args(args,argc,argv); + return eov; + } + auto& an = args.template get<"args">(); + auto& coupling = an.template get<"coupling">(); - // auto eo_args = saw::parse_args< + if(coupling.stl_view() == "hlbm"){ + return lbm_main<FloatT,DescT,method::Hlbm>(args); + //}else if(coupling.view() == "psm"){ + // return kmain<FloatT,DescT,coupling::Psm>(args); + } + + return saw::make_error<saw::err::critical>("Invalid coupling"); } + int main(int argc, char** argv){ - auto eov = lbm_main<FloatT,kel::lbm::sch::D2Q9, >(argc, argv); + auto eov = k_main(argc, argv); if(eov.is_error()){ auto& err = eov.get_error(); std::cerr<<"[Error] "<<err.get_category(); diff --git a/examples/poiseulle_particles_2d_gpu/step.hpp b/examples/poiseulle_particles_2d_gpu/step.hpp index aa0e382..37eb4e7 100644 --- a/examples/poiseulle_particles_2d_gpu/step.hpp +++ b/examples/poiseulle_particles_2d_gpu/step.hpp @@ -5,7 +5,7 @@ namespace kel { namespace lbm { -template<typename T, typename Desc> +template<typename T, typename Desc, typename Coll> saw::error_or<void> step( const converter<T>& conv, saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields, @@ -17,29 +17,44 @@ saw::error_or<void> step( 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::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){ + saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index; + for(uint64_t i = 0u; i < Desc::D; ++i){ + index.at({{i}}).set(idx[i]); + } + + hlbm_reset.apply(fields,macros,index,t_i); + }); + }).wait(); + + q.submit([&](acpp::sycl::handler& h){ + component<T,Desc,cmpt::HlbmOneParticle,encode::Sycl<saw::encode::Native>> hlbm_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]); + } - component<T,Desc,cmpt::HlbmParticle,encode::Sycl<saw::encode::Native>> particle; + hlbm_one_part.apply(fields,macros,particles,index,t_i,{16u}); + }); + }).wait(); // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.8}; component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb; - component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb; - - saw::data<sch::Scalar<T>> rho_b; - rho_b.at({}) = 1.0; - saw::data<sch::Vector<T,Desc::D>> vel_b; - vel_b.at({{0u}}) = 0.015; - - component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b}; component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ uint64_t target_t_i = 16u; if(t_i.get() < target_t_i){ - return 1.0 + (0.0002 / target_t_i) * t_i.get(); + return 1.0 + (0.01 / target_t_i) * t_i.get(); } - return 1.0002; + return 1.01; }() }; component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; @@ -63,12 +78,10 @@ saw::error_or<void> step( break; case 3u: flow_in.apply(fields,index,t_i); - // equi.apply(fields,index,t_i); collision.apply(fields,macros,index,t_i); break; case 4u: flow_out.apply(fields,index,t_i); - // equi.apply(fields,index,t_i); collision.apply(fields,macros,index,t_i); break; default: @@ -77,19 +90,6 @@ saw::error_or<void> step( }); }).wait(); - q.submit([&](acpp::sycl::handler& h){ - h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ - particle.apply(fields,macros,particles,{{0u}},t_i); - }); - }).wait(); - - // Step - /* - q.submit([&](acpp::sycl::handler& h){ - // h.depends_on(collision_ev); - }).wait(); - */ - return saw::make_void(); } |