diff options
Diffstat (limited to 'examples/poiseulle_particles_2d_gpu')
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/init.hpp | 6 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 8 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/step.hpp | 55 |
3 files changed, 34 insertions, 35 deletions
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 a4f08f0..fd6cdca 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -81,7 +81,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){ 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; } @@ -115,21 +115,21 @@ 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{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()){ diff --git a/examples/poiseulle_particles_2d_gpu/step.hpp b/examples/poiseulle_particles_2d_gpu/step.hpp index c1b148f..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,28 +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::HlbmOneParticle,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; - 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}; @@ -62,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: @@ -76,21 +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(); } |