7 files changed, 702 insertions, 0 deletions
diff --git a/modules/remote-hip/c++/SConscript b/modules/remote-hip/c++/SConscript
new file mode 100644
index 0000000..29b7ca4
--- /dev/null
+++ b/modules/remote-hip/c++/SConscript
@@ -0,0 +1,38 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+dev_sycl_env = env.Clone();
+
+dev_sycl_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+dev_sycl_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += dev_sycl_env.sources;
+env.headers += dev_sycl_env.headers;
+
+## Shared lib
+objects_shared = []
+dev_sycl_env.add_source_files(objects_shared, dev_sycl_env.sources, shared=True);
+env.library_shared = dev_sycl_env.SharedLibrary('#build/forstio-remote-sycl', [objects_shared]);
+
+## Static lib
+objects_static = []
+dev_sycl_env.add_source_files(objects_static, dev_sycl_env.sources, shared=False);
+env.library_static = dev_sycl_env.StaticLibrary('#build/forstio-remote-sycl', [objects_static]);
+
+# Set Alias
+env.Alias('library_remote_sycl', [env.library_shared, env.library_static]);
+
+env.targets += ['library_remote_sycl'];
+
+# Install
+env.Install('$prefix/lib/', [env.library_shared, env.library_static]);
+env.Install('$prefix/include/forstio/remote/sycl/', [dev_sycl_env.headers]);
diff --git a/modules/remote-hip/c++/common.hpp b/modules/remote-hip/c++/common.hpp
new file mode 100644
index 0000000..cf1ca93
--- /dev/null
+++ b/modules/remote-hip/c++/common.hpp
@@ -0,0 +1,22 @@
+#pragma once
+
+#include <forstio/remote/remote.hpp>
+#include <forstio/codec/data.hpp>
+#include <forstio/codec/id_map.hpp>
+
+namespace saw {
+namespace rmt {
+struct Hip {};
+}
+namespace encode {
+template<typename Inner>
+struct Hip {};
+}
+
+template<>
+class remote<rmt::Hip>;
+
+template<typename T>
+class device;
+
+}
diff --git a/modules/remote-hip/c++/data.hpp b/modules/remote-hip/c++/data.hpp
new file mode 100644
index 0000000..04c8111
--- /dev/null
+++ b/modules/remote-hip/c++/data.hpp
@@ -0,0 +1,18 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace saw {
+
+/**
+ * Generic wrapper class which stores data on the sycl side.
+ * Most of the times this will be a root object.
+ */
+template<typename Schema>
+class data<Schema, encode::Hip<encode::Native>> {
+private:
+public:
+	data(const data<Schema, encode::Native>& data__)
+	{}
+};
+}
diff --git a/modules/remote-hip/c++/device.hpp b/modules/remote-hip/c++/device.hpp
new file mode 100644
index 0000000..97204f0
--- /dev/null
+++ b/modules/remote-hip/c++/device.hpp
@@ -0,0 +1,19 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace saw {
+/**
+ * Represents a remote Sycl device.
+ */
+template<>
+class device<rmt::Sycl> final {
+private:
+public:
+	device() = default;
+
+	SAW_FORBID_COPY(device);
+	SAW_FORBID_MOVE(device);
+};
+
+}
diff --git a/modules/remote-hip/c++/remote.hpp b/modules/remote-hip/c++/remote.hpp
new file mode 100644
index 0000000..5bb41cf
--- /dev/null
+++ b/modules/remote-hip/c++/remote.hpp
@@ -0,0 +1,91 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace saw {
+
+template<>
+struct remote_address<rmt::Hip> {
+private:
+	uint64_t dev_id_;
+
+	SAW_FORBID_COPY(remote_address);
+	SAW_FORBID_MOVE(remote_address);
+public:
+	remote_address(uint64_t id):
+		dev_id_{id}
+	{}
+
+	uint64_t get_device_id() const {
+		return dev_id_;
+	}
+};
+
+template<>
+class remote<rmt::Hip> {
+private:
+	SAW_FORBID_COPY(remote);
+	SAW_FORBID_MOVE(remote);
+
+	struct key_t {
+		uint64_t device_id;
+		uint32_t sch_id;
+		uint32_t enc_id;
+		
+		bool operator<(const key_t& rhs) const {
+			if(device_id != rhs.device_id){
+				return device_id < rhs.device_id;
+			}
+			if(sch_id != rhs.sch_id){
+				return sch_id < rhs.sch_id;
+			}
+			if(enc_id != rhs.enc_id){
+				return enc_id < rhs.enc_id;
+			}
+			return false;
+		}
+	};
+
+	std::map<uint64_t, our<device<rmt::Hip>>> devs_;
+	std::map<key_t, ptr<i_data_server<rmt::Hip>>> reg_dat_srvs_;
+public:
+	/**
+	 * Default constructor
+	 */
+	remote(){}
+
+	/**
+	 * For now we don't need to specify the location since
+	 * we just create a default.
+	 */
+	conveyor<own<remote_address<rmt::Hip>>> resolve_address(uint64_t dev_id = 0u){
+		return heap<remote_address<rmt::Hip>>(dev_id);
+	}
+
+	/**
+	 * Parse address, but don't resolve it.
+	 */
+	error_or<own<remote_address<rmt::Hip>>> parse_address(uint64_t dev_id = 0u){
+		return heap<remote_address<rmt::Hip>>(dev_id);
+	}
+
+	/**
+	 * Spin up data server
+	 */
+	template<typename Schema, typename Encoding>
+	error_or<own<data_server<Schema, Encoding, rmt::Hip>>> data_listen(remote_address<rmt::Hip>& dev){
+		return heap<data_server<Schema, Encoding, rmt::Hip>>(dev);
+	}
+
+	/**
+	 * Spin up a rpc server
+	 */
+	template<typename Iface, typename Encoding>
+	rpc_server<Iface, Encoding, rmt::Hip> listen(remote_address<rmt::Hip>& dev, typename rpc_server<Iface, Encoding, rmt::Hip>::InterfaceT iface){
+		//using RpcServerT = rpc_server<Iface, Encoding, rmt::Hip>;
+		//using InterfaceT = typename RpcServerT::InterfaceT;
+		return {share<device<rmt::Hip>>(), std::move(iface)};
+	}
+};
+
+}
diff --git a/modules/remote-hip/c++/rpc.hpp b/modules/remote-hip/c++/rpc.hpp
new file mode 100644
index 0000000..f6b519b
--- /dev/null
+++ b/modules/remote-hip/c++/rpc.hpp
@@ -0,0 +1,276 @@
+#pragma once
+
+#include "common.hpp"
+#include "remote.hpp"
+#include "data.hpp"
+#include "device.hpp"
+#include "transfer.hpp"
+
+namespace saw {
+/**
+ * Remote data class for the Sycl backend.
+ */
+template<typename T, typename Encoding>
+class remote_data<T, Encoding, rmt::Hip> final {
+private:
+	/**
+	 * An identifier to the data being held on the remote
+	 */
+	id<T> data_id_;
+
+	/**
+	 * The sycl queue object
+	 */
+	cl::sycl::queue* queue_;
+public:
+	/**
+	 * Main constructor
+	 */
+	remote_data(id<T> data_id__, cl::sycl::queue& queue__):
+		data_id_{data_id__},
+		queue_{&queue__}
+	{}
+
+	/**
+	 * Destructor specifically designed to deallocate on the device.
+	 */
+	~remote_data(){}
+
+	SAW_FORBID_COPY(remote_data);
+	SAW_FORBID_MOVE(remote_data);
+	/**
+	remote_data(const id<T>& id, id_map<T, Encoding, rmt::Hip>& map, cl::sycl::queue& queue__):
+		id_{id},
+		map_{&map}
+	{}
+	*/
+
+	/**
+	 * Wait for the data
+	 */
+	error_or<data<T,Encoding>> wait(){
+		return make_error<err::not_implemented>();
+	}
+
+	/**
+	 * Request data asynchronously
+	 */
+	// conveyor<data<T,Encoding>> on_receive(); /// Stopped here
+};
+
+/**
+ * Meant to be a helper object which holds the allocated data on the sycl side
+ */
+//template<typename Schema, typename Encoding, typename Backend>
+//class device_data;
+
+/**
+ * This class helps in regards to the ownership on the server side
+template<typename Schema, typename Encoding>
+class device_data<Schema, Encoding, rmt::Hip> {
+private:
+	data<Schema,Encoding,storage::Default>* device_data_;
+	cl::sycl::queue* queue_;
+public:
+	device_data(data<Schema,Encoding,storage::Default>& device_data__, cl::sycl::queue& queue__):
+		device_data_{&device_data__},
+		queue_{&queue__}
+	{}
+	
+	~device_data(){
+		if(device_data_){
+			cl::sycl::free(device_data_,queue_);
+			device_data_ = nullptr;
+		}
+	}
+
+	SAW_FORBID_COPY(device_data);
+	SAW_FORBID_MOVE(device_data);
+};
+ */
+
+}
+// Maybe a helper impl tmpl file?
+namespace saw {
+
+
+namespace impl {
+template<typename Func>
+struct rpc_func_type_helper;
+
+template<typename Response, typename Request> 
+struct rpc_func_type_helper<schema::Function<Request, Response>>{
+	using type = tmpl_group<Response, Request>;
+};
+
+template<typename Iface>
+struct rpc_iface_type_helper {
+	using type = tmpl_group<>;
+};
+
+template<typename Func, string_literal K, typename... Functions, string_literal... Keys>
+struct rpc_iface_type_helper<schema::Interface<schema::Member<Func,K>,schema::Member<Functions,Keys>...>> {
+	using inner_type = typename rpc_func_type_helper<Func>::type;
+	using type = typename tmpl_concat<inner_type, typename rpc_iface_type_helper<schema::Interface<schema::Member<Functions,Keys>...>>::type>::type;
+};
+}
+
+/**
+ * Rpc Client class for the Sycl backend.
+ */
+template<typename Iface, typename Encoding>
+class rpc_client<Iface, Encoding, rmt::Hip> {
+public:
+private:
+	/**
+	 * Server this client is tied to
+	 */
+	rpc_server<Iface, Encoding, rmt::Hip>* srv_;
+
+	/**
+	 * TransferClient created from the internal RPC data server
+	 */
+	data_client<typename impl::rpc_iface_type_helper<Iface>::type, Encoding, rmt::Hip> data_client_;
+
+	/**
+	 * Generated some sort of id for the request.
+	 */
+public:
+	rpc_client(rpc_server<Iface, Encoding, rmt::Hip>& srv):
+		srv_{&srv},
+		data_client_{srv_->data_server}
+	{}
+
+	/**
+	 * Rpc call
+	 */
+	template<string_literal Name>
+	error_or<
+		id<
+			typename schema_member_type<Name, Iface>::type::ResponseT
+		>
+	> call(const data_or_id<typename schema_member_type<Name, Iface>::type::RequestT, Encoding>& input){
+		auto next_free_id = srv_->template next_free_id<typename schema_member_type<Name, Iface>::type::ResponseT>();
+		return srv_->template call<Name>(input, next_free_id);
+	}
+
+};
+
+/**
+ * Rpc Server class for the Sycl backend.
+ */
+template<typename Iface, typename Encoding>
+class rpc_server<Iface, Encoding, rmt::Hip> {
+public:
+	using InterfaceCtxT = cl::sycl::queue*;
+	using InterfaceT = interface<Iface, encode::Sycl<Encoding>, InterfaceCtxT>;
+
+private:
+	/**
+	 * Device instance enabling the use of the remote device.
+	 */
+	our<device<rmt::Hip>> device_;
+
+	using DataServerT = data_server<typename impl::rpc_iface_type_helper<Iface>::type, Encoding, rmt::Hip>;
+	/**
+	 * Data server storing the relevant data
+	 */
+	DataServerT* data_server_;
+
+	/**
+	 * The interface including the relevant context class.
+	 */
+	interface<Iface, Encoding, InterfaceCtxT> cl_interface_;
+
+public:
+
+	/**
+	 * Main constructor
+	 */
+	rpc_server(our<device<rmt::Hip>> dev__, DataServerT& data_server__, InterfaceT cl_iface):
+		device_{std::move(dev__)},
+		data_server_{&data_server__},
+		cl_interface_{std::move(cl_iface)}
+	{}
+	
+	/**
+	 * Ask which id the server prefers as the next one. Only available for fast requests on no roundtrip setups.
+	 */
+	/**
+	template<typename T>
+	id<T> next_free_id() const {
+		return std::get<id_map<T,Encoding,rmt::Hip>>(storage_.maps).next_free_id();
+	}
+	*/
+
+	/**
+	template<typename IdT, typename Storage>
+	remote_data<IdT, Encoding, rmt::Hip> request_data(id<IdT> dat_id){
+		return {dat_id, std::get<id_map<IdT,Encoding,rmt::Hip>>(storage_.maps), device_->get_handle()};
+	}
+	*/
+
+	/**
+	 * Rpc call based on the name
+	 */
+	template<string_literal Name>
+	error_or<
+		id<
+			typename schema_member_type<Name, Iface>::type::ResponseT
+		>
+	> call(data_or_id<typename schema_member_type<Name, Iface>::type::RequestT, Encoding> input, id<typename schema_member_type<Name,Iface>::type::ResponseT> rpc_id){
+		using FuncT = typename schema_member_type<Name, Iface>::type;
+
+		/**
+		 * Object needed if and only if the provided data type is not an id
+		 */
+		own<data<typename FuncT::RequestT, encode::Sycl<Encoding>>> dev_tmp_inp = nullptr;
+		/**
+		 * First check if it's data or an id.
+		 * If it's an id, check if it's registered within the storage and retrieve it.
+		 */
+		auto eoinp = [&,this]() -> error_or<data<typename FuncT::RequestT, encode::Sycl<Encoding>>* > {
+			if(input.is_id()){
+				// storage_.maps
+				auto eov = data_server_->template find<typename FuncT::RequestT>(input.get_id());
+				if(eov.is_error()){
+					return std::move(eov.get_error());
+				}
+				return eov.get_value();
+			} else {
+				auto& client_data = input.get_data();
+
+				auto eov = device_->template copy_to_device(client_data);
+				if(eov.is_error()){
+					return std::move(eov.get_error());
+				}
+				auto& val = eov.get_value();
+
+				dev_tmp_inp = heap<data<typename FuncT::RequestT, encode::Sycl<Encoding>>>(std::move(val));
+				device_->get_handle().wait();
+				return dev_tmp_inp.get();
+			}
+		}();
+		if(eoinp.is_error()){
+			return std::move(eoinp.get_error());
+		}
+		auto& inp = *(eoinp.get_value());
+
+		auto eod = cl_interface_.template call<Name>(inp, &(device_->get_handle()));
+
+		if(eod.is_error()){
+			return std::move(eod.get_error());
+		}
+
+		auto& val = eod.get_value();
+		/**
+		 * Store returned data in rpc storage
+		 */
+		auto eoid = data_server_->template insert<typename schema_member_type<Name, Iface>::type::RequestT>(std::move(val), rpc_id);
+		if(eoid.is_error()){
+			return std::move(eoid.get_error());
+		}
+		return rpc_id;
+	}
+};
+}
diff --git a/modules/remote-hip/c++/transfer.hpp b/modules/remote-hip/c++/transfer.hpp
new file mode 100644
index 0000000..8c2cc02
--- /dev/null
+++ b/modules/remote-hip/c++/transfer.hpp
@@ -0,0 +1,238 @@
+#pragma once
+
+#include "common.hpp"
+#include "data.hpp"
+#include "device.hpp"
+
+#include <forstio/error.hpp>
+#include <forstio/reduce_templates.hpp>
+#include <forstio/remote/transfer.hpp>
+
+namespace saw {
+
+template<typename Schema, typename Encoding>
+class data_server<Schema, Encoding, rmt::Hip> final : public i_data_server<rmt::Hip> {
+private:
+	our<device<rmt::Hip>> device_;
+
+	std::map<uint64_t, data<Schema, encode::Sycl<Encoding>>> values_;
+public:
+	data_server(our<device<rmt::Hip>> device__):
+		device_{std::move(device__)}
+	{}
+
+	error_or<void> send(const data<Schema,Encoding>& dat, id<Schema> store_id){
+		auto eo_val = device_->template copy_to_device(dat);
+		if(eo_val.is_error()){
+			auto& err = eo_val.get_error();
+			return std::move(err);
+		}
+		auto& val = eo_val.get_value();
+
+		try {
+			auto insert_res = values_.emplace(std::make_pair(store_id.get_value(), std::move(val)));
+			if(!insert_res.second){
+				return make_error<err::already_exists>();
+			}
+		}catch(const std::exception&){
+			return make_error<err::out_of_memory>();
+		}
+		return make_void();
+	}
+
+	error_or<void> allocate(const data<typename meta_schema<Schema>::MetaSchema, Encoding>& dat, id<Schema> store_id){
+		return make_error<err::not_implemented>("Allocate not implemented");
+		return make_void();
+	}
+	
+	error_or<data<Schema,Encoding>> receive(id<Schema> store_id){
+		return make_error<err::not_implemented>("Receive not implemented");
+	}
+	
+	error_or<void> erase(id<Schema> store_id){
+		return make_error<err::not_implemented>("Erase not implemented");
+		return make_void();
+	}
+	
+	error_or<ptr<data<Schema, encode::Sycl<Encoding>>>> find(id<Schema> store_id){
+		auto find_res = values_.find(store_id.get_value());
+		if(find_res == values_.end()){
+			return make_error<err::not_found>();
+		}
+
+		return {(find_res.second)};
+	}
+};
+
+template<typename... Schema, typename Encoding>
+class data_server<tmpl_group<Schema...>, Encoding, rmt::Hip> {
+private:
+	/**
+	 * Device context class
+	 */
+	our<device<rmt::Hip>> device_;
+
+	/**
+	 * Store for the data the server manages.
+	 */
+	typename impl::data_server_redux<encode::Sycl<Encoding>, typename tmpl_reduce<tmpl_group<Schema...>>::type >::type values_;
+public:
+	/**
+	 * Main constructor
+	 */
+	data_server(our<device<rmt::Hip>> device__):
+		device_{std::move(device__)}
+	{}
+
+	/**
+	 * Get data which we will store.
+	 */
+	template<typename Sch>
+	error_or<void> send(const data<Sch, Encoding>& dat, id<Sch> store_id){
+		auto& vals = std::get<std::unordered_map<uint64_t, data<Sch,encode::Sycl<Encoding>>>>(values_);
+		auto eoval = device_->template copy_to_device<Sch, Encoding>(dat);
+		if(eoval.is_error()){
+			auto& err = eoval.get_error();
+			return std::move(err);
+		}
+		auto& val = eoval.get_value();
+		try {
+			auto insert_res = vals.insert(std::make_pair(store_id.get_value(), std::move(val)));
+			if(!insert_res.second){
+				return make_error<err::already_exists>();
+			}
+		}catch ( std::exception& ){
+			return make_error<err::out_of_memory>();
+		}
+		return void_t{};
+	}
+
+	template<typename Sch>
+	error_or<void> allocate(const data<typename meta_schema<Sch>::MetaSchema, Encoding>& dat, id<Sch> store_id){
+		auto& vals = std::get<std::unordered_map<uint64_t, data<Sch,encode::Sycl<Encoding>>>>(values_);
+		auto eoval = device_->template allocate_on_device<Sch, Encoding>(dat);
+		if(eoval.is_error()){
+			auto& err = eoval.get_error();
+			return std::move(err);
+		}
+		auto& val = eoval.get_value();
+		try {
+			auto insert_res = vals.insert(std::make_pair(store_id.get_value(), std::move(val)));
+			if(!insert_res.second){
+				return make_error<err::already_exists>();
+			}
+		}catch ( std::exception& ){
+			return make_error<err::out_of_memory>();
+		}
+		return void_t{};
+	}
+
+	/**
+	 * Requests data from the server
+	 */
+	template<typename Sch>
+	error_or<data<Sch, Encoding>> receive(id<Sch> store_id){
+		auto& vals = std::get<std::unordered_map<uint64_t, data<Sch,encode::Sycl<Encoding>>>>(values_);
+		auto find_res = vals.find(store_id.get_value());
+		if(find_res == vals.end()){
+			return make_error<err::not_found>();
+		}
+		auto& dat = find_res->second;
+
+		auto eoval = device_->template copy_to_host<Sch, Encoding>(dat);
+		return eoval;
+	}
+
+	/**
+	 * Request an erase of the stored data
+	 */
+	template<typename Sch>
+	error_or<void> erase(id<Sch> store_id){
+		auto& vals = std::get<std::unordered_map<uint64_t, data<Sch,Encoding>>>(values_);
+		auto erase_op = vals.erase(store_id.get_value());
+		if(erase_op == 0u){
+			return make_error<err::not_found>();
+		}
+		return void_t{};
+	}
+
+	/**
+	 * Get the stored data on the server side for immediate use.
+	 * Insert operations may invalidate the pointer.
+	 */
+	template<typename Sch>
+	error_or<data<Sch, encode::Sycl<Encoding>>*> find(id<Sch> store_id){
+		auto& vals = std::get<std::unordered_map<uint64_t, data<Sch,Encoding>>>(values_);
+		auto find_res = vals.find(store_id.get_value());
+		if(find_res == vals.end()){
+			return make_error<err::not_found>();
+		}
+
+		return &(find_res.second);
+	}
+};
+
+/**
+ * Client for transporting data to remote and receiving data back
+ */
+template<typename... Schema, typename Encoding>
+class data_client<tmpl_group<Schema...>, Encoding, rmt::Hip> {
+private:
+	/**
+	 * Corresponding server for this client
+	 */
+	data_server<tmpl_group<Schema...>, Encoding, rmt::Hip>* srv_;
+
+	/**
+	 * The next id for identifying issues on the remote side.
+	 */
+	uint64_t next_id_;
+public:
+	/**
+	 * Main constructor
+	 */
+	data_client(data_server<tmpl_group<Schema...>, Encoding, rmt::Hip>& srv__):
+		srv_{&srv__},
+		next_id_{0u}
+	{}
+
+	/**
+	 * Send data to the remote.
+	 */
+	template<typename Sch>
+	error_or<id<Sch>> send(const data<Sch, Encoding>& dat){
+		id<Sch> dat_id{next_id_};
+		auto eov = srv_->send(dat, dat_id);
+		if(eov.is_error()){
+			auto& err = eov.get_error();
+			return std::move(err);
+		}
+
+		++next_id_;
+		return dat_id;
+	}
+
+	/**
+	 * Receive data
+	 */
+	template<typename Sch>
+	conveyor<data<Sch, Encoding>> receive(id<Sch> dat_id){
+		auto eov = srv_->receive(dat_id);
+		if(eov.is_error()){
+			auto& err = eov.get_error();
+			return std::move(err);
+		}
+
+		auto& val = eov.get_value();
+		return std::move(val);
+	}
+
+	/**
+	 * Erase data
+	 */
+	template<typename Sch>
+	error_or<void> erase(id<Sch> dat_id){
+		return srv_->erase(dat_id);
+	}
+};
+}