diff options
| -rw-r--r-- | examples/settling_cubes_2d_ibm_gpu/sim.cpp | 104 | ||||
| -rw-r--r-- | lib/core/c++/abstract/templates.hpp | 3 | ||||
| -rw-r--r-- | lib/core/c++/collision.hpp | 41 |
3 files changed, 41 insertions, 107 deletions
diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index 16b44c3..4501a54 100644 --- a/examples/settling_cubes_2d_ibm_gpu/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,43 +99,6 @@ 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(); } @@ -180,28 +124,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::BGKGuo, 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,21 +139,11 @@ 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; } @@ -255,13 +169,13 @@ saw::error_or<void> lbm_main(int argc, char** argv){ using dfi = df_info<T,Desc>; - auto eo_lbm_dir = lbm_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 / "settling_cubes_2d_ibm_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..f8312c0 100644 --- a/lib/core/c++/collision.hpp +++ b/lib/core/c++/collision.hpp @@ -115,39 +115,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 * 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_); } } }; |