From ba9c23e4177ab9309f69155601578b118b2fd782 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 4 Feb 2026 16:28:48 +0100 Subject: Weird missing vtk writes --- examples/poiseulle_particles_2d_gpu/sim.cpp | 72 +++++++++++++++++++++++------ 1 file changed, 57 insertions(+), 15 deletions(-) (limited to 'examples/poiseulle_particles_2d_gpu/sim.cpp') diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 1ef7584..3084bca 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -7,8 +7,8 @@ namespace kel { namespace lbm { -constexpr uint64_t dim_x = 256u; -constexpr uint64_t dim_y = 32u; +constexpr uint64_t dim_x = 1024u; +constexpr uint64_t dim_y = 512u; namespace sch { using namespace saw::schema; @@ -24,10 +24,25 @@ using ChunkStruct = Struct< Member, "dfs">, Member, "dfs_old"> >; + +template +using VelChunk = Chunk, 0u, dim_x, dim_y>; + +template +using RhoChunk = Chunk, 0u, dim_x, dim_y>; + +template +using MacroStruct = Struct< + Member, "velocity">, + Member, "density"> +>; } template -saw::error_or setup_initial_conditions(saw::data>& fields){ +saw::error_or setup_initial_conditions( + saw::data>& fields, + saw::data>& macros +){ auto& info_f = fields.template get<"info">(); // Set everything as walls iterator::apply( @@ -69,13 +84,16 @@ saw::error_or setup_initial_conditions(saw::data> ); // auto& df_f = fields.template get<"dfs">(); - saw::data rho{1}; - saw::data> vel{}; - auto eq = equilibrium(rho,vel); + auto& rho_f = macros.template get<"density">(); + auto& vel_f = macros.template get<"velocity">(); iterator::apply( [&](auto& index){ auto& df = df_f.at(index); + auto& rho = rho_f.at(index); + rho.at({}) = {1}; + auto& vel = vel_f.at(index); + auto eq = equilibrium(rho,vel); df = eq; }, @@ -87,7 +105,12 @@ saw::error_or setup_initial_conditions(saw::data> } template -saw::error_or step(saw::data>,encode::Sycl>& fields, saw::data t_i, device& dev){ +saw::error_or step( + saw::data>,encode::Sycl>& fields, + saw::data>,encode::Sycl>& macros, + saw::data t_i, + device& dev +){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); @@ -96,7 +119,7 @@ saw::error_or step(saw::data>,encode::Sy // Need nicer things to handle the flow. I see improvement here component> collision{0.6}; component> bb; - component,encode::Sycl> flow_in{1.01}; + component,encode::Sycl> flow_in{1.001}; component,encode::Sycl> flow_out{1.0}; h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ @@ -114,15 +137,15 @@ saw::error_or step(saw::data>,encode::Sy bb.apply(fields,index,t_i); break; case 2u: - collision.apply(fields,index,t_i); + collision.apply(fields,macros,index,t_i); break; case 3u: flow_in.apply(fields,index,t_i); - collision.apply(fields,index,t_i); + collision.apply(fields,macros,index,t_i); break; case 4u: flow_out.apply(fields,index,t_i); - collision.apply(fields,index,t_i); + collision.apply(fields,macros,index,t_i); break; default: break; @@ -181,6 +204,7 @@ saw::error_or lbm_main(int argc, char** argv){ // saw::data> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap>>(); + auto lbm_macro_data_ptr = saw::heap>>(); device dev; @@ -188,21 +212,29 @@ saw::error_or lbm_main(int argc, char** argv){ sycl_q.wait(); { - auto eov = setup_initial_conditions(*lbm_data_ptr); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr); if(eov.is_error()){ return eov; } } saw::data, encode::Sycl> lbm_sycl_data{sycl_q}; + saw::data, encode::Sycl> lbm_sycl_macro_data{sycl_q}; sycl_q.wait(); auto lsd_view = make_chunk_struct_view(lbm_sycl_data); + auto lsdm_view = make_chunk_struct_view(lbm_sycl_macro_data); { auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data); if(eov.is_error()){ return eov; } } + { + auto eov = dev.copy_to_device(*lbm_macro_data_ptr,lbm_sycl_macro_data); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); saw::data time_steps{256ul}; @@ -211,13 +243,19 @@ saw::error_or lbm_main(int argc, char** argv){ std::string file_name = "tmp/t_"; file_name += std::to_string(i.get()); file_name += ".vtk"; - auto eov = write_vtk_file(file_name, *lbm_data_ptr); + auto eov = write_vtk_file(file_name, *lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } + { + auto eov = step(lsd_view,lsdm_view,i,dev); if(eov.is_error()){ return eov; } } { - auto eov = step(lsd_view,i,dev); + auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); if(eov.is_error()){ return eov; } @@ -234,10 +272,14 @@ saw::error_or lbm_main(int argc, char** argv){ std::string file_name = "tmp/t_"; file_name += std::to_string(time_steps.get()); file_name += ".vtk"; - auto eov = write_vtk_file(file_name, *lbm_data_ptr); + auto eov = write_vtk_file(file_name, *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } + auto eov2 = write_vtk_file((std::string{"tmp/df_"}+std::to_string(time_steps.get())+std::string{".vtk"}), *lbm_data_ptr); + if(eov2.is_error()){ + return eov2; + } } /* -- cgit v1.2.3