Fixed and added Array for sycl

8 files changed, 75 insertions, 38 deletions

diff --git a/default.nix b/default.nix
index 011ca29..ee3c963 100644
--- a/default.nix
+++ b/default.nix
@@ -87,7 +87,7 @@ in rec {
 			build_examples = "true";
     };
 	 
-	 io_codec = pkgs.callPackage modules/io_codec/.nix/derivation.nix {
+		io_codec = pkgs.callPackage modules/io_codec/.nix/derivation.nix {
       inherit version;
       inherit forstio;
       inherit stdenv;
@@ -97,10 +97,10 @@ in rec {
 		};
 		
 		remote-thread = pkgs.callPackage modules/remote-thread/.nix/derivation.nix {
-      inherit version;
-      inherit forstio;
-      inherit stdenv;
-      inherit clang-tools;
+			inherit version;
+			inherit forstio;
+			inherit stdenv;
+			inherit clang-tools;
 		};
 		
 		remote = pkgs.callPackage modules/remote/.nix/derivation.nix {
@@ -128,7 +128,7 @@ in rec {
 			inherit clang-tools;
 			openmp = pkgs.llvmPackages_17.openmp;
 			
-			build_examples = "true";
+			build_examples = "false";
 			build_benchmarks = "true";
 		};
 		
diff --git a/modules/remote-sycl/.nix/derivation.nix b/modules/remote-sycl/.nix/derivation.nix
index 33317eb..bf40e89 100644
--- a/modules/remote-sycl/.nix/derivation.nix
+++ b/modules/remote-sycl/.nix/derivation.nix
@@ -11,7 +11,7 @@
 , python3
 , bash
 
-, build_examples ? "true"
+, build_examples ? "false"
 , build_benchmarks ? "true"
 }:
 
@@ -54,12 +54,12 @@ in stdenv.mkDerivation {
 		scons prefix=$out build_benchmarks=${build_benchmarks} build_examples=${build_examples} install
 	'';
 
-  doCheck = true;
+	doCheck = true;
 	checkPhase = ''
 		export ACPP_APPDB_DIR=.
-    scons test
-    ./bin/tests
-  '';
+		scons test
+		./bin/tests
+	'';
 
-  outputs = ["out" "dev"];
+	outputs = ["out" "dev"];
 }
diff --git a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
index c71582e..b939374 100644
--- a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
@@ -13,14 +13,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Sycl<saw::enc
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
 					
-					saw::data<sch::Float64> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float64> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo.get() == 1.1e12 ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo + foo * saw::data<sch::Float64>{1.7342345};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
@@ -31,14 +31,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Sycl<saw::enc
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
-					saw::data<sch::Float64> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float64> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo == saw::data<sch::Float64>{1.1e12} ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo +foo * saw::data<sch::Float64>{1.7342345};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
@@ -49,14 +49,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Sycl<saw::enc
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
-					saw::data<sch::Float32> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float32> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo == saw::data<sch::Float32>{1.1e12f} ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo + foo * saw::data<sch::Float32>{1.7342345f};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
@@ -74,14 +74,14 @@ saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::MixedArray>, saw::enco
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
 
-					saw::data<sch::Float64> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float64> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo.get() == 1.1e12 ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo + foo * saw::data<sch::Float64>{1.7342345};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
@@ -92,18 +92,18 @@ saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float32Array>, saw::en
 	return {
 		[&](saw::data<sch::Float32Array, saw::encode::Sycl<saw::encode::Native>>& in, cl::sycl::queue* cmd) -> saw::error_or<void> {
 			auto in_size = in.size();
-			float32_ev = cmd->submit([&](cl::sycl::handler& h){
+			ev = cmd->submit([&](cl::sycl::handler& h){
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
-					saw::data<sch::Float32> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float32> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo == saw::data<sch::Float32>{1.1e12f} ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo + foo * saw::data<sch::Float32>{1.7342345f};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
@@ -115,18 +115,18 @@ saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float64Array>, saw::en
 	return {
 		[&](saw::data<sch::Float64Array, saw::encode::Sycl<saw::encode::Native>>& in, cl::sycl::queue* cmd) -> saw::error_or<void> {
 			auto in_size = in.size();
-			float64_ev = cmd->submit([&](cl::sycl::handler& h){
+			ev = cmd->submit([&](cl::sycl::handler& h){
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size.get()), [=] (cl::sycl::id<1> it){
-					saw::data<sch::Float64> foo = {acc_buff[0u].at(it[0u]).get()};
+					saw::data<sch::Float64> foo = {acc_buff[it[0u]].get()};
 					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
 						if( foo == saw::data<sch::Float64>{1.1e12} ){
-							acc_buff[0u].at(it[0u]) = 0.f;
+							acc_buff[it[0u]] = 0.f;
 						}
 						foo = foo +foo * saw::data<sch::Float64>{1.7342345};
 					}
-					acc_buff[0u].at(it[0u]) = foo;
+					acc_buff[it[0u]] = foo;
 				});
 			});
 			return saw::void_t{};
diff --git a/modules/remote-sycl/benchmarks/mixed_precision.cpp b/modules/remote-sycl/benchmarks/mixed_precision.cpp
index d4b119b..4d37f5f 100644
--- a/modules/remote-sycl/benchmarks/mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/mixed_precision.cpp
@@ -140,7 +140,7 @@ int main(int argc, char** argv){
 		uint64_t time_float32 = std::numeric_limits<uint64_t>::max();
 		for(uint64_t runs_i = 0u; runs_i < runs; ++runs_i){
 
-			(std::cout<<'.').flush();
+			(std::cout<<''.').flush();
 
 			data<sch::MixedArray> mixed_host_data;
 			data<sch::Float64Array> float64_host_data;
diff --git a/modules/remote-sycl/c++/common.hpp b/modules/remote-sycl/c++/common.hpp
index 3859dad..822dae5 100644
--- a/modules/remote-sycl/c++/common.hpp
+++ b/modules/remote-sycl/c++/common.hpp
@@ -2,6 +2,7 @@
 
 #include <forstio/remote/remote.hpp>
 #include <forstio/codec/data.hpp>
+#include <forstio/codec/schema_hash.hpp>
 #include <forstio/codec/id_map.hpp>
 
 #include <AdaptiveCpp/CL/sycl.hpp>
diff --git a/modules/remote-sycl/c++/data.hpp b/modules/remote-sycl/c++/data.hpp
index 7763952..91f74f8 100644
--- a/modules/remote-sycl/c++/data.hpp
+++ b/modules/remote-sycl/c++/data.hpp
@@ -41,4 +41,40 @@ public:
 		return data_.template get_access<AccessMode>(h);
 	}
 };
+
+template<typename Sch, uint64_t Dim>
+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_;
+	data<schema::UInt64, encode::Native> size_;
+public:
+	data(const data<Schema, encode::Native>& host_data__):
+		data_{&host_data__.at({0u}),host_data__.size().get()},
+		size_{host_data__.size()}
+	{}
+
+	auto& get_handle() {
+		return data_;
+	}
+
+	const auto& get_handle() const {
+		return data_;
+	}
+
+	data<schema::UInt64, encode::Native> size() const {
+		return size_;
+	}
+
+	template<cl::sycl::access::mode AccessMode>
+	auto access(cl::sycl::handler& h){
+		return data_.template get_access<AccessMode>(h);
+	}
+
+	template<cl::sycl::access::mode AccessMode>
+	auto access(cl::sycl::handler& h) const {
+		return data_.template get_access<AccessMode>(h);
+	}
+};
 }
diff --git a/modules/remote-sycl/c++/remote.hpp b/modules/remote-sycl/c++/remote.hpp
index fd2f64a..65f645e 100644
--- a/modules/remote-sycl/c++/remote.hpp
+++ b/modules/remote-sycl/c++/remote.hpp
@@ -31,9 +31,9 @@ private:
 		std::array<uint64_t,3> data;
 
 		template<typename Schema, typename Encoding>
-		static key_t create(const remote_address<rmt::Loopback>& addr){
+		static key_t create(const remote_address<rmt::Sycl>& addr){
 			key_t k;
-			k.data = std::array<uint64_t,3>{addr.get_address_id().get(), schema_hash<Schema>::apply(), schema_hash<Encoding>::apply()};
+			k.data = std::array<uint64_t,3>{addr.get_device_id(), schema_hash<Schema>::apply(), schema_hash<Encoding>::apply()};
 
 			return k;
 		}
diff --git a/modules/remote-sycl/tests/mixed_precision.cpp b/modules/remote-sycl/tests/mixed_precision.cpp
index 42563fd..4a5218d 100644
--- a/modules/remote-sycl/tests/mixed_precision.cpp
+++ b/modules/remote-sycl/tests/mixed_precision.cpp
@@ -76,7 +76,7 @@ SAW_TEST("SYCL Mixed Test"){
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(test_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * data<schema::Float64>{2.0};
+					acc_buff[it[0u]] = acc_buff[it[0u]] * data<schema::Float64>{2.0};
 				});
 			});
 			return saw::void_t{};
@@ -135,7 +135,7 @@ SAW_TEST("SYCL Float Test"){
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(test_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * data<schema::Float32>{2.0};
+					acc_buff[it[0u]] = acc_buff[it[0u]] * data<schema::Float32>{2.0};
 				});
 			});
 			return saw::void_t{};
@@ -194,7 +194,7 @@ SAW_TEST("SYCL Double Test"){
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(test_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * data<schema::Float64>{2.0};
+					acc_buff[it[0u]] = acc_buff[it[0u]] * data<schema::Float64>{2.0};
 				});
 			});
 			return saw::void_t{};