Adding possible diverge

1 files changed, 70 insertions, 0 deletions

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{};
+		}
+	};
+}