prepare for hip setup

11 files changed, 925 insertions, 0 deletions

diff --git a/modules/remote-hip/.nix/derivation.nix b/modules/remote-hip/.nix/derivation.nix
new file mode 100644
index 0000000..c63dd19
--- /dev/null
+++ b/modules/remote-hip/.nix/derivation.nix
@@ -0,0 +1,60 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, version
+, forstio
+, openmp
+, keldu
+, ocl-icd
+, lld_15
+, python3
+, bash
+
+, build_examples ? "false"
+, build_benchmarks ? "false"
+}:
+
+let
+
+in stdenv.mkDerivation {
+  pname = "forstio-remote-sycl";
+  inherit version;
+  src = ./..;
+
+  enableParallelBuilding = true;
+
+  nativeBuildInputs = [
+		clang-tools
+		python3
+		scons
+  ];
+
+  buildInputs = [
+    forstio.core
+		forstio.codec
+		forstio.async
+		forstio.remote
+		keldu.adaptivecpp-dev
+		ocl-icd
+		openmp
+		lld_15
+	];
+
+	buildPhase = ''
+		scons build_benchmarks=${build_benchmarks} build_examples=${build_examples}
+	'';
+
+	installPhase = ''
+		scons prefix=$out build_benchmarks=${build_benchmarks} build_examples=${build_examples} install
+	'';
+
+  doCheck = true;
+	checkPhase = ''
+		export ACPP_APPDB_DIR=.
+    scons test
+    ./bin/tests
+  '';
+
+  outputs = ["out" "dev"];
+}
diff --git a/modules/remote-hip/SConstruct b/modules/remote-hip/SConstruct
new file mode 100644
index 0000000..78bf333
--- /dev/null
+++ b/modules/remote-hip/SConstruct
@@ -0,0 +1,91 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+    def isbasestring(s):
+        return isinstance(s,basestring)
+else:
+    def isbasestring(s):
+        return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+    if isbasestring(filetype):
+        dir_path = self.Dir('.').abspath
+        filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+    for path in filetype:
+        target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+        if shared:
+            target_name+='.os'
+            sources.append( self.SharedObject( target=target_name, source=path ) )
+        else:
+            target_name+='.o'
+            sources.append( self.StaticObject( target=target_name, source=path ) )
+    pass
+
+def isAbsolutePath(key, dirname, env):
+	assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+	args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+	help='Installation target location of build results and headers',
+	default='/usr/local/',
+	validator=isAbsolutePath
+)
+
+env_vars.Add(
+    BoolVariable('build_examples',
+        help='Build examples',
+        default=False
+    )
+);
+
+env_vars.Add(
+    BoolVariable('build_benchmarks',
+        help='Build benchmarks',
+        default=False
+    )
+);
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+    CPPDEFINES=['SAW_UNIX'],
+    CXXFLAGS=['-std=c++20','-g','-Wall','-Wextra'],
+    LIBS=['forstio-core'
+        ,'forstio-codec'
+        ,'forstio-async'
+        ,'forstio-remote'
+        ,'acpp-rt'
+        ,'OpenCL'
+        ,'omp'
+    ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('c++/SConscript')
+SConscript('tests/SConscript')
+#SConscript('examples/SConscript')
+#SConscript('benchmarks/SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
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);
+	}
+};
+}
diff --git a/modules/remote-hip/examples/SConscript b/modules/remote-hip/examples/SConscript
new file mode 100644
index 0000000..5966b1b
--- /dev/null
+++ b/modules/remote-hip/examples/SConscript
@@ -0,0 +1,39 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+
+
+examples_sycl_env = examples_env.Clone();
+examples_sycl_env['CXX'] = 'acpp';
+examples_sycl_env['CXXFLAGS'] += ['-O2'];
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+#sycl_objects = [];
+#examples_sycl_env.add_source_files(sycl_objects, ['sycl_basic_kernel.cpp'], shared=True);
+
+#objects_static = []
+#examples_env.sycl_basic = examples_env.Program('#bin/sycl_basic', ['sycl_basic.cpp', env.library_static, sycl_objects]);
+
+# Set Alias
+env.examples = [examples_env.sycl_basic];
+env.Alias('examples', env.examples);
+
+if env["build_examples"]:
+	env.targets += ['examples'];
+	env.Install('$prefix/bin/', env.examples);
+#endif
diff --git a/modules/remote-hip/tests/SConscript b/modules/remote-hip/tests/SConscript
new file mode 100644
index 0000000..f56f0bd
--- /dev/null
+++ b/modules/remote-hip/tests/SConscript
@@ -0,0 +1,33 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+test_cases_env = env.Clone();
+
+test_cases_env['CXX'] = 'acpp';
+
+test_cases_env.Append(LIBS=['forstio-test']);
+
+test_cases_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+test_cases_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += test_cases_env.sources;
+env.headers += test_cases_env.headers;
+
+objects_static = []
+test_cases_env.add_source_files(objects_static, test_cases_env.sources, shared=False);
+test_cases_env.program = test_cases_env.Program('#bin/tests', [objects_static, env.library_static]);
+
+# Set Alias
+env.Alias('test', test_cases_env.program);
+env.Alias('check', test_cases_env.program);
+
+# env.targets += ['test','check'];