Merge branch 'dev'

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">
 >;
 }