#pragma once #include "common.hpp" namespace kel { namespace lbm { template saw::error_or step( const converter& conv, saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, saw::data>,encode::Sycl>& particles, saw::data t_i, device& dev ){ static constexpr saw::data sub_steps{0u}; auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); auto& ext_force_f = macros.template get<"external_force">(); if constexpr ( std::is_same_v ){ q.submit([&](acpp::sycl::handler& h){ component> hlbm_reset; 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]); } hlbm_reset.apply(fields,macros,index,t_i); }); }).wait(); q.submit([&](acpp::sycl::handler& h){ component> hlbm_one_part; h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){ saw::data> index; for(uint64_t i = 0u; i < 1u; ++i){ index.at({{i}}).set(idx[i]); } hlbm_one_part.apply(fields,macros,particles,index,t_i,{16u}); }); }).wait(); q.submit([&](acpp::sycl::handler& h){ component> collision{1.0}; 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(); }else if constexpr ( std::is_same_v ){ q.submit([&](acpp::sycl::handler& h){ component> fplbm_reset; 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]); } fplbm_reset.apply(fields,macros,index,t_i); // TODO apply external force // auto& ext_force = ext_force_f.at(index); }); }).wait(); q.submit([&](acpp::sycl::handler& h){ saw::data> gf; gf.at({{1u}}) = conv.acceleration_si_to_lbm({-9.81}).handle(); component> fplbm_one_part{gf}; h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){ saw::data> index; for(uint64_t i = 0u; i < 1u; ++i){ index.at({{i}}).set(idx[i]); } fplbm_one_part.apply(fields,macros,particles,index,t_i,{16u}); }); }).wait(); // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ component> bgk{1.0}; component> collision{1.0}; component> bb; 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: { component,encode::Sycl> flow_in{ [&]() -> saw::data> { saw::data> vel; { auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to()).handle(); vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle(); vel.at({{1u}}) = 0.0f; } return vel; }() }; flow_in.apply(fields,index,t_i); } bgk.apply(fields,macros,index,t_i); break; case 4u: flow_out.apply(fields,index,t_i); bgk.apply(fields,macros,index,t_i); break; default: break; } }); }).wait(); }else if constexpr ( std::is_same_v ){ q.submit([&](acpp::sycl::handler& h){ component> fplbm_reset; 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]); } fplbm_reset.apply(fields,macros,index,t_i); }); }).wait(); q.submit([&](acpp::sycl::handler& h){ component> fplbm_one_part; h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){ saw::data> index; for(uint64_t i = 0u; i < 1u; ++i){ index.at({{i}}).set(idx[i]); } fplbm_one_part.apply(fields,macros,particles,index,t_i,{16u}); }); }).wait(); // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ component> bgk{1.0}; component> collision{1.0}; component> bb; 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: { component,encode::Sycl> flow_in{ [&]() -> saw::data> { saw::data> vel; { auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to()).handle(); vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle(); vel.at({{1u}}) = 0.0f; } return vel; }() }; flow_in.apply(fields,index,t_i); } bgk.apply(fields,macros,index,t_i); break; case 4u: flow_out.apply(fields,index,t_i); bgk.apply(fields,macros,index,t_i); break; default: break; } }); }).wait(); } return saw::make_void(); } } }