Going forward with testing sycl setup

3 files changed, 62 insertions, 84 deletions

diff --git a/default.nix b/default.nix
index b6f8f5b..bd3182a 100644
--- a/default.nix
+++ b/default.nix
@@ -12,10 +12,10 @@ let
     inherit clang-tools;
 	}).forstio;
 
-	adaptive-cpp = pkgs.callPackage ./.nix/adaptive-cpp.nix {
-			inherit stdenv;
-			llvmPackages = pkgs.llvmPackages_19;
-			lld = pkgs.lld_19;
+	adaptive-cpp-dev = pkgs.callPackage .nix/adaptive-cpp.nix {
+		inherit stdenv;
+		llvmPackages = pkgs.llvmPackages_19;
+		lld = pkgs.lld_19;
 	};
 
 	pname = "kel-lbm";
@@ -29,7 +29,7 @@ in rec {
 
 	examples = {
 		cavity_2d_gpu = pkgs.callPackage ./examples/cavity_2d_gpu/.nix/derivation.nix {
-			inherit pname version stdenv forstio adaptive-cpp;
+			inherit pname version stdenv forstio adaptive-cpp-dev;
 			kel-lbm = lbm;
 		};
 
@@ -45,6 +45,6 @@ in rec {
 	};
 
 	debug = {
-		inherit adaptive-cpp;
+		# inherit foo;
 	};
 }
diff --git a/examples/cavity_2d_gpu/.nix/derivation.nix b/examples/cavity_2d_gpu/.nix/derivation.nix
index 73b81b7..2748a18 100644
--- a/examples/cavity_2d_gpu/.nix/derivation.nix
+++ b/examples/cavity_2d_gpu/.nix/derivation.nix
@@ -3,9 +3,10 @@
 , scons
 , clang-tools
 , forstio
-, adaptive-cpp
+, keldu
 , pname
 , version
+, adaptive-cpp-dev
 , kel-lbm
 }:
 
@@ -20,14 +21,13 @@ stdenv.mkDerivation {
   ];
 
   buildInputs = [
-    forstio.core
-    forstio.async
-    forstio.codec
-    forstio.codec-unit
-    forstio.remote
-    forstio.remote-sycl
+		forstio.core
+		forstio.async
+		forstio.codec
+		forstio.codec-unit
+		forstio.remote
 		forstio.codec-json
-		adaptive-cpp
+		adaptive-cpp-dev
 		kel-lbm
   ];
 
diff --git a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
index c4836d0..ed7042f 100644
--- a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
+++ b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
@@ -1,7 +1,6 @@
 #include <kel/lbm/lbm.hpp>
 
 #include <forstio/codec/data.hpp>
-
 #include <AdaptiveCpp/sycl/sycl.hpp>
 
 #include <iostream>
@@ -17,7 +16,7 @@ struct Sycl {
 template<typename Schema>
 class data<Schema, encode::Sycl<encode::Native>> {
 private:
-	cl::sycl::buffer<data<Schema, encode::Native>> data_;
+	acpp::sycl::buffer<data<Schema, encode::Native>> data_;
 	data<schema::UInt64, encode::Native> size_;
 public:
 	data(const data<Schema, encode::Native>& data__):
@@ -37,13 +36,13 @@ public:
 		return size_;
 	}
 
-	template<cl::sycl::access::mode AccessMode>
-	auto access(cl::sycl::handler& h){
+	template<acpp::sycl::access::mode AccessMode>
+	auto access(acpp::sycl::handler& h){
 		return data_.template get_access<AccessMode>(h);
 	}
 
-	template<cl::sycl::access::mode AccessMode>
-	auto access(cl::sycl::handler& h) const {
+	template<acpp::sycl::access::mode AccessMode>
+	auto access(acpp::sycl::handler& h) const {
 		return data_.template get_access<AccessMode>(h);
 	}
 };
@@ -53,7 +52,7 @@ class data<schema::Array<Sch, Dim>, encode::Sycl<encode::Native>> {
 public:
 	using Schema = schema::Array<Sch,Dim>;
 private:
-	cl::sycl::buffer<data<Sch, encode::Native>> data_;
+	acpp::sycl::buffer<data<Sch, encode::Native>> data_;
 	data<schema::UInt64, encode::Native> size_;
 public:
 	data(const data<Schema, encode::Native>& host_data__):
@@ -73,13 +72,13 @@ public:
 		return size_;
 	}
 
-	template<cl::sycl::access::mode AccessMode>
-	auto access(cl::sycl::handler& h){
+	template<acpp::sycl::access::mode AccessMode>
+	auto access(acpp::sycl::handler& h){
 		return data_.template get_access<AccessMode>(h);
 	}
 
-	template<cl::sycl::access::mode AccessMode>
-	auto access(cl::sycl::handler& h) const {
+	template<acpp::sycl::access::mode AccessMode>
+	auto access(acpp::sycl::handler& h) const {
 		return data_.template get_access<AccessMode>(h);
 	}
 };
@@ -184,7 +183,6 @@ public:
 	}
 };
 
-constexpr size_t dim_size = 2;
 constexpr size_t dim_x = 128;
 constexpr size_t dim_y = 128;
 
@@ -227,7 +225,6 @@ void set_geometry(saw::data<sch::CavityFieldD2Q9>& lattice){
 	 */
 	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
 		auto& info = info_field.at(index);
-
 		info({0u}).set(3u);
 
 	}, {{0u,1u}}, {{meta.at({0u}), 2u}}, {{2u,0u}});
@@ -265,26 +262,29 @@ struct sycl_component_context {
 	acpp::sycl::handler* handler = nullptr;
 };
 
+template<typename T, typename Desc>
 void lbm_step(
-	saw::data<sch::CellField<sch::D2Q9,sch::CellInfo<sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>>& info,
-	saw::data<sch::CellField<sch::D2Q9,kel::lbm::sch::DfCell  <sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>>& dfs,
-	saw::data<sch::CellField<sch::D2Q9,kel::lbm::sch::DfCell  <sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>>& dfs_old,
-	bool even_step,
+	acpp::sycl::buffer<saw::data<sch::CellInfo<Desc>>>& info,
+	acpp::sycl::buffer<saw::data<sch::DfCell<Desc>>>& dfs,
+	acpp::sycl::buffer<saw::data<sch::DfCell<Desc>>>& dfs_old,
+	bool is_even,
 	uint64_t time_step,
 	acpp::sycl::queue& sycl_q
 ){
   using namespace kel::lbm;
   using namespace acpp;
 
-  using dfi = df_info<sch::T, sch::D2Q9>;
+  using dfi = df_info<T,Desc>;
 
-  //component<sch::T, sch::D2Q9, cmpt::BGK> coll{0.59};
-  //component<sch::T, sch::D2Q9, cmpt::BounceBack> bb;
-  component<sch::T, sch::D2Q9, cmpt::MovingWall> bb_lid;
+  //component<T, Desc, cmpt::BGK> coll{0.59};
+  //component<T, Desc, cmpt::BounceBack> bb;
+  component<T, Desc, cmpt::MovingWall> bb_lid;
   bb_lid.lid_vel = {0.1, 0.0};
 
+  /*
   const size_t Nx = latt.template get_dim_size<0>().get();
   const size_t Ny = latt.template get_dim_size<1>().get();
+  */
 
   // Submit collision kernel
   sycl_q.submit([&](sycl::handler& cgh) {
@@ -294,67 +294,45 @@ void lbm_step(
     auto dfs_old_acc = dfs_old.template get_access<sycl::access::mode::read_write>(cgh);
 
     cgh.parallel_for(
-      sycl::range<2>{Nx, Ny}, [=](sycl::id<2> idx) {
+      sycl::range<2>{dim_x, dim_y}, [=](sycl::id<2> idx) {
         size_t i = idx[0];
         size_t j = idx[1];
 
+		  size_t acc_id = i * dim_x + j;
+
         // Read cell info
-        auto& info = info_acc[idx];
+        auto& info = info_acc[acc_id];
 
         switch (info({0u}).get()) {
           case 1u: {
             // bb.apply(latt_acc, {i, j}, time_step);
-				// bool is_even = ((time_step % 2) == 0);
 
-				auto& dfs_old = (is_even) ? dfs_old_acc[idx] : dfs_acc[idx];
-				// auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
+						auto& dfs_old = (is_even) ? dfs_old_acc[acc_id] : dfs_acc[acc_id];
+						// auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
+						auto df_cpy = dfs_old.copy();
 
-				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 = 1u; i < Desc::Q; ++i){
+							dfs_old({i}) = df_cpy({dfi::opposite_index.at(i)});
+						}
 
-				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()));
-				}
             break;
           }
           case 2u: {
             // coll.apply(latt_acc, {i, j}, time_step);
-				// bool is_even = ((time_step % 2) == 0);
 
-				auto& dfs_old = (is_even) ? dfs_old_acc[idx] : dfs_acc[idx];
-				// auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
+						auto& dfs_old = (is_even) ? dfs_old_acc[acc_id] : dfs_acc[acc_id];
 
-				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);
+						saw::data<T> rho;
+						saw::data<sch::FixedArray<T,Desc::D>> vel;
+						compute_rho_u<T,Desc>(dfs_old,rho,vel);
+						auto eq = equilibrium<T,Desc>(rho,vel);
 
-				using dfi = df_info<T,Descriptor>;
-
-				auto& force = cell.template get<"force">();
-
-				for(uint64_t i = 0u; i < Descriptor::Q; ++i){
-					// saw::data<T> ci_min_u{0};
-					saw::data<T> ci_dot_u{0};
-					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])};
+						for(uint64_t i = 0u; i < Desc::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()));
+						}
+						break;
 					}
-					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));
-
-					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({d});
-					}
-
-
-					dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}) ) + F_i * (saw::data<T>{1} - saw::data<T>{0.5f} * frequency_);
-				}
-            break;
-          }
           default:
             // Do nothing
             break;
@@ -375,7 +353,7 @@ void lbm_step(
         auto& info_new = cell.template get<"info">();
 
         if (info_new({0u}).get() > 0u && info_new({0u}).get() != 3u) {
-          for (uint64_t k = 0u; k < sch::D2Q9::Q; ++k) {
+          for (uint64_t k = 0u; k < Desc::Q; ++k) {
             auto dir = dfi::directions[dfi::opposite_index[k]];
             auto& cell_dir_old = latt_acc({i + dir[0], j + dir[1]});
 
@@ -401,8 +379,8 @@ void lbm_step(
 int main(){
 	using namespace kel::lbm;
 
-	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> dim{{dim_x,info_field.at dim_y}};
-	const dim_size = dim_x * dim_y;
+	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> dim{{dim_x, dim_y}};
+	constexpr size_t dim_size = dim_x * dim_y;
 
 	saw::data<sch::CavityFieldD2Q9>> lattice{dim};
 
@@ -435,9 +413,9 @@ int main(){
 	auto& dfs_old_field = lattice.template get<"dfs_old">();
 
 	acpp::sycl::queue sycl_q{acpp::sycl::default_selector_v, acpp::sycl::property::queue::in_order{}};
-	saw::data<sch::CellField<Desc,sch::CellInfo<sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>> sycl_info_field{info_field};
-	saw::data<sch::CellField<Desc,sch::DfCell<sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>> sycl_dfs_field{dfs_field};
-	saw::data<sch::CellField<Desc,sch::DfCell<sch::D2Q9>>, saw::encode::Sycl<saw::encode::Native>> sycl_dfs_old_field{dfs_old_field};
+	acpp::sycl::buffer<saw::data<sch::CellInfo<sch::D2Q9>>> info_sycl{info_field.internal_data(), info_field.internal_size()};
+	acpp::sycl::buffer<saw::data<sch::DfCell<sch::D2Q9>>> dfs_sycl{dfs_field.internal_data(), dfs_field.internal_size()};
+	acpp::sycl::buffer<saw::data<sch::DfCell<sch::D2Q9>>> dfs_old_sycl{dfs_old_field.internal_data(), dfs_old_field.internal_size()};
 
 	/**
 	 * Timeloop
@@ -458,7 +436,7 @@ int main(){
 			write_vtk_file(vtk_f_name, macros);
 		}
 
-		lbm_step(sycl_info_field, sycl_dfs_field, sycl_dfs_old_field, (i%2u == 0u), i, sycl_q);
+		lbm_step<sch::T,sch::D2Q9>(info_sycl, dfs_sycl, dfs_old_sycl, (i%2u == 0u), i, sycl_q);
 	}
 	return 0;
 }