diff options
| -rw-r--r-- | examples/poiseulle_3d_gpu/sim.cpp | 10 | ||||
| -rw-r--r-- | lib/core/c++/collision.hpp | 9 | ||||
| -rw-r--r-- | lib/core/c++/descriptor.hpp | 58 | ||||
| -rw-r--r-- | lib/core/tests/memory.cpp | 4 |
4 files changed, 52 insertions, 29 deletions
diff --git a/examples/poiseulle_3d_gpu/sim.cpp b/examples/poiseulle_3d_gpu/sim.cpp index 624566f..df4647e 100644 --- a/examples/poiseulle_3d_gpu/sim.cpp +++ b/examples/poiseulle_3d_gpu/sim.cpp @@ -59,6 +59,7 @@ using MacroStruct = Struct< template<typename T, typename Desc> saw::error_or<void> setup_initial_conditions( + converter<T>& conv, saw::data<sch::ChunkStruct<T,Desc>>& fields, saw::data<sch::MacroStruct<T,Desc>>& macros, saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_size>>& particles @@ -148,6 +149,7 @@ saw::error_or<void> setup_initial_conditions( template<typename T, typename Desc> saw::error_or<void> step( + converter<T>& conv, 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::UInt64> t_i, @@ -159,7 +161,7 @@ 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 - component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.6}; + component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{conv.template kinematic_viscosity_si_to_tau<Desc>({0.1}).handle()}; // component<T,Desc,cmpt::HLBM,encode::Sycl<saw::encode::Native>> collision{0.6}; component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb; @@ -255,7 +257,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } auto& lbm_dir = eo_lbm_dir.get_value(); - auto out_dir = lbm_dir / "poiseulle_particles_3d_gpu"; + auto out_dir = lbm_dir / "poiseulle_3d_gpu"; { std::error_code ec; @@ -302,7 +304,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,*lbm_particle_data_ptr); + auto eov = setup_initial_conditions<T,Desc>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); if(eov.is_error()){ return eov; } @@ -355,7 +357,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ }*/ } { - auto eov = step<T,Desc>(lsd_view,lsdm_view,i,dev); + auto eov = step<T,Desc>(conv,lsd_view,lsdm_view,i,dev); if(eov.is_error()){ return eov; } diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp index 349833f..f259c9f 100644 --- a/lib/core/c++/collision.hpp +++ b/lib/core/c++/collision.hpp @@ -18,12 +18,17 @@ struct BGKGuo {}; template<typename T, typename Descriptor, typename Encode> class component<T, Descriptor, cmpt::BGK, Encode> { private: - typename saw::native_data_type<T>::type relaxation_; + saw::data<T> relaxation_; saw::data<T> frequency_; public: component(typename saw::native_data_type<T>::type relaxation__): + relaxation_{{relaxation__}}, + frequency_{saw::data<T>{1} / relaxation_} + {} + + component(const saw::data<T>& relaxation__): relaxation_{relaxation__}, - frequency_{typename saw::native_data_type<T>::type(1) / relaxation_} + frequency_{saw::data<T>{1} / relaxation_} {} using Component = cmpt::BGK; diff --git a/lib/core/c++/descriptor.hpp b/lib/core/c++/descriptor.hpp index e38daee..73f0cce 100644 --- a/lib/core/c++/descriptor.hpp +++ b/lib/core/c++/descriptor.hpp @@ -54,6 +54,22 @@ struct CellFieldStruct { template<typename T, typename Desc> class df_info{}; +/* +namespace impl { +template<typename Desc> +struct df_ct_helper { + template<uint64_t i> + static constexpr uint64_t apply_i(const std::array<std::array<int32_t,Desc::D>,Desc::Q>& dirs, const std::array<T,Desc::D+1u>& inp){ + if constexpr ( i < Desc::Q ){ + for(uint64_t j = 0u; j < Desc::D; ++j){ + } + } + return 0u; + } +}; +} +*/ + template<typename T> class df_info<T,sch::Descriptor<1,3>> { public: @@ -214,34 +230,34 @@ public: static constexpr std::array<typename saw::native_data_type<T>::type,Q> weights = { 8./27., // 1D - 2./27., - 2./27., - // 2D - 2./27., - 1./9., 1./36., 1./36., + // 2D 1./36., + 1./54., + 1./54., 1./36., + 1./54., + 1./54., // 3D - 8./27., - 1./9., - 1./9., - 1./9., - 1./9., - 1./36., - 1./36., 1./36., + 1./54., + 1./54., + 1./54., + 1./216., + 1./216., + 1./54., + 1./216., + 1./216., 1./36., - 8./27., - 1./9., - 1./9., - 1./9., - 1./9., - 1./36., - 1./36., - 1./36., - 1./36. + 1./54., + 1./54., + 1./54., + 1./216., + 1./216., + 1./54., + 1./216., + 1./216. }; static constexpr std::array<uint64_t,Q> opposite_index = { diff --git a/lib/core/tests/memory.cpp b/lib/core/tests/memory.cpp index cdc6f8c..d59d212 100644 --- a/lib/core/tests/memory.cpp +++ b/lib/core/tests/memory.cpp @@ -46,13 +46,13 @@ SAW_TEST("Memory Estimate Scalar"){ SAW_TEST("Memory Estimate Chunk"){ using namespace kel::lbm; - SAW_EXPECT((memory_estimate<sch::TChunk>().get() == 480u), std::string{"TChunk isn't 480 bytes "} + std::to_string(memory_estimate<sch::TChunk>().get()) ); + SAW_EXPECT((memory_estimate<sch::TChunk>().get() == 480u), std::string{"TChunk isn't 480 bytes, but is "} + std::to_string(memory_estimate<sch::TChunk>().get()) ); } SAW_TEST("Memory Estimate Struct of Chunk"){ using namespace kel::lbm; - SAW_EXPECT((memory_estimate<sch::TChunkStruct>().get() == 960u), std::string{"TChunk isn't 480 bytes "} + std::to_string(memory_estimate<sch::TChunk>().get()) ); + SAW_EXPECT((memory_estimate<sch::TChunkStruct>().get() == 960u), std::string{"TChunkStruct isn't 960 bytes, but is "} + std::to_string(memory_estimate<sch::TChunk>().get()) ); } } |