#pragma once #include #include #include #include namespace saw { namespace rmt { struct Sycl {}; } template<> class remote; /** * Remote data class for the Sycl backend. */ template class remote_data { private: /** * Id representing the remote data */ id id_; /** * Storage for the */ id_map* map_; public: /** * Main constructor */ remote_data(const id& id, id_map& map, cl::sycl::queue& queue__): id_{id}, map_{&map} {} /** * Wait for the data */ error_or> wait(){ auto eov = map_->find(id_); if(eov.is_error()){ auto& err = eov.get_error(); return std::move(err); } auto& val = eov.get_value(); std::cout<<"Values Sycl in Map: "<size()<template copy_to_host(); if(eocop.is_error()){ return eocop; } return eocop.get_value(); } } /** * Request data asynchronously */ conveyor> on_receive(); /// Stopped here }; /** * template class data, encode::Native, rmt::Sycl> { public: using Schema = schema::Primitive; using NativeType = typename native_data_type::type; private: NativeType val_; public: data(NativeType val__): val_{val__} {} NativeType get(){ return val_; } }; */ template class data, encode::Native, rmt::Sycl> { public: using Schema = schema::Array; private: uint64_t total_length_; data* device_data_; cl::sycl::queue* queue_; static_assert(is_primitive::value, "Only supports primitives for now"); static_assert(D==1u, "For now we only support 1D Arrays"); public: data(uint64_t size, cl::sycl::queue& q__): total_length_{size}, device_data_{cl::sycl::malloc_device>(size, q__)}, queue_{&q__} { if(!device_data_){ total_length_ = 0u; return; } queue_->wait(); } template data(const data& from, cl::sycl::queue& q__): total_length_{from.size()}, device_data_{cl::sycl::malloc_device>(from.size(), q__)}, queue_{&q__} { if(!device_data_){ total_length_ = 0u; return; } queue_->template copy>(&from.at(0), device_data_, total_length_); queue_->wait(); } data(const data& from): total_length_{from.size()}, device_data_{nullptr}, queue_{from.queue_} { if(total_length_ == 0u || !queue_){ return; } device_data_ = cl::sycl::malloc_device>(from.size(), *queue_); // device_data_ = cl::sycl::malloc_device::type>(from.size(), *queue_); if(!device_data_){ total_length_ = 0u; return; } queue_->template copy>(from.device_data_, device_data_, total_length_); } data(data&& rhs): total_length_{rhs.total_length_}, device_data_{rhs.device_data_}, queue_{rhs.queue_} { rhs.total_length_ = 0u; rhs.device_data_ = nullptr; rhs.queue_ = nullptr; } data& operator=(data&& rhs){ total_length_ = rhs.total_length_; device_data_ = rhs.device_data_; queue_ = rhs.queue_; rhs.total_length_ = 0u; rhs.device_data_ = nullptr; rhs.queue_ = nullptr; return *this; } ~data(){ // free data if(device_data_){ /// SYCL FREE cl::sycl::free(device_data_, *queue_); } } /** * Allocate appropriate meta data and then copy to host */ template error_or> copy_to_host() const { data data_{total_length_}; /// TODO Check success queue_->template copy>(device_data_, &data_.at(0), total_length_); queue_->wait(); return data_; } data& at(uint64_t i){ return device_data_[i]; } uint64_t size() const { return total_length_; } }; namespace impl { template struct rpc_id_map_helper { static_assert(always_false, "Only supports Interface schema types."); }; template struct rpc_id_map_helper, Encoding, Storage> { std::tuple...> maps; }; } template class device; /** * Represents a remote Sycl device. * */ template<> class device { private: cl::sycl::queue cmd_queue_; public: /** * Copy data to device */ template error_or> copy_to_device(const data>& host_data){ (void) host_data; return make_error(); } /** * Copy data to host */ template error_or> copy_to_host(const data>& device_data){ (void) device_data; return make_error(); } }; /** * Device data transport */ /** * Rpc Client class for the Sycl backend. */ template class rpc_client { public: private: /** * Server this client is tied to */ rpc_server* srv_; /** * Generated some sort of id for the request. */ public: rpc_client(rpc_server& srv): srv_{&srv} {} /** * Rpc call */ template error_or< id< typename schema_member_type::type::ResponseT > > call(const data_or_id::type::RequestT, Encoding, Storage>& 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; private: /** * Device instance enabling the use of the remote device. */ device* device_; /** * The interface including the relevant context class. */ interface cl_interface_; /** * Basic storage for response data. */ impl::rpc_id_map_helper storage_; public: /** * Main constructor */ rpc_server(device& dev__, interface cl_iface): device_{&dev__}, cl_interface_{std::move(cl_iface)}, storage_{} {} /** * 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_.cmd_queue_}; } /** * Rpc call based on the name */ template error_or< id< typename schema_member_type::type::ResponseT > > call(data_or_id::type::RequestT, Encoding, ClientAllocation> 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& inner_map = std::get> (storage_.maps); auto eov = inner_map.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(); dev_tmp_inp = heap>(client_data, cmd_queue_); cmd_queue_.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, &cmd_queue_); if(eod.is_error()){ return std::move(eod.get_error()); } auto& val = eod.get_value(); /** * Store returned data in rpc storage */ auto& inner_map = std::get::type::RequestT, Encoding,rmt::Sycl>> (storage_.maps); auto eoid = inner_map.insert_as(std::move(val), rpc_id); if(eoid.is_error()){ return std::move(eoid.get_error()); } return rpc_id; } }; template<> struct remote_address { private: remote* ctx_; SAW_FORBID_COPY(remote_address); SAW_FORBID_MOVE(remote_address); public: remote_address(remote& r_ctx): ctx_{&r_ctx} {} }; template<> class remote { private: SAW_FORBID_COPY(remote); SAW_FORBID_MOVE(remote); public: /** * Default constructor */ remote(){} /** * For now we don't need to specify the location since * we just create a default. */ conveyor>> resolve_address(){ return heap>(*this); } /** * Connect to a device */ device connect_device(const remote_address&){ return {}; } /** * Spin up a rpc server */ template rpc_server listen(const device& dev, typename rpc_server::InterfaceT iface){ using RpcServerT = rpc_server; using InterfaceT = typename RpcServerT::InterfaceT; return {dev, std::move(iface)}; } }; }