Double free, but constexpr things compile nowdev

4 files changed, 71 insertions, 63 deletions

diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index bb1fca5..4a28541 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -8,7 +8,7 @@ namespace kel {
 namespace lbm {
 
 constexpr uint64_t dim_x = 256u;
-constexpr uint64_t dim_y = 16u;
+constexpr uint64_t dim_y = 32u;
 
 namespace sch {
 using namespace saw::schema;
@@ -20,8 +20,9 @@ using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>;
 
 template<typename T, typename Desc>
 using ChunkStruct = Struct<
-	Member<InfoChunk, "info">
-	//,Member<DfChunk<T,Desc>, "dfs">
+	Member<InfoChunk, "info">,
+	Member<DfChunk<T,Desc>, "dfs">,
+	Member<DfChunk<T,Desc>, "dfs_old">
 >;
 }
 
@@ -66,7 +67,6 @@ saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>
 		{{dim_x, dim_y}},
 		{{0u,1u}}
 	);
-/*
 	//
 	auto& df_f = fields.template get<"dfs">();
 	saw::data<T> rho{1};
@@ -82,7 +82,6 @@ saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>
 		{},// 0-index
 		df_f.get_dims()
 	);
-*/
 
 	return saw::make_void();
 }
@@ -90,6 +89,19 @@ saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>
 template<typename T, typename Desc>
 saw::error_or<void> step(saw::data<sch::ChunkStruct<T,Desc>,encode::Sycl<saw::encode::Native>>& fields, saw::data<sch::UInt64> t_i, device& dev){
 	auto& q = dev.get_handle();
+	auto& info_f = fields.template get<"info">();
+
+	q.submit([&](acpp::sycl::handler& h){
+		component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> collision{0.6};
+		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;
+			for(uint64_t i = 0u; i < Desc::D; ++i){
+				index.at({{i}}).set(idx[i]);
+			}
+
+			collision.apply(fields,index,t_i);
+		});
+	}).wait();
 
 	return saw::make_void();
 }
@@ -132,11 +144,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	}
 	{
 
-	std::cout<<"Hey"<<std::endl;
-
 	saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q};
 	sycl_q.wait();
-	std::cout<<"Hey2"<<std::endl;
 	{
 		auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data);
 		if(eov.is_error()){
@@ -144,7 +153,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 		}
 	}
 	sycl_q.wait();
-	std::cout<<"Hey3"<<std::endl;
 
 	for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{32ul}; ++i){
 		auto eov = step<T,Desc>(lbm_sycl_data,i,dev);
@@ -155,7 +163,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	}
 	sycl_q.wait();
 
-	std::cout<<"Hey4"<<std::endl;
 	/*
 	iterator<Desc::D>::apply(
 		[&](auto& index){
diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp
index 2f20601..85c3af9 100644
--- a/lib/core/c++/collision.hpp
+++ b/lib/core/c++/collision.hpp
@@ -45,43 +45,19 @@ public:
 	 * Raw setup
 	 */
 	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){
+	void apply(const saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step) const {
 		bool is_even = ((time_step.get() % 2) == 0);
-		auto& cell = field.at(index);
 
-		auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-		// auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
+		// auto& dfs_f = (is_even) ? field.template get<"dfs">() : field.template get<"dfs_old">();
+		auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
 
 		saw::data<T> rho;
 		saw::data<sch::FixedArray<T,Descriptor::D>> vel;
-		compute_rho_u<T,Descriptor>(dfs_old,rho,vel);
-		auto eq = equilibrium<T,Descriptor>(rho,vel);
-
-		for(uint64_t i = 0u; i < Descriptor::Q; ++i){
-			dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}));
-			// dfs_old({i}).set(dfs_old({i}).get() + (1.0 / relaxation_) * (eq.at(i).get() - dfs_old({i}).get()));
-		}
-	}
-
-	template<typename CellStructSchema>
-	void apply(
-		saw::data<CellStructSchema, Encode>* field, 
-		const saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>>& meta,
-		const saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>>& index,
-		saw::data<sch::UInt64> time_step
-	){
-		bool is_even = ((time_step.get() % 2) == 0);
-		auto& cell = field[0u];
-		
-		auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-
-		saw::data<T> rho;
-		saw::data<sch::FixedArray<T,Descriptor::D>> vel;
-		compute_rho_u<T,Descriptor>(dfs_old,rho,vel);
+		compute_rho_u<T,Descriptor>(dfs_old_f.at(index),rho,vel);
 		auto eq = equilibrium<T,Descriptor>(rho,vel);
 
 		for(uint64_t i = 0u; i < Descriptor::Q; ++i){
-			dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}));
+			dfs_old_f.at(index).at({i}) = dfs_old_f.at(index).at({i}) + frequency_ * (eq.at(i) - dfs_old_f.at(index).at({i}));
 		}
 	}
 };
diff --git a/lib/core/c++/macroscopic.hpp b/lib/core/c++/macroscopic.hpp
index 51368e9..1532873 100644
--- a/lib/core/c++/macroscopic.hpp
+++ b/lib/core/c++/macroscopic.hpp
@@ -4,6 +4,31 @@
 
 namespace kel {
 namespace lbm {
+template<typename T, typename Desc>
+void compute_rho_u (
+		const saw::data<sch::FixedArray<T,Desc::Q>>& dfs,
+		saw::data<T>& rho,
+		saw::data<sch::FixedArray<T,Desc::D>>& vel
+	)
+{
+	using dfi = df_info<T, Desc>;
+
+	rho = 0;
+	for(size_t j = 0; j < Desc::D; ++j){
+		vel.at({{j}}).set(0);
+	}
+
+	for(size_t j = 0; j < Desc::Q; ++j){
+		rho = rho + dfs.at({j});
+		for(size_t i = 0; i < Desc::D; ++i){
+			vel.at({{i}}) = vel.at({{i}}) + dfi::directions[j][i] * dfs.at({j}).get();
+		}
+	}
+
+	for(size_t i = 0; i < Desc::D; ++i){
+		vel.at({i}) = vel.at({i}) / rho;
+	}
+}
 /**
  * Calculate the macroscopic variables rho and u in Lattice Units.
  */
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
index cffeb38..d5adb95 100644
--- a/lib/sycl/c++/data.hpp
+++ b/lib/sycl/c++/data.hpp
@@ -25,15 +25,14 @@ public:
 	using Schema = schema::FixedArray<Sch,Dims...>;
 private:
 	acpp::sycl::queue* q_;
-	data<Sch>* values_;
+	data<Sch,Encode>* values_;
 
-	SAW_FORBID_COPY(data);
+	//SAW_FORBID_COPY(data);
 public:
 	data(acpp::sycl::queue& q__):
 		q_{&q__},
 		values_{nullptr}
 	{
-		std::cout<<"Hey: "<<ct_multiply<uint64_t,Dims...>::value<<std::endl;
 		values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_);
 		SAW_ASSERT(values_ and q_);
 	}
@@ -49,14 +48,18 @@ public:
 	}
 
 	static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() {
-		return {std::array<uint64_t, sizeof...(Dims)>{Dims...}};
+		return saw::data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>>{{Dims...}};
 	}
 
-	data<Sch>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
+	constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
 		return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
 	}
 	
-	constexpr data<Sch,encode::Native>* flat_data() {
+	constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index) const{
+		return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+	}
+
+	constexpr data<Sch,Encode>* flat_data() {
 		return values_;
 	}
 };
@@ -79,7 +82,7 @@ public:
 		return values_.at(index);
 	}
 
-	const data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+	data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
 		return values_.at(index);
 	}
 
@@ -87,7 +90,7 @@ public:
 		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
 	}
 
-	data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+	data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
 		std::decay_t<decltype(index)> ind;
 		for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
 			ind.at({i}) = index.at({i}) + Ghost;
@@ -95,7 +98,7 @@ public:
 		return values_.at(ind);
 	}
 
-	const data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+	data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
 		std::decay_t<decltype(index)> ind;
 		for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
 			ind.at({i}) = index.at({i}) + Ghost;
@@ -116,12 +119,11 @@ public:
 	}
 };
 
-template<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
-class data<schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode> > final {
+template<typename... Members, typename Encode>
+class data<schema::Struct<Members...>, kel::lbm::encode::Sycl<Encode> > final {
 public:
-	static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}};
-	using StorageT = std::tuple<data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>...>;
-	using Schema = schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>;
+	using StorageT = std::tuple<data<typename Members::ValueType,kel::lbm::encode::Sycl<Encode>>...>;
+	using Schema = schema::Struct<Members...>;
 private:
 
 	/**
@@ -133,24 +135,24 @@ private:
 	template<std::size_t... Is>
 	constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>):
 		members_{(static_cast<void>(Is), q)...}
-	{
-
-	}
+	{}
 public:
 	data(acpp::sycl::queue& q__):
-		data{q__, std::make_index_sequence<sizeof...(Sch)>{}}
-//		members_(q__)//data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>(q__)...)
+		data{q__, std::make_index_sequence<sizeof...(Members)>{}}
 	{
 		q__.wait();
 	}
 
 	template<saw::string_literal K>
 	auto& get(){
-		return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
+		return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
+	}
+	
+	template<saw::string_literal K>
+	const auto& get() const {
+		return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
 	}
 };
-
-
 }
 
 namespace kel {
@@ -175,8 +177,6 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final {
 
 			static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
 
-			std::cout<<host_member_data.flat_size().get()<<" "<<std::endl;
-
 			q.submit([&](acpp::sycl::handler& h){
 				h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get());
 			}).wait();