diff options
Diffstat (limited to 'modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp')
| -rw-r--r-- | modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp | 29 |
1 files changed, 25 insertions, 4 deletions
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{}; |