diff options
Diffstat (limited to 'examples/poiseulle_particles_2d_ibm_gpu')
| -rw-r--r-- | examples/poiseulle_particles_2d_ibm_gpu/sim.cpp | 87 |
1 files changed, 64 insertions, 23 deletions
diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp index 9905906..ea49d51 100644 --- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp @@ -54,12 +54,21 @@ using MacroStruct = Struct< //using ParticleArray = Array< // Particle<T,Desc::D> //>; + +template<typename T, typename Desc> +using ParticleSpheroidGroup = ParticleGroup< + T, + Desc::D, + sch::ParticleCollisionSpheroid<T,2.0f> +>; + } template<typename T, typename Desc> saw::error_or<void> setup_initial_conditions( saw::data<sch::ChunkStruct<T,Desc>>& fields, - saw::data<sch::MacroStruct<T,Desc>>& macros + saw::data<sch::MacroStruct<T,Desc>>& macros, + saw::data<sch::ParticleSpheroidGroup<T,Desc>>& particles ){ auto& info_f = fields.template get<"info">(); // Set everything as walls @@ -136,25 +145,41 @@ saw::error_or<void> setup_initial_conditions( df_f.get_dims(), {{1u,1u}} ); - - iterator<Desc::D>::apply( - [&](auto& index){ - saw::data<sch::Vector<T,Desc::D>> middle, ind_vec; - middle.at({{0u}}) = dim_x * 0.25; - middle.at({{1u}}) = dim_y * 0.5; - ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to<T>(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>(); + // Particles + { + saw::data<sch::Scalar<T>> dense_p; + dense_p.at({}).set(1); - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - info_f.at(index) = 1u; - } - }, - {},// 0-index - df_f.get_dims() - ); + auto& spheroid_grp = particles; + + spheroid_grp = create_spheroid_particle_group<T,Desc::D,2.0f>( + 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(); } @@ -163,6 +188,7 @@ template<typename T, typename Desc> saw::error_or<void> step( saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields, saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros, + saw::data<sch::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& particles, saw::data<sch::UInt64> t_i, device& dev ){ @@ -172,9 +198,10 @@ saw::error_or<void> step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here - saw::data<sch::Vector<T,Desc::D>> f; - f.at({{0u}}) = 0.0; - f.at({{1u}}) = -1.0; + saw::data<sch::Vector<T,Desc::D>> f; + f.at({{0u}}) = 0.0; + f.at({{1u}}) = -1.0; + component<T,Desc,cmpt::BGKGuo, encode::Sycl<saw::encode::Native>> collision{0.65,f}; component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb; component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb; @@ -237,6 +264,11 @@ saw::error_or<void> step( // 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){ + + }); + }).wait(); return saw::make_void(); } @@ -285,6 +317,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ // 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>>>(); auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>(); + auto lbm_particle_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>(); std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl; @@ -310,7 +343,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ sycl_q.wait(); { - auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr); + auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_ptr); if(eov.is_error()){ return eov; } @@ -324,6 +357,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q}; saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q}; + saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_particle{sycl_q}; sycl_q.wait(); { @@ -338,9 +372,16 @@ saw::error_or<void> lbm_main(int argc, char** argv){ 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<sch::UInt64> time_steps{16u*4096ul}; auto& info_f = lsd_view.template get<"info">(); @@ -348,7 +389,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){ // BC + Collision { - auto eov = step<T,Desc>(lsd_view,lsdm_view,i,dev); + auto eov = step<T,Desc>(lsd_view,lsdm_view,lsdp_view,i,dev); if(eov.is_error()){ return eov; } |