Merge branch 'dev'

12 files changed, 102 insertions, 166 deletions

diff --git a/default.nix b/default.nix
index d3692b6..04b2616 100644
--- a/default.nix
+++ b/default.nix
@@ -36,7 +36,7 @@ let
 		inherit stdenv llvmPackages lld;
 
 		config = {
-			rocm = false;
+			rocm = true;
 		};
 	};
 
@@ -86,7 +86,7 @@ in rec {
 		};
 	};
 
-	examples = {
+	ex = {
 		cavity_2d_gpu = pkgs.callPackage ./examples/cavity_2d_gpu/.nix/derivation.nix {
 			inherit pname version stdenv forstio adaptive-cpp;
 			inherit kel;
@@ -165,11 +165,11 @@ in rec {
 		poiseulle_particles_2d_all_gpu = pkgs.symlinkJoin {
 			name = "poiseulle_particles_2d_all_gpu";
 			paths = [
-				examples.poiseulle_particles_2d_bgk_gpu
-				examples.poiseulle_particles_2d_psm_gpu
-				examples.poiseulle_particles_2d_hlbm_gpu
-				examples.poiseulle_particles_2d_fplbm_gpu
-				examples.poiseulle_particles_2d_ibm_gpu
+				ex.poiseulle_particles_2d_bgk_gpu
+				ex.poiseulle_particles_2d_psm_gpu
+				ex.poiseulle_particles_2d_hlbm_gpu
+				ex.poiseulle_particles_2d_fplbm_gpu
+				ex.poiseulle_particles_2d_ibm_gpu
 			];
 		};
 		
@@ -241,7 +241,7 @@ in rec {
 		settling_particles = pkgs.symlinkJoin {
 			name = "kel-lbm-settling_particles-${version}";
 			paths = [
-				examples.settling_cubes_2d_ibm_gpu
+				ex.settling_cubes_2d_ibm_gpu
 			];
 		};
 	};
@@ -255,36 +255,36 @@ in rec {
 			inherit kel;
 		};
 
-		examples = pkgs.symlinkJoin {
+		ex = pkgs.symlinkJoin {
 			name = "kel-lbm-sims-${version}";
 			paths = [
 				kel.lbm.core
 				kel.lbm.sycl
-				examples.meta
-				examples.poiseulle_particles_2d_bgk_gpu
-				examples.poiseulle_particles_2d_psm_gpu
-				examples.poiseulle_particles_2d_hlbm_gpu
-				examples.poiseulle_particles_2d_fplbm_gpu
-				examples.poiseulle_particles_2d_ibm_gpu
+				ex.meta
+				ex.poiseulle_particles_2d_bgk_gpu
+				ex.poiseulle_particles_2d_psm_gpu
+				ex.poiseulle_particles_2d_hlbm_gpu
+				ex.poiseulle_particles_2d_fplbm_gpu
+				ex.poiseulle_particles_2d_ibm_gpu
 			];
 		};
 
 		poiseulle_particles_2d = pkgs.symlinkJoin {
 			name = "kel-lbm-sims-poiseulle_particles_2d-${version}";
 			paths = [
-				examples.poiseulle_particles_2d_bgk_gpu
-				examples.poiseulle_particles_2d_psm_gpu
-				examples.poiseulle_particles_2d_hlbm_gpu
-				examples.poiseulle_particles_2d_fplbm_gpu
+				ex.poiseulle_particles_2d_bgk_gpu
+				ex.poiseulle_particles_2d_psm_gpu
+				ex.poiseulle_particles_2d_hlbm_gpu
+				ex.poiseulle_particles_2d_fplbm_gpu
 			];
 		};
 		stokes_drag_2d = pkgs.symlinkJoin {
 			name = "kel-lbm-sims-stokes_drag_2d-${version}";
 			paths = [
-				examples.stokes_drag_particle_2d_bgk_gpu
-				examples.stokes_drag_particle_2d_psm_gpu
-				examples.stokes_drag_particle_2d_hlbm_gpu
-				examples.stokes_drag_particle_2d_fplbm_gpu
+				ex.stokes_drag_particle_2d_bgk_gpu
+				ex.stokes_drag_particle_2d_psm_gpu
+				ex.stokes_drag_particle_2d_hlbm_gpu
+				ex.stokes_drag_particle_2d_fplbm_gpu
 			];
 		};
 
@@ -297,5 +297,5 @@ in rec {
 		};
 	};
 
-	default = release.examples;
+	default = release.ex;
 }
diff --git a/examples/one_particle_sedimentation_2d/common.hpp b/examples/one_particle_sedimentation_2d/common.hpp
index 37e0402..954557d 100644
--- a/examples/one_particle_sedimentation_2d/common.hpp
+++ b/examples/one_particle_sedimentation_2d/common.hpp
@@ -14,15 +14,17 @@ namespace kel {
 namespace lbm {
 
 /**
- * length 2.2m
- * height 0.41m
+ * length 20mm
+ * height 80mm
  * viscosity 10^-3
- * particle diameter 0.1m
- * position at (0.16m,0.2m)
+ * particle radius 1.25mm
+ * position at (10mm,40mm)
+ * particle density 0.00125g/mm³
+ * fluid density 0.001g/mm³
  */
 
-constexpr uint64_t dim_y = 512ul;
-constexpr uint64_t dim_x = 256ul;
+constexpr uint64_t dim_y = 800ul;
+constexpr uint64_t dim_x = 200ul;
 
 constexpr uint64_t particle_amount = 1ul;
 
diff --git a/examples/one_particle_sedimentation_2d/init.hpp b/examples/one_particle_sedimentation_2d/init.hpp
index a35c6e8..d9b8a37 100644
--- a/examples/one_particle_sedimentation_2d/init.hpp
+++ b/examples/one_particle_sedimentation_2d/init.hpp
@@ -87,8 +87,8 @@ saw::error_or<void> init(
 			auto& prb = p.template get<"rigid_body">();
 			auto& p_pos = prb.template get<"position">();
 			auto& p_posold = prb.template get<"position_old">();
-			p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.16}}).handle();
-			p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.2}}).handle();
+			p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.01}}).handle();
+			p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.04}}).handle();
 
 			p_posold = p_pos;
 		}
diff --git a/examples/one_particle_sedimentation_2d/sim.cpp b/examples/one_particle_sedimentation_2d/sim.cpp
index 7c9a56c..b0b374c 100644
--- a/examples/one_particle_sedimentation_2d/sim.cpp
+++ b/examples/one_particle_sedimentation_2d/sim.cpp
@@ -46,7 +46,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 
 	converter<T> conv {
 		// delta_x
-		{{0.41 / dim_y}},
+		{{0.02 / dim_x}},
 		// delta_t
 		{{1.067e-4}}
 	};
diff --git a/examples/one_particle_sedimentation_2d/step.hpp b/examples/one_particle_sedimentation_2d/step.hpp
index 07f51d1..20b7f4b 100644
--- a/examples/one_particle_sedimentation_2d/step.hpp
+++ b/examples/one_particle_sedimentation_2d/step.hpp
@@ -30,7 +30,6 @@ saw::error_or<void> step(
 				for(uint64_t i = 0u; i < Desc::D; ++i){
 					index.at({{i}}).set(idx[i]);
 				}
-
 				hlbm_reset.apply(fields,macros,index,t_i);
 			});
 		}).wait();
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
index b9b3778..59f61d9 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
@@ -48,7 +48,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 		// delta_x
 		{{0.41 / dim_y}},
 		// delta_t
-		{{1.067e-4}}
+		{{1.067e-4 * 0.5}}
 	};
 	
 	print_lbm_meta<T,Desc>(conv,{1e-3},{1e-4},{0.5 * dim_y});
@@ -122,7 +122,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 	auto lsdm_view = make_view(lbm_sycl_macro_data);
 	auto lsdp_view = make_view(lbm_sycl_particle_data);
 	
-	saw::data<sch::UInt64> time_steps{64u*1024ul};
+	saw::data<sch::UInt64> time_steps{128u*1024ul};
 
 	auto& info_f = lsd_view.template get<"info">();
 
@@ -136,6 +136,7 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 		}
 
 		sycl_q.wait();
+		/*
 		if(i.get() % 64u == 0u){
 			{
 				auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
@@ -144,12 +145,13 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 				}
 			}
 			{
-				auto eov = write_csv_file(out_dir,"stdkr_2d",i.get(), *lbm_macro_data_ptr);
+				auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+an.template get<"coupling">().stl_string(),i.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;
@@ -189,7 +191,12 @@ saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 			return eov;
 		}
 	}
-
+	{
+		auto eov = write_csv_file(out_dir,"stdkr_2d",time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 	sycl_q.wait();
 	return saw::make_void();
 }
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
index cbdfd10..beb123b 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
@@ -19,8 +19,6 @@ saw::error_or<void> step(
 	auto& q = dev.get_handle();
 	auto& info_f = fields.template get<"info">();
 	auto& porous_f = macros.template get<"porosity">();
-
-
 	
 	if constexpr ( std::is_same_v<Coll, method::Hlbm> ){
 		q.submit([&](acpp::sycl::handler& h){
@@ -48,7 +46,7 @@ saw::error_or<void> step(
 		}).wait();
 
 		q.submit([&](acpp::sycl::handler& h){
-			component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.75};
 			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
 
@@ -74,9 +72,19 @@ saw::error_or<void> step(
 							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
 								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
 									saw::data<sch::Vector<T,Desc::D>> vel;
+									saw::data<T> scale;
+									{
+										scale.set([&](){
+											if(t_i.get() > 1024u){
+												return 1.0f;
+											}
+
+											return t_i.get() / 1024.0f;
+										}());
+									}
 									{
 										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
-										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle() * scale;
 										vel.at({{1u}}) = 0.0f;
 									}
 									return vel;
@@ -123,8 +131,8 @@ saw::error_or<void> step(
 
 		// auto coll_ev = 
 		q.submit([&](acpp::sycl::handler& h){
-			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{1.0};
-			component<T,Desc,cmpt::FpLbm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{0.75};
+			component<T,Desc,cmpt::FpLbm,encode::Sycl<saw::encode::Native>> collision{0.75};
 			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
 
@@ -199,8 +207,8 @@ saw::error_or<void> step(
 
 		// auto coll_ev = 
 		q.submit([&](acpp::sycl::handler& h){
-			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{1.0};
-			component<T,Desc,cmpt::Psm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{0.75};
+			component<T,Desc,cmpt::Psm,encode::Sycl<saw::encode::Native>> collision{0.75};
 			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
 
diff --git a/modules/core/c++/abstract/schema.hpp b/modules/core/c++/abstract/schema.hpp
index 8bc23b4..89d9c2e 100644
--- a/modules/core/c++/abstract/schema.hpp
+++ b/modules/core/c++/abstract/schema.hpp
@@ -42,12 +42,6 @@ using Float16 = Primitive<FloatingPoint,2u>;
 using Float32 = Primitive<FloatingPoint,4u>;
 using Float64 = Primitive<FloatingPoint,8u>;
 
-template<typename Sch, uint64_t StorageBytes, uint64_t InterfaceBytes>
-struct MixedPrecision {
-	using Meta = Void;
-	using Inner = Sch;
-};
-
 template<typename T, uint64_t... Dims>
 struct FixedArray {
 	using Meta = Void;
diff --git a/modules/core/c++/abstract/stringify_schema.hpp b/modules/core/c++/abstract/stringify_schema.hpp
index 3d6b1d1..dd218d0 100644
--- a/modules/core/c++/abstract/stringify_schema.hpp
+++ b/modules/core/c++/abstract/stringify_schema.hpp
@@ -10,36 +10,36 @@ template<class T>
 struct schema_to_string_helper {
 	static_assert(always_false<T>, "schema_to_string_helper specialization not handled.");
 
-	static constexpr string_literal apply(){
-		return {};
+	static constexpr auto apply(){
+		return string_literal{""};
 	}
 };
 
 template<>
 struct schema_to_string_helper<sch::SignedInteger> final {
-	static constexpr string_literal apply(){
-		return "SignedInteger";
+	static constexpr auto apply(){
+		return string_literal{"SignedInteger"};
 	}
 };
 
 template<>
 struct schema_to_string_helper<sch::UnsignedInteger> final {
-	static constexpr string_literal apply(){
-		return "UnsignedInteger";
+	static constexpr auto apply(){
+		return string_literal{"UnsignedInteger"};
 	}
 };
 
 template<>
 struct schema_to_string_helper<sch::FloatingPoint> final {
-	static constexpr string_literal apply(){
-		return "FloatingPoint";
+	static constexpr auto apply(){
+		return string_literal{"FloatingPoint"};
 	}
 };
 
 template<typename T, uint64_t N>
 struct schema_to_string_helper<sch::Primitive<T,N>> final {
-	static constexpr string_literal apply(){
-		return "Primitive<"+schema_to_string_helper<T>::apply()+","+N+">";
+	static constexpr auto apply(){
+		return string_literal{"Primitive<"}+schema_to_string_helper<T>::apply()+string_literal{","}+N+string_literal{">"};
 	}
 };
 }
diff --git a/modules/core/c++/fplbm.hpp b/modules/core/c++/fplbm.hpp
index 382d067..d8c075b 100644
--- a/modules/core/c++/fplbm.hpp
+++ b/modules/core/c++/fplbm.hpp
@@ -92,7 +92,7 @@ public:
 		auto& force_f = macros.template get<"force">();
 		auto& force = force_f.at(index);
 
-		vel = vel + force * half;
+		vel = vel + force * half / rho;
 
 		// compute_rho_u<T,Descriptor>(dfs,rho,vel);
 		auto eq = equilibrium<T,Descriptor>(rho,vel);
@@ -204,119 +204,18 @@ public:
 
 			// vel_s is technically time the density of the particle?
 
-			force = ( vel_s * rho - vel * rho ) /* two */ * flip_por;
+			force = ( vel_s * rho - vel * rho ) * two * flip_por;
 
 			force_p = force_p - force;
 		}, aabb.template get<"a">(), aabb.template get<"b">());
 
 		auto& pirb_acc = pirb.template get<"acceleration">();
-		pirb_acc = force_p / pg.template get<"total_mass">().at({});
+		pirb_acc = force_p / (two * pg.template get<"total_mass">().at({}));
 
 		for(saw::data<sch::UInt64> i{0u}; i < sub_steps; ++i){
 			verlet_step_lambda<T,Descriptor::D>(pi,sss);
 		}
 	}
 };
-
-template<typename T, typename Descriptor, typename Encode>
-class component<T, Descriptor, cmpt::FpLbmOneParticleNoVelocity, Encode> final {
-private:
-	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_{saw::data<T>{1} / relaxation_}
-	{}
-
-	using Component = cmpt::FpLbmOneParticleNoVelocity;
-
-	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);
-
-		auto& rho_f = macros.template get<"density">();
-		auto& vel_f = macros.template get<"momentum">();
-		auto& por_f = macros.template get<"porosity">();
-
-		// Temporary find a better way
-		auto& force_f = macros.template get<"force">();
-
-
-		/**
-		 * The other methods all work with a flipped porosity.
-		 * For compat reasons this is also flipped
-		 */
-		auto porosity = por_f.at(index);
-		saw::data<sch::Scalar<T>> one;
-		one.at({}) = 1.0;
-		auto flip_porosity = one - porosity;
-
-		saw::data<sch::Scalar<T>>& rho = rho_f.at(index);
-
-		saw::data<sch::Scalar<T>> half;
-		half.at({}).set(0.5);
-		saw::data<sch::Vector<T,Descriptor::D>>& vel = vel_f.at(index);// + total_force * ( half / rho );
-
-		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>;
-
-		saw::data<sch::Scalar<T>> min_two;
-		min_two.at({}).set(-2);
-
-		auto& force = force_f.at(index);
-		force = vel * rho * min_two * flip_porosity;
-
-		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<sch::Vector<T,Descriptor::D>> ci;
-			for(uint64_t d = 0u; d < Descriptor::D; ++d){
-				ci.at({{d}}).set(static_cast<typename saw::native_data_type<T>::type>(dfi::directions[i][d]));
-			}
-			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]);
-
-			/*
-			saw::data<sch::Scalar<T>> F_i_sum;
-			for(uint64_t d = 0u; d < Descriptor::D; ++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;
-			}
-			*/
-			auto term1 = (ci-vel) * dfi_inv_cs2;
-			auto term2 = ci * (ci_dot_u * dfi_inv_cs2 * dfi_inv_cs2);
-
-			auto force_projection = saw::math::dot(term1 + term2, force);
-
-			auto F_i = w * force_projection;
-
-			dfs.at({i}) = dfs.at({i})
-									+ frequency_
-										* (eq.at(i) - dfs.at({i}) )
-									+ F_i.at({})
-										* (saw::data<T>{1} - saw::data<T>{0.5f} * frequency_);
-		}
-	}
-
-};
 }
 }
diff --git a/modules/core/c++/macroscopic.hpp b/modules/core/c++/macroscopic.hpp
index 768fef0..9db1538 100644
--- a/modules/core/c++/macroscopic.hpp
+++ b/modules/core/c++/macroscopic.hpp
@@ -33,5 +33,30 @@ void compute_rho_u (
 		vel().at({{i}}) = vel().at({{i}}) / rho().at({});
 	}
 }
+
+/**
+ * Calculate the macroscopic variables rho and u in Lattice Units.
+ */
+template<typename T, typename Desc>
+void compute_rho_momentum (
+		const saw::data<sch::FixedArray<T,Desc::Q>>& dfs,
+		saw::ref<saw::data<sch::Scalar<T>>> rho,
+		saw::ref<saw::data<sch::Vector<T,Desc::D>>> mom
+	)
+{
+	using dfi = df_info<T, Desc>;
+
+	rho().at({}).set(0);
+	for(uint64_t i = 0; i < Desc::D; ++i){
+		mom().at({{i}}).set(0);
+	}
+
+	for(size_t j = 0; j < Desc::Q; ++j){
+		rho().at({}) = rho().at({}) + dfs.at({j});
+		for(size_t i = 0; i < Desc::D; ++i){
+			mom().at({{i}}) = mom().at({{i}}) + saw::data<T>{static_cast<typename saw::native_data_type<T>::type>(dfi::directions[j][i])} * dfs.at({j});
+		}
+	}
+}
 }
 }
diff --git a/modules/core/tests/schema.cpp b/modules/core/tests/schema.cpp
index db82b97..6649bad 100644
--- a/modules/core/tests/schema.cpp
+++ b/modules/core/tests/schema.cpp
@@ -18,6 +18,7 @@ SAW_TEST("Abstract/Schema/Ptr"){
 }
 */
 
+/*
 SAW_TEST("Abstract/SchemaStringify/PrimitiveKind"){
 	using namespace kel::lbm;
 
@@ -25,4 +26,5 @@ SAW_TEST("Abstract/SchemaStringify/PrimitiveKind"){
 
 	static constexpr auto str = schema_to_string<TSch>::value;
 }
+*/
 }