#include #include #include #include #include #include #include namespace kel { namespace lbm { constexpr uint64_t dim_x = 512ul; constexpr uint64_t dim_y = dim_x * 2ul; constexpr uint64_t particle_amount = 16ul; 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 saw::error_or setup_initial_conditions( saw::data>& fields, saw::data>& macros, saw::data, particle_amount>>& 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}} ); // 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() ); return saw::make_void(); } template void add_particles(saw::data>>& particles){ ////// TODO } template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, saw::data, particle_amount>>& particles, saw::data t_i, device& dev ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); { auto& p = particles.at({{0u}}); auto& p_coll = p.template get<"collision">(); 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> collision{0.8}; component> bb; 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; default: break; } }); }).wait(); // Step /* q.submit([&](acpp::sycl::handler& h){ // h.depends_on(collision_ev); }).wait(); */ return saw::make_void(); } } } template saw::error_or lbm_main(int argc, char** argv){ using namespace kel::lbm; using dfi = df_info; auto eo_lbm_env = lbm_directory(); if(eo_lbm_env.is_error()){ return std::move(eo_lbm_env.get_error()); } auto& lbm_env = eo_lbm_env.get_value(); auto out_dir = lbm_env.data_dir / "settling_cubes_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.05},{0.01},{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_data_ptr = saw::heap, particle_amount>>>(); std::cout<<"Estimated Bytes: "<,sch::MacroStruct>().get()<(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_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, particle_amount>, encode::Sycl> lbm_sycl_particle_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; } } { auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data); if(eov.is_error()){ return eov; } } sycl_q.wait(); 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,*lbm_particle_data_ptr,i,dev); if(eov.is_error()){ return eov; } } 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; } } } */ /* { { auto eov = dev.copy_to_host(lbm_sycl_data,*lbm_data_ptr); if(eov.is_error()){ return eov; } } { auto eov = write_csv_file(out_dir,"lbm",i.get(), *lbm_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<<" - "<