Adding possible diverge

3 files changed, 93 insertions, 0 deletions

diff --git a/modules/codec-json/tests/codec-json.cpp b/modules/codec-json/tests/codec-json.cpp
index c5f08ec..ab33f5a 100644
--- a/modules/codec-json/tests/codec-json.cpp
+++ b/modules/codec-json/tests/codec-json.cpp
@@ -378,6 +378,17 @@ SAW_TEST("Struct Array Struct read and write"){
 
 	native.template get<"banana">() = {4u};
 	data<schema::TestStructArrayStruct, encode::Json> json;
+	/**
+	 * Usually I would do something like this
+	{
+		auto& banana = native.template get<"banana">();
+		{
+			banana.at(0).get<"foo">().set(5);
+			banana.at(0).get<"bar">().set("baz");
+		}
+	}
+	* But this also works.
+	*/
 
 	native.template get<"banana">().at(0).get<"foo">().set(5);
 	native.template get<"banana">().at(0).get<"bar">().set("baz");
diff --git a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
index 2c87714..c71582e 100644
--- a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
@@ -63,3 +63,73 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Sycl<saw::enc
 		}
 	};
 }
+
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::MixedArray>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_mixed_float_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity){
+	return {
+		[&](saw::data<sch::MixedArray, saw::encode::Sycl<saw::encode::Native>>& in, cl::sycl::queue* cmd) -> saw::error_or<void> {
+			auto in_size = in.size();
+
+			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()};
+					for(uint64_t i = 0; i < arithmetic_intensity; ++i){
+						if( foo.get() == 1.1e12 ){
+							acc_buff[0u].at(it[0u]) = 0.f;
+						}
+						foo = foo + foo * saw::data<sch::Float64>{1.7342345};
+					}
+					acc_buff[0u].at(it[0u]) = foo;
+				});
+			});
+			return saw::void_t{};
+		}
+	};
+}
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float32Array>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_half_float_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity){
+	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){
+				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()};
+					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;
+						}
+						foo = foo + foo * saw::data<sch::Float32>{1.7342345f};
+					}
+					acc_buff[0u].at(it[0u]) = foo;
+				});
+			});
+			return saw::void_t{};
+		}
+	};
+}
+
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float64Array>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_full_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity){
+	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){
+				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()};
+					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;
+						}
+						foo = foo +foo * saw::data<sch::Float64>{1.7342345};
+					}
+					acc_buff[0u].at(it[0u]) = foo;
+				});
+			});
+			return saw::void_t{};
+		}
+	};
+}
diff --git a/modules/remote-sycl/benchmarks/mixed_precision.hpp b/modules/remote-sycl/benchmarks/mixed_precision.hpp
index fc3fc46..e21a2e6 100644
--- a/modules/remote-sycl/benchmarks/mixed_precision.hpp
+++ b/modules/remote-sycl/benchmarks/mixed_precision.hpp
@@ -24,6 +24,18 @@ using MixedPrecisionBenchmarkInterface = Interface<
 	Member<Function<Float64Array,Void>, "float64">,
 	Member<Function<Float32Array,Void>, "float32">
 >;
+
+/**
+ * This is more of a presentation
+ */
+template<typename FloatArraySchema>
+using FloatPrecisionBenchmarkInterface = Interface<
+	Member<Function<FloatArraySchema,Void>, "float_arr">
+>;
 }
 
 saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_mixed_precision(cl::sycl::event& mixed_ev, cl::sycl::event& float64_ev, cl::sycl::event& float32_ev, uint64_t& arithmetic_intensity);
+
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::MixedArray>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_mixed_float_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity);
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float32Array>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_half_float_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity);
+saw::interface<sch::FloatPrecisionBenchmarkInterface<sch::Float64Array>, saw::encode::Sycl<saw::encode::Native>, cl::sycl::queue*> listen_full_precision(cl::sycl::event& ev, uint64_t& arithmetic_intensity);