summaryrefslogtreecommitdiff
path: root/examples/poiseulle_particles_2d_gpu/step.hpp
diff options
context:
space:
mode:
Diffstat (limited to 'examples/poiseulle_particles_2d_gpu/step.hpp')
-rw-r--r--examples/poiseulle_particles_2d_gpu/step.hpp54
1 files changed, 27 insertions, 27 deletions
diff --git a/examples/poiseulle_particles_2d_gpu/step.hpp b/examples/poiseulle_particles_2d_gpu/step.hpp
index aa0e382..37eb4e7 100644
--- a/examples/poiseulle_particles_2d_gpu/step.hpp
+++ b/examples/poiseulle_particles_2d_gpu/step.hpp
@@ -5,7 +5,7 @@
namespace kel {
namespace lbm {
-template<typename T, typename Desc>
+template<typename T, typename Desc, typename Coll>
saw::error_or<void> step(
const converter<T>& conv,
saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields,
@@ -17,29 +17,44 @@ saw::error_or<void> step(
auto& q = dev.get_handle();
auto& info_f = fields.template get<"info">();
auto& porous_f = macros.template get<"porosity">();
+
+ q.submit([&](acpp::sycl::handler& h){
+ component<T,Desc,cmpt::HlbmReset,encode::Sycl<saw::encode::Native>> hlbm_reset;
+ h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+ saw::data<sch::FixedArray<sch::UInt64,Desc::D>> 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<T,Desc,cmpt::HlbmOneParticle,encode::Sycl<saw::encode::Native>> hlbm_one_part;
+
+ h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+ saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+ for(uint64_t i = 0u; i < 1u; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
- component<T,Desc,cmpt::HlbmParticle,encode::Sycl<saw::encode::Native>> particle;
+ hlbm_one_part.apply(fields,macros,particles,index,t_i,{16u});
+ });
+ }).wait();
// auto coll_ev =
q.submit([&](acpp::sycl::handler& h){
component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.8};
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
- component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
-
- saw::data<sch::Scalar<T>> rho_b;
- rho_b.at({}) = 1.0;
- saw::data<sch::Vector<T,Desc::D>> vel_b;
- vel_b.at({{0u}}) = 0.015;
-
- component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
[&](){
uint64_t target_t_i = 16u;
if(t_i.get() < target_t_i){
- return 1.0 + (0.0002 / target_t_i) * t_i.get();
+ return 1.0 + (0.01 / target_t_i) * t_i.get();
}
- return 1.0002;
+ return 1.01;
}()
};
component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -63,12 +78,10 @@ saw::error_or<void> step(
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);
- // equi.apply(fields,index,t_i);
collision.apply(fields,macros,index,t_i);
break;
default:
@@ -77,19 +90,6 @@ saw::error_or<void> step(
});
}).wait();
- q.submit([&](acpp::sycl::handler& h){
- h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){
- particle.apply(fields,macros,particles,{{0u}},t_i);
- });
- }).wait();
-
- // Step
- /*
- q.submit([&](acpp::sycl::handler& h){
- // h.depends_on(collision_ev);
- }).wait();
- */
-
return saw::make_void();
}