diff options
| -rw-r--r-- | examples/poiseulle_particles_2d_bgk_gpu/sim.cpp | 5 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp | 34 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 7 | ||||
| -rw-r--r-- | lib/core/c++/boundary.hpp | 14 | ||||
| -rw-r--r-- | lib/core/c++/grid.hpp | 2 |
5 files changed, 27 insertions, 35 deletions
diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp index f5d3f72..6bf4f84 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp @@ -181,9 +181,8 @@ saw::error_or<void> step( 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::HLBM,encode::Sycl<saw::encode::Native>> collision{0.65}; - // component<T,Desc,cmpt::PSM,encode::Sycl<saw::encode::Native>> collision{0.65}; 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; @@ -216,7 +215,7 @@ saw::error_or<void> step( case 0u: break; case 1u: - bb.apply(fields,index,t_i); + abb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp index 1992ede..a5ef912 100644 --- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -200,14 +200,11 @@ saw::error_or<void> step( auto& p_rad = p_coll.template get<"radius">(); } - // 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::HLBM,encode::Sycl<saw::encode::Native>> collision{0.65}; - // component<T,Desc,cmpt::PSM,encode::Sycl<saw::encode::Native>> collision{0.65}; + component<T,Desc,cmpt::HLBM,encode::Sycl<saw::encode::Native>> collision{0.65}; 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; @@ -240,7 +237,7 @@ saw::error_or<void> step( case 0u: break; case 1u: - bb.apply(fields,index,t_i); + abb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); @@ -286,7 +283,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } auto& lbm_dir = eo_lbm_dir.get_value(); - auto out_dir = lbm_dir / "poiseulle_particles_2d_gpu"; + auto out_dir = lbm_dir / "poiseulle_particles_2d_hlbm_gpu"; { std::error_code ec; @@ -372,8 +369,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } sycl_q.wait(); - saw::data<sch::UInt64> time_steps{4096ul}; - + saw::data<sch::UInt64> time_steps{16u*4096ul}; auto& info_f = lsd_view.template get<"info">(); for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){ @@ -385,20 +381,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } sycl_q.wait(); - { - { - auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = write_vtk_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } /* { { @@ -443,6 +425,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){ // After Loop sycl_q.wait(); { + auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } + { auto eov = write_vtk_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr); if(eov.is_error()){ return eov; diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index ff636ca..a8aeabb 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -201,13 +201,10 @@ 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::HLBM,encode::Sycl<saw::encode::Native>> collision{0.65}; component<T,Desc,cmpt::PSM,encode::Sycl<saw::encode::Native>> collision{0.65}; 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; @@ -240,7 +237,7 @@ saw::error_or<void> step( case 0u: break; case 1u: - bb.apply(fields,index,t_i); + abb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); diff --git a/lib/core/c++/boundary.hpp b/lib/core/c++/boundary.hpp index 3503681..adb473d 100644 --- a/lib/core/c++/boundary.hpp +++ b/lib/core/c++/boundary.hpp @@ -114,10 +114,18 @@ public: // Technically use .copy() auto& dfs_old = dfs_old_f.at(index); auto df_cpy = dfs_old; + + static_assert(Descriptor::D == 2u and Descriptor::Q == 9u, "Some parts are hard coded sadly"); - for(uint64_t i = 0u; i < Descriptor::Q; ++i){ - dfs_old.at({i}) = df_cpy.at({dfi::opposite_index.at(i)}); - } + dfs_old.at({0u}) = df_cpy.at({0u}); + dfs_old.at({1u}) = df_cpy.at({1u}); + dfs_old.at({2u}) = df_cpy.at({2u}); + dfs_old.at({3u}) = df_cpy.at({4u}); + dfs_old.at({4u}) = df_cpy.at({3u}); + dfs_old.at({5u}) = df_cpy.at({7u}); + dfs_old.at({6u}) = df_cpy.at({8u}); + dfs_old.at({7u}) = df_cpy.at({5u}); + dfs_old.at({8u}) = df_cpy.at({6u}); } }; diff --git a/lib/core/c++/grid.hpp b/lib/core/c++/grid.hpp index be86e18..c9a3b05 100644 --- a/lib/core/c++/grid.hpp +++ b/lib/core/c++/grid.hpp @@ -27,8 +27,8 @@ public: } }; +template<typename Schema> void clean_grid(saw::data<Schema>& info_field){ - } } } |