diff options
| -rw-r--r-- | default.nix | 5 | ||||
| -rw-r--r-- | examples/settling_cubes_2d_ibm_gpu/.nix/derivation.nix (renamed from examples/settling_cubes_2d_ibm/.nix/derivation.nix) | 0 | ||||
| -rw-r--r-- | examples/settling_cubes_2d_ibm_gpu/SConscript (renamed from examples/settling_cubes_2d_ibm/SConscript) | 0 | ||||
| -rw-r--r-- | examples/settling_cubes_2d_ibm_gpu/SConstruct (renamed from examples/settling_cubes_2d_ibm/SConstruct) | 0 | ||||
| -rw-r--r-- | examples/settling_cubes_2d_ibm_gpu/sim.cpp (renamed from examples/settling_cubes_2d_ibm/sim.cpp) | 111 | ||||
| -rw-r--r-- | lib/core/c++/abstract/templates.hpp | 3 | ||||
| -rw-r--r-- | lib/core/c++/collision.hpp | 54 | ||||
| -rw-r--r-- | lib/core/c++/environment.hpp | 54 | ||||
| -rw-r--r-- | lib/core/c++/particle/particle.hpp | 5 |
9 files changed, 122 insertions, 110 deletions
diff --git a/default.nix b/default.nix index 2232e47..05ca001 100644 --- a/default.nix +++ b/default.nix @@ -171,6 +171,11 @@ in rec { inherit pname version stdenv forstio adaptive-cpp; inherit kel; }; + + settling_cubes_2d_ibm_gpu = pkgs.callPackage ./examples/settling_cubes_2d_ibm_gpu/.nix/derivation.nix { + inherit pname version stdenv forstio adaptive-cpp; + inherit kel; + }; heterogeneous_computing = pkgs.callPackage ./examples/heterogeneous_computing/.nix/derivation.nix { inherit pname version stdenv forstio adaptive-cpp; diff --git a/examples/settling_cubes_2d_ibm/.nix/derivation.nix b/examples/settling_cubes_2d_ibm_gpu/.nix/derivation.nix index d7f138b..d7f138b 100644 --- a/examples/settling_cubes_2d_ibm/.nix/derivation.nix +++ b/examples/settling_cubes_2d_ibm_gpu/.nix/derivation.nix diff --git a/examples/settling_cubes_2d_ibm/SConscript b/examples/settling_cubes_2d_ibm_gpu/SConscript index ae08056..ae08056 100644 --- a/examples/settling_cubes_2d_ibm/SConscript +++ b/examples/settling_cubes_2d_ibm_gpu/SConscript diff --git a/examples/settling_cubes_2d_ibm/SConstruct b/examples/settling_cubes_2d_ibm_gpu/SConstruct index 0611b67..0611b67 100644 --- a/examples/settling_cubes_2d_ibm/SConstruct +++ b/examples/settling_cubes_2d_ibm_gpu/SConstruct diff --git a/examples/settling_cubes_2d_ibm/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index 411efa7..05a4a09 100644 --- a/examples/settling_cubes_2d_ibm/sim.cpp +++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp @@ -45,7 +45,8 @@ using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>; template<typename T, typename Desc> using MacroStruct = Struct< Member<VelChunk<T,Desc>, "velocity">, - Member<RhoChunk<T>, "density"> + Member<RhoChunk<T>, "density">, + Member<VectorChunk<T,Desc>, "force"> >; //template<typename T, typename Desc> @@ -79,26 +80,6 @@ saw::error_or<void> setup_initial_conditions( info_f.get_dims(), {{1u,1u}} ); - - // Inflow - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(3u); - }, - {{0u,0u}}, - {{1u,dim_y}}, - {{0u,1u}} - ); - - // Outflow - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(4u); - }, - {{dim_x-1u,0u}}, - {{dim_x, dim_y}}, - {{0u,1u}} - ); // auto& df_f = fields.template get<"dfs_old">(); auto& rho_f = macros.template get<"density">(); @@ -118,47 +99,15 @@ saw::error_or<void> setup_initial_conditions( df_f.get_dims() ); - iterator<Desc::D>::apply( - [&](auto& index){ - auto& df = df_f.at(index); - auto& rho = rho_f.at(index); - rho.at({}) = {1}; - auto& vel = vel_f.at(index); - if(info_f.at(index).get() == 2u){ - vel.at({{0u}}) = 0.0; - } - auto eq = equilibrium<T,Desc>(rho,vel); - - df = eq; - }, - {},// 0-index - 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>(); - - 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() - ); - return saw::make_void(); } template<typename T, typename Desc> +void add_particles(saw::data<sch::Array<sch::Particle<T,Desc::D>>>& particles){ + ////// TODO +} + +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, @@ -180,28 +129,8 @@ 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.65}; + component<T,Desc,cmpt::BGKGuo, 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 = 64u; - if(t_i.get() < target_t_i){ - return 1.0 + (0.0015 / target_t_i) * t_i.get(); - } - return 1.0015; - }() - }; - component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; - 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; @@ -215,25 +144,17 @@ saw::error_or<void> step( case 0u: break; case 1u: - abb.apply(fields,index,t_i); + bb.apply(fields,index,t_i); break; case 2u: collision.apply(fields,macros,index,t_i); 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: break; } }); + + }).wait(); @@ -255,13 +176,13 @@ saw::error_or<void> lbm_main(int argc, char** argv){ using dfi = df_info<T,Desc>; - auto eo_lbm_dir = output_directory(); - if(eo_lbm_dir.is_error()){ - return std::move(eo_lbm_dir.get_error()); + auto eo_lbm_env = lbm_directory(); + if(eo_lbm_env.is_error()){ + return std::move(eo_lbm_env.get_error()); } - auto& lbm_dir = eo_lbm_dir.get_value(); + auto& lbm_env = eo_lbm_env.get_value(); - auto out_dir = lbm_dir / "poiseulle_particles_2d_bgk_gpu"; + auto out_dir = lbm_env.data_dir / "settling_cubes_2d_ibm_gpu"; { std::error_code ec; diff --git a/lib/core/c++/abstract/templates.hpp b/lib/core/c++/abstract/templates.hpp index d4d4ace..f675a99 100644 --- a/lib/core/c++/abstract/templates.hpp +++ b/lib/core/c++/abstract/templates.hpp @@ -1,4 +1,7 @@ #pragma once namespace kel { +namespace lbm { + +} } diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp index ef6bfe2..df45bbe 100644 --- a/lib/core/c++/collision.hpp +++ b/lib/core/c++/collision.hpp @@ -95,10 +95,21 @@ class component<T, Descriptor, cmpt::BGKGuo, Encode> { private: typename saw::native_data_type<T>::type relaxation_; saw::data<T> frequency_; + saw::data<sch::Vector<T,Descriptor::D>> global_force_; public: component(typename saw::native_data_type<T>::type relaxation__): relaxation_{relaxation__}, - frequency_{typename saw::native_data_type<T>::type(1) / relaxation_} + frequency_{typename saw::native_data_type<T>::type(1) / relaxation_}, + global_force_{} + {} + + component( + typename saw::native_data_type<T>::type relaxation__, + saw::data<sch::Vector<T,Descriptor::D>>& global_force__ + ): + relaxation_{relaxation__}, + frequency_{typename saw::native_data_type<T>::type(1) / relaxation_}, + global_force_{global_force__} {} using Component = cmpt::BGKGuo; @@ -115,39 +126,56 @@ public: static constexpr saw::string_literal after = ""; static constexpr saw::string_literal before = "streaming"; - template<typename CellFieldSchema> - void apply(saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64, Descriptor::D>> index, saw::data<sch::UInt64> time_step){ + template<typename CellFieldSchema, typename MacroFieldSchema> + void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step) const { + // void apply(saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64, Descriptor::D>> index, saw::data<sch::UInt64> time_step){ bool is_even = ((time_step.get() % 2) == 0); auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">(); auto& dfs = dfs_old_f.at(index); - saw::data<sch::Scalar<T>> rho; - saw::data<sch::Vector<T,Descriptor::D>> vel; + auto& rho_f = macros.template get<"density">(); + auto& vel_f = macros.template get<"velocity">(); + + saw::data<sch::Scalar<T>>& rho = rho_f.at(index); + saw::data<sch::Vector<T,Descriptor::D>>& vel = vel_f.at(index); + compute_rho_u<T,Descriptor>(dfs_old_f.at(index),rho,vel); auto eq = equilibrium<T,Descriptor>(rho,vel); using dfi = df_info<T,Descriptor>; - auto& force_f = field.template get<"force">(); + auto& force_f = macros.template get<"force">(); auto& force = force_f.at(index); + + saw::data<sch::Scalar<T>> dfi_inv_cs2; + dfi_inv_cs2.at({}).set(dfi::inv_cs2); for(uint64_t i = 0u; i < Descriptor::Q; ++i){ // saw::data<T> ci_min_u{0}; - saw::data<T> ci_dot_u{0}; + saw::data<sch::Vector<T,Descriptor::D>> ci; for(uint64_t d = 0u; d < Descriptor::D; ++d){ - ci_dot_u = ci_dot_u + vel.at({{d}}) * saw::data<T>{static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d])}; + ci.at({{d}}).set(static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d])); } - saw::data<T> F_i{0};// = saw::data<T>{dfi::weights[i]} * (ci_dot_F * dfi::inv_cs2 + ci_dot_u * ci_dot_F * (dfi::inv_cs2 * dfi::inv_cs2)); + auto ci_dot_u = saw::math::dot(ci,vel); + + // saw::data<sch::Vector<T,Descriptor::D>> F_i; + // F_i = f * (c_i - u * ics2 + <c_i,u> * c_i * ics2 * ics2) * w_i; + saw::data<sch::Scalar<T>> w; + w.at({}).set(dfi::weights[i]); + auto F_i_d = saw::math::dot((force+global_force_) * w, (ci - vel * dfi_inv_cs2 + ci * ci_dot_u * dfi_inv_cs2 * dfi_inv_cs2 )); + /* + saw::data<sch::Scalar<T>> F_i_sum; for(uint64_t d = 0u; d < Descriptor::D; ++d){ - F_i = F_i + - saw::data<T>{dfi::weights[i]} * ((saw::data<T>{static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d])} - - vel.at({d}) ) * dfi::inv_cs2 + ci_dot_u * saw::data<T>{static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d])} * dfi::inv_cs2 * dfi::inv_cs2 ) * force.at({d}); + saw::data<sch::Scalar<T>> F_i_d; + F_i_d.at({}) = F_i.at({{d}}); + F_i_sum = F_i_sum + F_i_d; } + */ - dfs.at({i}) = dfs.at({i}) + frequency_ * (eq.at(i) - dfs.at({i}) ) + F_i * (saw::data<T>{1} - saw::data<T>{0.5f} * frequency_); + dfs.at({i}) = dfs.at({i}) + frequency_ * (eq.at(i) - dfs.at({i}) ) + F_i_d.at({}) * (saw::data<T>{1} - saw::data<T>{0.5f} * frequency_); } } }; diff --git a/lib/core/c++/environment.hpp b/lib/core/c++/environment.hpp index 9e587d0..27d8f3a 100644 --- a/lib/core/c++/environment.hpp +++ b/lib/core/c++/environment.hpp @@ -7,18 +7,70 @@ namespace kel { namespace lbm { +struct environment { + std::filesystem::path lbm_dir; + std::filesystem::path data_dir; +}; + +saw::error_or<environment> lbm_directory(){ + namespace fs = std::filesystem; + + const char* home_dir = std::getenv("HOME"); + if(not home_dir){ + return saw::make_error<saw::err::not_found>("Couldn't find home dir"); + } + + environment env; + + env.lbm_dir = std::filesystem::path{home_dir} / ".lbm"; + + { + env.data_dir = env.lbm_dir / "data"; + // LBM Data Location + auto lbm_dir_config = env.lbm_dir / "data_dir_location.txt"; + if( fs::exists(lbm_dir_config) && fs::is_regular_file(lbm_dir_config) ){ + std::ifstream file{lbm_dir_config}; + + if(file.is_open()){ + std::stringstream buffer; + buffer <<file.rdbuf(); + env.data_dir = fs::path{buffer.str()}; + } + } + } + + return env; +} + /** * Returns the default output directory. * Located outside the project dir because dispatching build jobs with output data in the git directory * also copies simulated data which takes a long time. */ saw::error_or<std::filesystem::path> output_directory(){ + namespace fs = std::filesystem; + const char* home_dir = std::getenv("HOME"); if(not home_dir){ return saw::make_error<saw::err::not_found>("Couldn't find home dir"); } - return std::filesystem::path{home_dir} / ".lbm"; + auto lbm_dir = std::filesystem::path{home_dir} / ".lbm"; + + { + auto lbm_dir_config = lbm_dir / "data_dir_location.txt"; + if( fs::exists(lbm_dir_config) && fs::is_regular_file(lbm_dir_config) ){ + std::ifstream file{lbm_dir_config}; + + if(file.is_open()){ + std::stringstream buffer; + buffer <<file.rdbuf(); + return fs::path{buffer.str()}; + } + } + } + + return lbm_dir / "data"; } } diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index 8ef590d..f8104fd 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -43,12 +43,15 @@ using ParticleCollisionSpheroid = Struct< Member<Scalar<T>, "radius"> >; +tem + template<typename T, uint64_t D, typename CollisionType = ParticleCollisionSpheroid<T>> using Particle = Struct< Member<ParticleRigidBody<T,D>, "rigid_body">, Member<CollisionType, "collision">, // Problem is that dynamic data would two layered - // Member<FixedArray<Float64,D,D,D>, "mask">, + // Member<Array<Float64,D>, "mask">, + Member<UInt64, "mask_id">, Member<Scalar<T>, "mass"> >; } |