From 53738bc556fc13c80dacdb56633ef8b7441e6566 Mon Sep 17 00:00:00 2001 From: Claudius 'keldu' Holeksa Date: Fri, 23 Aug 2024 17:15:18 +0200 Subject: prepare for hip setup --- modules/remote-hip/.nix/derivation.nix | 60 +++++++ modules/remote-hip/SConstruct | 91 +++++++++++ modules/remote-hip/c++/SConscript | 38 +++++ modules/remote-hip/c++/common.hpp | 22 +++ modules/remote-hip/c++/data.hpp | 18 +++ modules/remote-hip/c++/device.hpp | 19 +++ modules/remote-hip/c++/remote.hpp | 91 +++++++++++ modules/remote-hip/c++/rpc.hpp | 276 +++++++++++++++++++++++++++++++++ modules/remote-hip/c++/transfer.hpp | 238 ++++++++++++++++++++++++++++ modules/remote-hip/examples/SConscript | 39 +++++ modules/remote-hip/tests/SConscript | 33 ++++ 11 files changed, 925 insertions(+) create mode 100644 modules/remote-hip/.nix/derivation.nix create mode 100644 modules/remote-hip/SConstruct create mode 100644 modules/remote-hip/c++/SConscript create mode 100644 modules/remote-hip/c++/common.hpp create mode 100644 modules/remote-hip/c++/data.hpp create mode 100644 modules/remote-hip/c++/device.hpp create mode 100644 modules/remote-hip/c++/remote.hpp create mode 100644 modules/remote-hip/c++/rpc.hpp create mode 100644 modules/remote-hip/c++/transfer.hpp create mode 100644 modules/remote-hip/examples/SConscript create mode 100644 modules/remote-hip/tests/SConscript 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 +#include +#include + +namespace saw { +namespace rmt { +struct Hip {}; +} +namespace encode { +template +struct Hip {}; +} + +template<> +class remote; + +template +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 +class data> { +private: +public: + data(const data& 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 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 { +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 { +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>> devs_; + std::map>> reg_dat_srvs_; +public: + /** + * Default constructor + */ + remote(){} + + /** + * For now we don't need to specify the location since + * we just create a default. + */ + conveyor>> resolve_address(uint64_t dev_id = 0u){ + return heap>(dev_id); + } + + /** + * Parse address, but don't resolve it. + */ + error_or>> parse_address(uint64_t dev_id = 0u){ + return heap>(dev_id); + } + + /** + * Spin up data server + */ + template + error_or>> data_listen(remote_address& dev){ + return heap>(dev); + } + + /** + * Spin up a rpc server + */ + template + rpc_server listen(remote_address& dev, typename rpc_server::InterfaceT iface){ + //using RpcServerT = rpc_server; + //using InterfaceT = typename RpcServerT::InterfaceT; + return {share>(), 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 +class remote_data final { +private: + /** + * An identifier to the data being held on the remote + */ + id data_id_; + + /** + * The sycl queue object + */ + cl::sycl::queue* queue_; +public: + /** + * Main constructor + */ + remote_data(id 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& id, id_map& map, cl::sycl::queue& queue__): + id_{id}, + map_{&map} + {} + */ + + /** + * Wait for the data + */ + error_or> wait(){ + return make_error(); + } + + /** + * Request data asynchronously + */ + // conveyor> on_receive(); /// Stopped here +}; + +/** + * Meant to be a helper object which holds the allocated data on the sycl side + */ +//template +//class device_data; + +/** + * This class helps in regards to the ownership on the server side +template +class device_data { +private: + data* device_data_; + cl::sycl::queue* queue_; +public: + device_data(data& 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 +struct rpc_func_type_helper; + +template +struct rpc_func_type_helper>{ + using type = tmpl_group; +}; + +template +struct rpc_iface_type_helper { + using type = tmpl_group<>; +}; + +template +struct rpc_iface_type_helper,schema::Member...>> { + using inner_type = typename rpc_func_type_helper::type; + using type = typename tmpl_concat...>>::type>::type; +}; +} + +/** + * Rpc Client class for the Sycl backend. + */ +template +class rpc_client { +public: +private: + /** + * Server this client is tied to + */ + rpc_server* srv_; + + /** + * TransferClient created from the internal RPC data server + */ + data_client::type, Encoding, rmt::Hip> data_client_; + + /** + * Generated some sort of id for the request. + */ +public: + rpc_client(rpc_server& srv): + srv_{&srv}, + data_client_{srv_->data_server} + {} + + /** + * Rpc call + */ + template + error_or< + id< + typename schema_member_type::type::ResponseT + > + > call(const data_or_id::type::RequestT, Encoding>& input){ + auto next_free_id = srv_->template next_free_id::type::ResponseT>(); + return srv_->template call(input, next_free_id); + } + +}; + +/** + * Rpc Server class for the Sycl backend. + */ +template +class rpc_server { +public: + using InterfaceCtxT = cl::sycl::queue*; + using InterfaceT = interface, InterfaceCtxT>; + +private: + /** + * Device instance enabling the use of the remote device. + */ + our> device_; + + using DataServerT = data_server::type, Encoding, rmt::Hip>; + /** + * Data server storing the relevant data + */ + DataServerT* data_server_; + + /** + * The interface including the relevant context class. + */ + interface cl_interface_; + +public: + + /** + * Main constructor + */ + rpc_server(our> 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 + id next_free_id() const { + return std::get>(storage_.maps).next_free_id(); + } + */ + + /** + template + remote_data request_data(id dat_id){ + return {dat_id, std::get>(storage_.maps), device_->get_handle()}; + } + */ + + /** + * Rpc call based on the name + */ + template + error_or< + id< + typename schema_member_type::type::ResponseT + > + > call(data_or_id::type::RequestT, Encoding> input, id::type::ResponseT> rpc_id){ + using FuncT = typename schema_member_type::type; + + /** + * Object needed if and only if the provided data type is not an id + */ + own>> 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>* > { + if(input.is_id()){ + // storage_.maps + auto eov = data_server_->template find(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>>(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(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::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 +#include +#include + +namespace saw { + +template +class data_server final : public i_data_server { +private: + our> device_; + + std::map>> values_; +public: + data_server(our> device__): + device_{std::move(device__)} + {} + + error_or send(const data& dat, id 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(); + } + }catch(const std::exception&){ + return make_error(); + } + return make_void(); + } + + error_or allocate(const data::MetaSchema, Encoding>& dat, id store_id){ + return make_error("Allocate not implemented"); + return make_void(); + } + + error_or> receive(id store_id){ + return make_error("Receive not implemented"); + } + + error_or erase(id store_id){ + return make_error("Erase not implemented"); + return make_void(); + } + + error_or>>> find(id store_id){ + auto find_res = values_.find(store_id.get_value()); + if(find_res == values_.end()){ + return make_error(); + } + + return {(find_res.second)}; + } +}; + +template +class data_server, Encoding, rmt::Hip> { +private: + /** + * Device context class + */ + our> device_; + + /** + * Store for the data the server manages. + */ + typename impl::data_server_redux, typename tmpl_reduce>::type >::type values_; +public: + /** + * Main constructor + */ + data_server(our> device__): + device_{std::move(device__)} + {} + + /** + * Get data which we will store. + */ + template + error_or send(const data& dat, id store_id){ + auto& vals = std::get>>>(values_); + auto eoval = device_->template copy_to_device(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(); + } + }catch ( std::exception& ){ + return make_error(); + } + return void_t{}; + } + + template + error_or allocate(const data::MetaSchema, Encoding>& dat, id store_id){ + auto& vals = std::get>>>(values_); + auto eoval = device_->template allocate_on_device(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(); + } + }catch ( std::exception& ){ + return make_error(); + } + return void_t{}; + } + + /** + * Requests data from the server + */ + template + error_or> receive(id store_id){ + auto& vals = std::get>>>(values_); + auto find_res = vals.find(store_id.get_value()); + if(find_res == vals.end()){ + return make_error(); + } + auto& dat = find_res->second; + + auto eoval = device_->template copy_to_host(dat); + return eoval; + } + + /** + * Request an erase of the stored data + */ + template + error_or erase(id store_id){ + auto& vals = std::get>>(values_); + auto erase_op = vals.erase(store_id.get_value()); + if(erase_op == 0u){ + return make_error(); + } + return void_t{}; + } + + /** + * Get the stored data on the server side for immediate use. + * Insert operations may invalidate the pointer. + */ + template + error_or>*> find(id store_id){ + auto& vals = std::get>>(values_); + auto find_res = vals.find(store_id.get_value()); + if(find_res == vals.end()){ + return make_error(); + } + + return &(find_res.second); + } +}; + +/** + * Client for transporting data to remote and receiving data back + */ +template +class data_client, Encoding, rmt::Hip> { +private: + /** + * Corresponding server for this client + */ + data_server, 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, Encoding, rmt::Hip>& srv__): + srv_{&srv__}, + next_id_{0u} + {} + + /** + * Send data to the remote. + */ + template + error_or> send(const data& dat){ + id 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 + conveyor> receive(id 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 + error_or erase(id 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']; -- cgit v1.2.3