From ff5535b51730974b7933dd93e140579f3232a275 Mon Sep 17 00:00:00 2001
From: "Claudius \"keldu\" Holeksa" <mail@keldu.de>
Date: Thu, 4 Jul 2024 11:45:02 +0200
Subject: Ammending kernel work

---
 .../benchmarks/kernel_mixed_precision.cpp          | 29 +++++++++++++++++++---
 1 file changed, 25 insertions(+), 4 deletions(-)

(limited to 'modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp')

diff --git a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
index 1c82361..591ded2 100644
--- a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
@@ -1,6 +1,6 @@
 #include "mixed_precision.hpp"
 
-saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::rmt::Sycl, cl::sycl::queue*> listen_mixed_precision(cl::sycl::event& mixed_ev, cl::sycl::event& float64_ev, cl::sycl::event& float32_ev){
+saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::rmt::Sycl, 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){
 	return {
 		/**
 		 * Mixed
@@ -12,7 +12,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float64>{1.7342345};
+					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{};
@@ -23,7 +30,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float64>{1.7342345};
+					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{};
@@ -34,7 +48,14 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::
 				auto acc_buff = in.template access<cl::sycl::access::mode::read_write>(h);
 
 				h.parallel_for(cl::sycl::range<1>(in_size), [=] (cl::sycl::id<1> it){
-					acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float32>{1.7342345f};
+					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{};
-- 
cgit v1.2.3