diff options
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 53 | ||||
| -rw-r--r-- | lib/core/c++/geometry.hpp | 13 | ||||
| -rw-r--r-- | lib/core/c++/geometry/poiseulle_channel.hpp | 15 | ||||
| -rw-r--r-- | lib/core/c++/particle/blur.hpp | 1 |
4 files changed, 44 insertions, 38 deletions
diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 4bc9dfb..e7aec8e 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -259,22 +259,6 @@ saw::error_or<void> step( }); }).wait(); - q.submit([&](acpp::sycl::handler& h){ - component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; - - 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]); - } - - auto info = info_f.at(index); - - if(info.get() > 0u){ - stream.apply(fields,index,t_i); - } - }); - }).wait(); // Step /* @@ -380,35 +364,56 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } sycl_q.wait(); - saw::data<sch::UInt64> time_steps{4096ul}; + saw::data<sch::UInt64> time_steps{128ul}; + + auto& info_f = lbm_sycl_data.template get<"info">(); for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){ + // BC + Collision { - auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); + auto eov = step<T,Desc>(lsd_view,lsdm_view,*lbm_particle_data_ptr,i,dev); if(eov.is_error()){ return eov; } } + sycl_q.wait(); { - { - auto eov = write_vtk_file(out_dir,"t",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } + auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; } } { - auto eov = step<T,Desc>(lsd_view,lsdm_view,*lbm_particle_data_ptr,i,dev); + auto eov = write_vtk_file(out_dir,"t",time_steps.get(), *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } } + // Stream + sycl_q.submit([&](acpp::sycl::handler& h){ + component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; + + 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]); + } + + auto info = info_f.at(index); + + if(info.get() > 0u){ + stream.apply(fields,index,t_i); + } + }); + }).wait(); wait.poll(); if(print_status){ std::cout<<"Status: "<<i.get()<<" of "<<time_steps.get()<<" - "<<(i.template cast_to<sch::Float64>().get() * 100 / time_steps.get())<<"%"<<std::endl; print_status = false; } } + + // After Loop sycl_q.wait(); { auto eov = write_vtk_file(out_dir,"t",time_steps.get(), *lbm_macro_data_ptr); diff --git a/lib/core/c++/geometry.hpp b/lib/core/c++/geometry.hpp index c8a48a6..1c5d0a3 100644 --- a/lib/core/c++/geometry.hpp +++ b/lib/core/c++/geometry.hpp @@ -12,19 +12,6 @@ struct geometry { } }; */ -namespace cmpt { -struct PoiseulleChannel; -} - -template<typename Schema, typename Desc, typename Encode> -class component<Schema, Desc, cmpt::PoiseulleChannel, Encode> final { -private: -public: - template<typename CellFieldSchema> - void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>){ - auto& info_f = field.template get<"info">(); - } -}; // Ghost - 0 // Wall - 1 diff --git a/lib/core/c++/geometry/poiseulle_channel.hpp b/lib/core/c++/geometry/poiseulle_channel.hpp index f675a99..f719ec4 100644 --- a/lib/core/c++/geometry/poiseulle_channel.hpp +++ b/lib/core/c++/geometry/poiseulle_channel.hpp @@ -1,7 +1,22 @@ #pragma once +#include "../geometry.hpp" + namespace kel { namespace lbm { +namespace cmpt { +struct PoiseulleChannel; +} + +template<typename Schema, typename Desc, typename Encode> +class component<Schema, Desc, cmpt::PoiseulleChannel, Encode> final { +private: +public: + template<typename CellFieldSchema> + void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>){ + auto& info_f = field.template get<"info">(); + } +}; } } diff --git a/lib/core/c++/particle/blur.hpp b/lib/core/c++/particle/blur.hpp index a304e8d..7b93ae9 100644 --- a/lib/core/c++/particle/blur.hpp +++ b/lib/core/c++/particle/blur.hpp @@ -31,7 +31,6 @@ void blur_mask(saw::data<sch::Array<T,D>>& p_mask){ p_mask = blurred_mask; } - } } } |