#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{};
		}
	};
}