#include #include #include #include #include #include #include #include #include namespace kel { namespace lbm { constexpr uint64_t dim_y = 256ul; constexpr uint64_t dim_x = dim_y * 20ul; constexpr uint64_t particle_amount = 1ul; namespace sch { using namespace saw::schema; using InfoChunk = Chunk; template using DfChunk = Chunk, 1u, dim_x, dim_y>; template using ScalarChunk = Chunk, 0u, dim_x, dim_y>; template using VectorChunk = Chunk, 0u, dim_x, dim_y>; template using ChunkStruct = Struct< Member, 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">, Member, "force"> >; //template //using ParticleArray = Array< // Particle //>; template using ParticleSpheroidGroup = ParticleGroup< T, Desc::D, sch::ParticleCollisionSpheroid >; } template saw::error_or setup_initial_conditions( saw::data>& fields, saw::data>& macros, saw::data>& particles ){ auto& info_f = fields.template get<"info">(); // Set everything as walls iterator::apply( [&](auto& index){ info_f.at(index).set(1u); }, {}, info_f.get_dims(), {} ); // Fluid iterator::apply( [&](auto& index){ info_f.at(index).set(2u); }, {}, info_f.get_dims(), {{1u,1u}} ); // Inflow iterator::apply( [&](auto& index){ info_f.at(index).set(3u); }, {{0u,0u}}, {{1u,dim_y}}, {{0u,1u}} ); // Outflow iterator::apply( [&](auto& index){ info_f.at(index).set(4u); }, {{dim_x-1u,0u}}, {{dim_x, dim_y}}, {{0u,1u}} ); // auto& df_f = fields.template get<"dfs_old">(); 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; }, {},// 0-index df_f.get_dims() ); 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); if(info_f.at(index).get() == 2u){ vel.at({{0u}}) = 0.0; } auto eq = equilibrium(rho,vel); df = eq; }, {},// 0-index df_f.get_dims(), {{1u,1u}} ); // Particles { saw::data> dense_p; dense_p.at({}).set(1); auto& spheroid_grp = particles; spheroid_grp = create_spheroid_particle_group( dense_p, {32u} ); { auto& p = spheroid_grp.template get<"particles">(); p = {{{1u}}}; iterator<1u>::apply( [&](auto& index){ auto& p_ind = p.at(index); auto& p_rb = p_ind.template get<"rigid_body">(); auto& p_pos = p_rb.template get<"position">(); auto& p_pos_old = p_rb.template get<"position_old">(); // Set position p_pos.at({{0u}}) = dim_x * 0.25; p_pos.at({{1u}}) = dim_y * 0.5; p_pos_old = p_pos; }, {}, p.meta() ); } } return saw::make_void(); } template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, saw::data>,encode::Sycl>& particles, saw::data t_i, device& dev ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here saw::data> f; f.at({{0u}}) = 0.0; f.at({{1u}}) = -1.0; component> collision{0.65,f}; component> bb; component,encode::Sycl> abb; saw::data> rho_b; rho_b.at({}) = 1.0; saw::data> vel_b; vel_b.at({{0u}}) = 0.015; component> equi{rho_b,vel_b}; component,encode::Sycl> flow_in{ [&](){ uint64_t target_t_i = 32u; if(t_i.get() < target_t_i){ return 1.0 + (0.0002 / target_t_i) * t_i.get(); } return 1.0002; }() }; component,encode::Sycl> flow_out{1.0}; h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; for(uint64_t i = 0u; i < Desc::D; ++i){ index.at({{i}}).set(idx[i]); } auto info = info_f.at(index); switch(info.get()){ case 0u: break; case 1u: 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); collision.apply(fields,macros,index,t_i); break; case 4u: flow_out.apply(fields,index,t_i); collision.apply(fields,macros,index,t_i); break; default: break; } }); }).wait(); // Step /* q.submit([&](acpp::sycl::handler& h){ // h.depends_on(collision_ev); }).wait(); */ q.submit([&](acpp::sycl::handler& h){ h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ auto& vel = macros.template get<"velocity">(); auto& ps = particles; auto& mask = ps.template get<"mask">(); auto& dense = ps.template get<"density">().at({}); auto& com = ps.template get<"center_of_mass">(); auto& parts = ps.template get<"particles">(); auto& p_i = parts.at({{0u}}); auto& p_i_rb = p_i.template get<"rigid_body">(); /// 0. Iterate over mask and calculate position in LBM grid /// In this case it's simple since I'm too lazy to do scaling and rotation /// Technically scale => rotate => translate /// Here it's only translate auto& p_i_rb_pos = p_i_rb.template get<"position">(); iterator::apply([&](const auto& index){ /// Calculate the shift from the mask saw::data> index_shift; for(uint64_t i = 0u; i < Desc::D; ++i){ index_shift.at({{i}}) = index.at({i}).template cast_to() - com.at({{i}}); } /// TODO 1. Calculate force pickup from neigbouring u_vel cells auto inter_vel = n_linear_interpolate(vel,index_shift); /// TODO 3. Distribute force to fluid }, {}, mask.meta()); }); }).wait(); return saw::make_void(); } } } template saw::error_or lbm_main(int argc, char** argv){ using namespace kel::lbm; /* auto eo_lbm_args = setup_lbm_env,sch::Tuple<>>>(argc,argv); if(eo_lbm_args.is_error()){ return std::move(eo_lbm_args.get_error()); } auto& lbm_args = eo_lbm_args.get_value(); */ using dfi = df_info; auto eo_lbm_dir = output_directory(); if(eo_lbm_dir.is_error()){ return std::move(eo_lbm_dir.get_error()); } auto& lbm_dir = eo_lbm_dir.get_value(); auto out_dir = lbm_dir / "poiseulle_particles_2d_ibm_gpu"; { std::error_code ec; std::filesystem::create_directories(out_dir,ec); if(ec != std::errc{}){ return saw::make_error("Could not create output directory"); } } converter conv { // delta_x {{1.0}}, // delta_t {{1.0}} }; print_lbm_meta(conv,{0.1},{1e-4},{0.4 * dim_y}); // saw::data> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap>>(); auto lbm_macro_data_ptr = saw::heap>>(); auto lbm_particle_ptr = saw::heap>>(); std::cout<<"Estimated Bytes: "<,sch::MacroStruct>().get()<(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_ptr); if(eov.is_error()){ return eov; } } { auto eov = write_vtk_file(out_dir,"initial_state",0u,*lbm_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}; saw::data, encode::Sycl> lbm_sycl_particle{sycl_q}; sycl_q.wait(); { 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; } } { auto eov = dev.copy_to_device(*lbm_particle_ptr,lbm_sycl_particle); if(eov.is_error()){ return eov; } } sycl_q.wait(); auto lsd_view = make_view(lbm_sycl_data); auto lsdm_view = make_view(lbm_sycl_macro_data); auto lsdp_view = make_view(lbm_sycl_particle); saw::data time_steps{16u*4096ul}; auto& info_f = lsd_view.template get<"info">(); for(saw::data i{0u}; i < time_steps and krun; ++i){ // BC + Collision { auto eov = step(lsd_view,lsdm_view,lsdp_view,i,dev); if(eov.is_error()){ return eov; } } 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()){ return eov; } } { auto eov = write_vtk_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } } } */ 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> stream; h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; for(uint64_t i = 0u; i < Desc::D; ++i){ index.at({{i}}).set(idx[i]); } auto info = info_f.at(index); if(info.get() > 0u){ stream.apply(lsd_view,index,i); } }); }).wait(); wait.poll(); if(print_status){ std::cout<<"Status: "<().get() * 100 / time_steps.get())<<"%"<(argc, argv); if(eov.is_error()){ auto& err = eov.get_error(); std::cerr<<"[Error] "< 0u){ std::cerr<<" - "<