diff options
| -rw-r--r-- | examples/poiseulle_particles_2d_bgk_gpu/sim.cpp | 13 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp | 16 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 12 | ||||
| -rw-r--r-- | lib/core/c++/write_vtk.hpp | 4 | ||||
| -rw-r--r-- | lib/sycl/c++/data.hpp | 10 |
5 files changed, 31 insertions, 24 deletions
diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp index 69e94c3..b8ac35e 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp @@ -183,11 +183,11 @@ saw::error_or<void> step( component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ [&](){ - uint64_t target_t_i = 64u; + uint64_t target_t_i = 32u; if(t_i.get() < target_t_i){ - return 1.0 + (0.0015 / target_t_i) * t_i.get(); + return 1.0 + (0.0002 / target_t_i) * t_i.get(); } - return 1.0015; + return 1.0002; }() }; component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; @@ -212,6 +212,7 @@ 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: @@ -266,7 +267,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ {{1.0}} }; - print_lbm_meta<T,Desc>(conv,{0.05},{0.01},{0.4 * dim_y}); + print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{0.4 * dim_y}); // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>(); @@ -340,7 +341,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } sycl_q.wait(); - { + /* + if(i.get() % 128u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); if(eov.is_error()){ @@ -354,6 +356,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } } + */ /* { { diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp index df57bbc..7cfd567 100644 --- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -191,9 +191,9 @@ saw::error_or<void> step( [&](){ uint64_t target_t_i = 64u; if(t_i.get() < target_t_i){ - return 1.0 + (0.0015 / target_t_i) * t_i.get(); + return 1.0 + (0.0002 / target_t_i) * t_i.get(); } - return 1.0015; + return 1.0002; }() }; component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; @@ -211,19 +211,19 @@ saw::error_or<void> step( case 0u: break; case 1u: - abb.apply(fields,index,t_i); + bb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); break; case 3u: - //flow_in.apply(fields,index,t_i); - equi.apply(fields,index,t_i); + 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); + flow_out.apply(fields,index,t_i); + // equi.apply(fields,index,t_i); collision.apply(fields,macros,index,t_i); break; default: @@ -274,7 +274,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ {{1.0}} }; - print_lbm_meta<T,Desc>(conv,{0.05},{0.01},{0.4 * dim_y}); + print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{0.4 * dim_y}); // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>(); diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index 7768ca4..75d1a14 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -207,9 +207,9 @@ saw::error_or<void> step( [&](){ uint64_t target_t_i = 64u; if(t_i.get() < target_t_i){ - return 1.0 + (0.0015 / target_t_i) * t_i.get(); + return 1.0 + (0.0002 / target_t_i) * t_i.get(); } - return 1.0015; + return 1.0002; }() }; component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; @@ -227,7 +227,7 @@ saw::error_or<void> step( case 0u: break; case 1u: - abb.apply(fields,index,t_i); + bb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); @@ -289,12 +289,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){ converter<T> conv { // delta_x - {{1e-3}}, + {{1.0}}, // delta_t - {{1e-06}} + {{1.0}} }; - print_lbm_meta<T,Desc>(conv,{0.1},{1e-2},{0.4 * dim_y *1e-3}); + print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{0.4 * dim_y}); // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>(); diff --git a/lib/core/c++/write_vtk.hpp b/lib/core/c++/write_vtk.hpp index f925361..e852172 100644 --- a/lib/core/c++/write_vtk.hpp +++ b/lib/core/c++/write_vtk.hpp @@ -136,7 +136,7 @@ struct lbm_vtk_writer<sch::Vector<T,D>> { template<typename T> struct lbm_vtk_writer<sch::Scalar<T>> { static saw::error_or<void> apply_header(std::ostream& vtk_file, std::string_view name){ - vtk_file<<"SCALARS "<<name<<" float\n"; + vtk_file<<"SCALARS "<<name<<" float 1\n"; vtk_file<<"LOOKUP_TABLE default\n"; return saw::make_void(); } @@ -206,7 +206,7 @@ struct lbm_vtk_writer<sch::Struct<sch::Member<sch::Chunk<MemberT,Ghost,Dims...>, if constexpr (sizeof...(MemberT) > 0u){ // POINT DATA { - vtk_file << "POINT_DATA " << pd_size.get() <<"\n"; + vtk_file << "\nPOINT_DATA " << pd_size.get() <<"\n"; } // HEADER TO BODY diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 71627de..9f43848 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -381,7 +381,7 @@ public: return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; } - auto flat_data(){ + auto flat_data() const { return values_.flat_data(); } @@ -741,8 +741,10 @@ struct sycl_copy_helper<sch::Chunk<Sch,Ghost,Dims...>, Encode> final { static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + auto flat_size = host_data.flat_size(); + q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, saw::ct_multiply<uint64_t,Dims...>::value); + h.copy(sycl_ptr,host_ptr, flat_size.get()); }).wait(); return saw::make_void(); } @@ -753,8 +755,10 @@ struct sycl_copy_helper<sch::Chunk<Sch,Ghost,Dims...>, Encode> final { static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + auto flat_size = host_data.flat_size(); + q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, saw::ct_multiply<uint64_t,Dims...>::value); + h.copy(host_ptr,sycl_ptr, flat_size.get()); }).wait(); return saw::make_void(); } |
