diff options
author | Claudius "keldu" Holeksa <mail@keldu.de> | 2024-07-02 19:46:02 +0200 |
---|---|---|
committer | Claudius "keldu" Holeksa <mail@keldu.de> | 2024-07-02 19:46:02 +0200 |
commit | 25e05907f0292310eaae27a032db0ee274413874 (patch) | |
tree | 283de0ebb6b61add2221436a77bb09e2ff101080 /modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp | |
parent | e51d2b1c0493dfd30d1622c8a0628ecf98c92f1c (diff) |
Preparing benchmark work
Diffstat (limited to 'modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp')
-rw-r--r-- | modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp | 42 |
1 files changed, 42 insertions, 0 deletions
diff --git a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp new file mode 100644 index 0000000..0ac9756 --- /dev/null +++ b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp @@ -0,0 +1,42 @@ +#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){ + return { + /** + * Mixed + */ + [&](saw::data<sch::MixedArray, saw::encode::Native, saw::rmt::Sycl>& in, cl::sycl::queue* cmd) -> saw::error_or<void> { + + mixed_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()), [=] (cl::sycl::id<1> it){ + acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float64>{2.0}; + }); + }); + return saw::void_t{}; + }, + [&](saw::data<sch::Float64Array, saw::encode::Native, saw::rmt::Sycl>& in, cl::sycl::queue* cmd) -> saw::error_or<void> { + + 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()), [=] (cl::sycl::id<1> it){ + acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float64>{2.0}; + }); + }); + return saw::void_t{}; + }, + [&](saw::data<sch::Float32Array, saw::encode::Native, saw::rmt::Sycl>& in, cl::sycl::queue* cmd) -> saw::error_or<void> { + + 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()), [=] (cl::sycl::id<1> it){ + acc_buff[0u].at(it[0u]) = acc_buff[0u].at(it[0u]) * saw::data<sch::Float32>{2.0f}; + }); + }); + return saw::void_t{}; + } + }; +} |