diff options
| -rw-r--r-- | modules/codec-json/tests/codec-json.cpp | 11 | ||||
| -rw-r--r-- | modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp | 70 | ||||
| -rw-r--r-- | modules/remote-sycl/benchmarks/mixed_precision.hpp | 12 |
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); |