diff options
Diffstat (limited to 'examples')
| -rw-r--r-- | examples/poiseulle_particles_2d_bgk_gpu/sim.cpp | 2 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp | 49 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp | 47 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 17 | ||||
| -rw-r--r-- | examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp | 60 | ||||
| -rw-r--r-- | examples/stokes_drag_particle_2d_psm_gpu/sim.cpp | 39 |
6 files changed, 92 insertions, 122 deletions
diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp index 73964b7..ef4639e 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp @@ -171,7 +171,7 @@ saw::error_or<void> step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here - component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.65}; + component<T,Desc,cmpt::BGK, 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; diff --git a/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp index 93d4b46..3d16884 100644 --- a/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp @@ -51,6 +51,7 @@ template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, Member<RhoChunk<T>, "density">, + Member<VelChunk<T,Desc>, "force">, Member<ScalarChunk<T,Desc>, "porosity"> >; @@ -144,6 +145,7 @@ saw::error_or<void> setup_initial_conditions( {{1u,1u}} ); + /* iterator<Desc::D>::apply( [&](auto& index){ saw::data<sch::Vector<T,Desc::D>> middle, ind_vec; @@ -162,6 +164,7 @@ saw::error_or<void> setup_initial_conditions( {},// 0-index df_f.get_dims() ); + */ return saw::make_void(); } @@ -179,16 +182,20 @@ saw::error_or<void> step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - component<T,Desc,cmpt::FpLbmOneParticleNoVelocity,encode::Sycl<saw::encode::Native>> collision{0.65}; + component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> bgk{0.8}; + component<T,Desc,cmpt::FpLbmOneParticleNoVelocity,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}; + saw::data<sch::Scalar<T>> eps; + eps.at({}).set(1.5f); + saw::data<sch::Scalar<T>> rad; + rad.at({}).set(dim_y*0.1); + saw::data<sch::Vector<T,Desc::D>> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.25; + p_pos.at({{1u}}) = dim_y * 0.5; + } + component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps}; component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ @@ -217,17 +224,18 @@ saw::error_or<void> step( bb.apply(fields,index,t_i); break; case 2u: + opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); break; case 3u: flow_in.apply(fields,index,t_i); // equi.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); // equi.apply(fields,index,t_i); - collision.apply(fields,macros,index,t_i); + bgk.apply(fields,macros,index,t_i); break; default: break; @@ -349,20 +357,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } 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_csv_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } + // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; @@ -402,6 +397,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){ 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(); diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp index b09aa64..b1bd37d 100644 --- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -48,6 +48,7 @@ template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, Member<RhoChunk<T>, "density">, + Member<VelChunk<T,Desc>, "force">, Member<ScalarChunk<T,Desc>, "porosity"> >; @@ -140,7 +141,8 @@ saw::error_or<void> setup_initial_conditions( df_f.get_dims(), {{1u,1u}} ); - + + /* iterator<Desc::D>::apply( [&](auto& index){ saw::data<sch::Vector<T,Desc::D>> middle, ind_vec; @@ -159,6 +161,7 @@ saw::error_or<void> setup_initial_conditions( {},// 0-index df_f.get_dims() ); + */ return saw::make_void(); } @@ -176,16 +179,23 @@ saw::error_or<void> step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.65}; + component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> bgk{0.8}; + 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}; + saw::data<sch::Scalar<T>> eps; + eps.at({}).set(1.5f); + saw::data<sch::Scalar<T>> rad; + rad.at({}).set(dim_y*0.1); + saw::data<sch::Vector<T,Desc::D>> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.25; + p_pos.at({{1u}}) = dim_y * 0.5; + } + component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps}; + + component<T,Desc,cmpt::ForceGather, encode::Sycl<saw::encode::Native>> fg; component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ @@ -214,32 +224,19 @@ saw::error_or<void> step( bb.apply(fields,index,t_i); break; case 2u: + opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); - { - saw::data<sch::Vector<T,Desc::D>> middle; - middle.at({{0u}}) = dim_x * 0.25; - middle.at({{1u}}) = dim_y * 0.5; - - saw::data<sch::Vector<T,Desc::D>> ind_vec; - ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to<T>(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>(); - - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - porous_f.at(index).at({}) = 0.0; - } - } + fg.apply(fields,macros,index,t_i); break; case 3u: flow_in.apply(fields,index,t_i); // equi.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); // equi.apply(fields,index,t_i); - collision.apply(fields,macros,index,t_i); + bgk.apply(fields,macros,index,t_i); break; default: break; diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index b3ea7ae..7d3c6f5 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -49,6 +49,7 @@ template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, Member<RhoChunk<T>, "density">, + Member<VelChunk<T,Desc>, "force">, Member<ScalarChunk<T,Desc>, "porosity"> >; @@ -216,17 +217,9 @@ saw::error_or<void> step( auto& porous_f = macros.template get<"porosity">(); 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::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}; - saw::data<sch::Scalar<T>> eps; eps.at({}).set(1.5f); @@ -250,6 +243,7 @@ saw::error_or<void> step( }; component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; + component<T,Desc,cmpt::ForceGather, encode::Sycl<saw::encode::Native>> fg; 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; @@ -268,16 +262,17 @@ saw::error_or<void> step( case 2u: opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); + fg.apply(fields,macros,index,t_i); break; case 3u: flow_in.apply(fields,index,t_i); //equi.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); // equi.apply(fields,index,t_i); - collision.apply(fields,macros,index,t_i); + bgk.apply(fields,macros,index,t_i); break; case 5u: // Corners diff --git a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp index a8ba21f..81498e2 100644 --- a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp +++ b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp @@ -13,7 +13,7 @@ namespace lbm { constexpr uint64_t dim_y = 1024ul; constexpr uint64_t dim_x = dim_y * 2ul; -constexpr uint64_t particle_amount = 1ul; +// constexpr uint64_t particle_amount = 1ul; namespace sch { using namespace saw::schema; @@ -47,6 +47,7 @@ using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>; template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, + Member<VelChunk<T,Desc>, "force">, Member<RhoChunk<T>, "density">, Member<ScalarChunk<T,Desc>, "porosity"> >; @@ -158,25 +159,6 @@ saw::error_or<void> setup_initial_conditions( {{1u,1u}} ); - iterator<Desc::D>::apply( - [&](auto& index){ - saw::data<sch::Vector<T,Desc::D>> middle, ind_vec; - middle.at({{0u}}) = dim_x * 0.5; - middle.at({{1u}}) = dim_y * 0.5; - - ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to<T>(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>(); - - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - porous_f.at(index).at({}) = 0.0; - } - }, - {},// 0-index - df_f.get_dims() - ); - return saw::make_void(); } @@ -196,12 +178,16 @@ saw::error_or<void> step( 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}; + saw::data<sch::Scalar<T>> eps; + eps.at({}).set(1.5f); + saw::data<sch::Scalar<T>> rad; + rad.at({}).set(dim_y*0.1*0.25); + saw::data<sch::Vector<T,Desc::D>> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.5; + p_pos.at({{1u}}) = dim_y * 0.5; + } + component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps}; component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ @@ -230,6 +216,7 @@ saw::error_or<void> step( abb.apply(fields,index,t_i); break; case 2u: + opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); break; case 3u: @@ -385,20 +372,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } */ - 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_csv_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } + // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; @@ -438,6 +412,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){ 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(); diff --git a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp index 0ff6e6b..6942cc7 100644 --- a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp +++ b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp @@ -48,6 +48,7 @@ template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, Member<RhoChunk<T>, "density">, + Member<VelChunk<T,Desc>, "force">, Member<ScalarChunk<T,Desc>, "porosity"> >; @@ -196,12 +197,16 @@ saw::error_or<void> step( 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}; + saw::data<sch::Scalar<T>> eps; + eps.at({}).set(1.5f); + saw::data<sch::Scalar<T>> rad; + rad.at({}).set(dim_y*0.1*0.25); + saw::data<sch::Vector<T,Desc::D>> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.5; + p_pos.at({{1u}}) = dim_y * 0.5; + } + component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps}; component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ @@ -230,6 +235,7 @@ saw::error_or<void> step( abb.apply(fields,index,t_i); break; case 2u: + opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); break; case 3u: @@ -385,20 +391,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } */ - 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_csv_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; @@ -438,7 +430,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){ 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(); } |
