From 25e05907f0292310eaae27a032db0ee274413874 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 2 Jul 2024 19:46:02 +0200 Subject: Preparing benchmark work --- modules/remote-sycl/benchmarks/SConscript | 39 ++++++++ .../benchmarks/kernel_mixed_precision.cpp | 42 ++++++++ modules/remote-sycl/benchmarks/mixed_precision.cpp | 110 +++++++++++++++++++++ modules/remote-sycl/benchmarks/mixed_precision.hpp | 27 +++++ 4 files changed, 218 insertions(+) create mode 100644 modules/remote-sycl/benchmarks/SConscript create mode 100644 modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp create mode 100644 modules/remote-sycl/benchmarks/mixed_precision.cpp create mode 100644 modules/remote-sycl/benchmarks/mixed_precision.hpp (limited to 'modules/remote-sycl/benchmarks') diff --git a/modules/remote-sycl/benchmarks/SConscript b/modules/remote-sycl/benchmarks/SConscript new file mode 100644 index 0000000..9976e0e --- /dev/null +++ b/modules/remote-sycl/benchmarks/SConscript @@ -0,0 +1,39 @@ +#!/bin/false + +import os +import os.path +import glob + + +Import('env') + +dir_path = Dir('.').abspath + +# Environment for base library +benchmarks_env = env.Clone(); + + +benchmarks_sycl_env = benchmarks_env.Clone(); +benchmarks_sycl_env['CXX'] = 'acpp'; +benchmarks_sycl_env['CXXFLAGS'] += ['-O2']; + +benchmarks_env.sources = sorted(glob.glob(dir_path + "/*.cpp")) +benchmarks_env.headers = sorted(glob.glob(dir_path + "/*.hpp")) + +env.sources += benchmarks_env.sources; +env.headers += benchmarks_env.headers; + +sycl_objects = []; +benchmarks_sycl_env.add_source_files(sycl_objects, ['kernel_mixed_precision.cpp'], shared=True); + +objects_static = [] +benchmarks_env.sycl_basic = benchmarks_env.Program('#bin/benchmark_mixed_precision', ['mixed_precision.cpp', env.library_static, sycl_objects]); + +# Set Alias +env.benchmarks = [benchmarks_env.sycl_basic]; +env.Alias('benchmarks', env.benchmarks); + +if env["build_benchmarks"]: + env.targets += ['benchmarks']; + env.Install('$prefix/bin/', env.benchmarks); +#endif 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 listen_mixed_precision(cl::sycl::event& mixed_ev, cl::sycl::event& float64_ev, cl::sycl::event& float32_ev){ + return { + /** + * Mixed + */ + [&](saw::data& in, cl::sycl::queue* cmd) -> saw::error_or { + + mixed_ev = cmd->submit([&](cl::sycl::handler& h){ + auto acc_buff = in.template access(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{2.0}; + }); + }); + return saw::void_t{}; + }, + [&](saw::data& in, cl::sycl::queue* cmd) -> saw::error_or { + + float64_ev = cmd->submit([&](cl::sycl::handler& h){ + auto acc_buff = in.template access(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{2.0}; + }); + }); + return saw::void_t{}; + }, + [&](saw::data& in, cl::sycl::queue* cmd) -> saw::error_or { + + float32_ev = cmd->submit([&](cl::sycl::handler& h){ + auto acc_buff = in.template access(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{2.0f}; + }); + }); + return saw::void_t{}; + } + }; +} diff --git a/modules/remote-sycl/benchmarks/mixed_precision.cpp b/modules/remote-sycl/benchmarks/mixed_precision.cpp new file mode 100644 index 0000000..e63a814 --- /dev/null +++ b/modules/remote-sycl/benchmarks/mixed_precision.cpp @@ -0,0 +1,110 @@ +#include "./mixed_precision.hpp" +#include + + +int main(){ + using namespace saw; + + constexpr uint64_t max_test_size = 1024ul * 1024ul * 512ul; + + std::random_device r; + std::default_random_engine e1{r()}; + std::uniform_real_distribution<> dis{-1.0,1.0}; + + + saw::event_loop loop; + saw::wait_scope wait{loop}; + + remote rmt; + + own> rmt_addr{}; + + rmt.resolve_address().then([&](auto addr){ + rmt_addr = std::move(addr); + }).detach(); + + wait.poll(); + if(!rmt_addr){ + return -1; + } + + data mixed_host_data; + data float64_host_data; + data float32_host_data; + + cl::sycl::event mixed_ev; + cl::sycl::event float32_ev; + cl::sycl::event float64_ev; + + auto sycl_iface = listen_mixed_precision(mixed_ev, float64_ev, float32_ev); + + auto time_eval = [](std::string_view tag, cl::sycl::event& ev){ + auto end = ev.get_profiling_info(); + auto start = ev.get_profiling_info(); + + std::cout<<"Elapsed "<get_device(); + + /** + * Warmup + */ + { + uint64_t test_size = max_test_size; + mixed_host_data = {test_size}; + float64_host_data = {test_size}; + float32_host_data = {test_size}; + for(uint64_t i = 0; i < test_size; ++i){ + double gen_num = dis(e1); + mixed_host_data.at(i) = static_cast(gen_num); + float64_host_data.at(i) = static_cast(gen_num); + float32_host_data.at(i) = static_cast(gen_num); + } + data mixed_device_data{mixed_host_data}; + data float64_device_data{float64_host_data}; + data float32_device_data{float32_host_data}; + + sycl_iface.template call<"float64_32">(mixed_device_data, &(device.get_handle())); + device.get_handle().wait(); + sycl_iface.template call<"float64">(float64_device_data, &(device.get_handle())); + device.get_handle().wait(); + sycl_iface.template call<"float32">(float32_device_data, &(device.get_handle())); + device.get_handle().wait(); + + } + + /** + * Benchmark + */ + for(uint64_t test_size = 1ul; test_size < max_test_size; test_size *= 2ul){ + device.get_handle().wait(); + mixed_host_data = {test_size}; + float64_host_data = {test_size}; + float32_host_data = {test_size}; + for(uint64_t i = 0; i < test_size; ++i){ + double gen_num = dis(e1); + mixed_host_data.at(i) = static_cast(gen_num); + float64_host_data.at(i) = static_cast(gen_num); + float32_host_data.at(i) = static_cast(gen_num); + } + data mixed_device_data{mixed_host_data}; + data float64_device_data{float64_host_data}; + data float32_device_data{float32_host_data}; + + sycl_iface.template call<"float64_32">(mixed_device_data); + device.get_handle().wait(); + sycl_iface.template call<"float64">(float64_device_data); + device.get_handle().wait(); + sycl_iface.template call<"float32">(float32_device_data); + device.get_handle().wait(); + + std::cout<<"\nSize: "< +>; + +using Float64Array = Array< + Float64 +>; + +using Float32Array = Array< + Float32 +>; + +using MixedPrecisionBenchmarkInterface = Interface< + Member, "float64_32">, + Member, "float64">, + Member, "float32"> +>; +} + +saw::interface listen_mixed_precision(cl::sycl::event& mixed_ev, cl::sycl::event& float64_ev, cl::sycl::event& float32_ev); -- cgit v1.2.3