From 0290e02fb1e4d3492a166e6eff3210100251f33a Mon Sep 17 00:00:00 2001
From: "Claudius \"keldu\" Holeksa" <mail@keldu.de>
Date: Tue, 2 Jul 2024 20:29:02 +0200
Subject: Fixed running benchmarks

---
 .../benchmarks/kernel_mixed_precision.cpp          | 13 +++---
 modules/remote-sycl/benchmarks/mixed_precision.cpp | 53 ++++++++++++----------
 2 files changed, 36 insertions(+), 30 deletions(-)

(limited to 'modules/remote-sycl/benchmarks')

diff --git a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
index 0ac9756..c17c137 100644
--- a/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/kernel_mixed_precision.cpp
@@ -6,33 +6,34 @@ saw::interface<sch::MixedPrecisionBenchmarkInterface, saw::encode::Native, saw::
 		 * Mixed
 		 */
 		[&](saw::data<sch::MixedArray, saw::encode::Native, saw::rmt::Sycl>& in, cl::sycl::queue* cmd) -> saw::error_or<void> {
-	
+			uint64_t in_size = in.size();
+
 			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){
+				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> {
-	
+			uint64_t 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()), [=] (cl::sycl::id<1> it){
+				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> {
-	
+			uint64_t 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()), [=] (cl::sycl::id<1> it){
+				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};
 				});
 			});
diff --git a/modules/remote-sycl/benchmarks/mixed_precision.cpp b/modules/remote-sycl/benchmarks/mixed_precision.cpp
index e63a814..b979b0c 100644
--- a/modules/remote-sycl/benchmarks/mixed_precision.cpp
+++ b/modules/remote-sycl/benchmarks/mixed_precision.cpp
@@ -1,11 +1,12 @@
 #include "./mixed_precision.hpp"
 #include <forstio/codec/schema.hpp>
 
+#include <sstream>
 
 int main(){
 	using namespace saw;
 
-	constexpr uint64_t max_test_size = 1024ul * 1024ul * 512ul;
+	constexpr uint64_t max_test_size = 1024ul * 1024ul * 256ul;
 
 	std::random_device r;
 	std::default_random_engine e1{r()};
@@ -28,21 +29,21 @@ int main(){
 		return -1;
 	}
 	
-	data<sch::MixedArray> mixed_host_data;
-	data<sch::Float64Array> float64_host_data;
-	data<sch::Float32Array> 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 sycl_iface = listen_mixed_precision(mixed_ev, float64_ev, float32_ev);
+		
+	data<sch::MixedArray> mixed_host_data;
+	data<sch::Float64Array> float64_host_data;
+	data<sch::Float32Array> float32_host_data;
+
+	auto time_eval = [](std::stringstream& sstr, cl::sycl::event& ev){
 		auto end = ev.get_profiling_info<cl::sycl::info::event_profiling::command_end>();
 		auto start = ev.get_profiling_info<cl::sycl::info::event_profiling::command_start>();
 
-		std::cout<<"Elapsed "<<tag<<" kernel time: "<< (end-start) / 1.0e9 << " seconds\n";
+		sstr<<(end-start) / 1.0e9;
 	};
 	
 	auto& device = rmt_addr->get_device();
@@ -50,8 +51,9 @@ int main(){
 	/**
 	 * Warmup
 	 */
-	{
-		uint64_t test_size = max_test_size;
+	std::cout<<"Warming up ..."<<std::endl;
+	for(uint64_t test_size = 1024ul; test_size < max_test_size; test_size *= 2ul){
+	
 		mixed_host_data = {test_size};
 		float64_host_data = {test_size};
 		float32_host_data = {test_size};
@@ -66,19 +68,20 @@ int main(){
 		data<sch::Float32Array, encode::Native, rmt::Sycl> 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();
-
 	}
 
+	std::cout<<"Benchmark starting ..."<<std::endl;
 	/**
 	 * Benchmark
 	 */
+	std::stringstream sstr;
 	for(uint64_t test_size = 1ul; test_size < max_test_size; test_size *= 2ul){
-		device.get_handle().wait();
+		data<sch::MixedArray> mixed_host_data;
+		data<sch::Float64Array> float64_host_data;
+		data<sch::Float32Array> float32_host_data;
 		mixed_host_data = {test_size};
 		float64_host_data = {test_size};
 		float32_host_data = {test_size};
@@ -92,19 +95,21 @@ int main(){
 		data<sch::Float64Array, encode::Native, rmt::Sycl> float64_device_data{float64_host_data};
 		data<sch::Float32Array, encode::Native, rmt::Sycl> float32_device_data{float32_host_data};
 		
-		sycl_iface.template call<"float64_32">(mixed_device_data);
+		sstr<<test_size<<",\t";
+		sycl_iface.template call<"float64_32">(mixed_device_data, &(device.get_handle()));
 		device.get_handle().wait();
-		sycl_iface.template call<"float64">(float64_device_data);
+		time_eval(sstr, mixed_ev);
+		sstr<<",\t";
+		sycl_iface.template call<"float64">(float64_device_data, &(device.get_handle()));
 		device.get_handle().wait();
-		sycl_iface.template call<"float32">(float32_device_data);
+		time_eval(sstr, float64_ev);
+		sstr<<",\t";
+		sycl_iface.template call<"float32">(float32_device_data, &(device.get_handle()));
 		device.get_handle().wait();
-		
-		std::cout<<"\nSize: "<<test_size<<'\n';
-		time_eval("Mixed", mixed_ev);
-		time_eval("Float32", float32_ev);
-		time_eval("Float64", float64_ev);
+		time_eval(sstr, float32_ev);
+		sstr<<'\n';
 	}
-	std::cout<<std::endl;
+	std::cout<<sstr.str()<<std::endl;
 
 	return 0;
 }
-- 
cgit v1.2.3