From 86b06a3fee2cd7635a9ab486e2a35bdf1e81ce38 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Fri, 21 Jun 2024 19:44:34 +0200 Subject: Moving forward with basic test for sycl --- default.nix | 2 +- modules/codec/c++/interface.hpp | 22 +++++-- modules/codec/c++/schema.hpp | 6 ++ modules/codec/tests/codec.cpp | 18 ++++-- modules/io_codec/c++/transfer.hpp | 9 +++ modules/remote-sycl/.nix/derivation.nix | 2 +- modules/remote-sycl/c++/remote.hpp | 71 +++++++++++------------ modules/remote-sycl/tests/calculator.cpp | 69 ---------------------- modules/remote-sycl/tests/calculator.foo | 69 ++++++++++++++++++++++ modules/remote-sycl/tests/sycl_basics.cpp | 96 +++++++++++++++++++++++++++++++ 10 files changed, 247 insertions(+), 117 deletions(-) create mode 100644 modules/io_codec/c++/transfer.hpp delete mode 100644 modules/remote-sycl/tests/calculator.cpp create mode 100644 modules/remote-sycl/tests/calculator.foo create mode 100644 modules/remote-sycl/tests/sycl_basics.cpp diff --git a/default.nix b/default.nix index 79f7777..681daf3 100644 --- a/default.nix +++ b/default.nix @@ -79,7 +79,7 @@ in rec { inherit clang-tools; openmp = pkgs.llvmPackages_15.openmp; - build_examples = "true"; + build_examples = "false"; }; tools = pkgs.callPackage modules/tools/.nix/derivation.nix { diff --git a/modules/codec/c++/interface.hpp b/modules/codec/c++/interface.hpp index 0f41f55..e1c9a12 100644 --- a/modules/codec/c++/interface.hpp +++ b/modules/codec/c++/interface.hpp @@ -12,28 +12,40 @@ template +struct FuncReturnTypeHelper { + using Type = data; +}; + +template +struct FuncReturnTypeHelper { + using Type = void; +}; + template struct FuncTypeHelper { - using Type = std::function(data&, Ctx)>; + using Type = std::function::Type>(data&, Ctx)>; }; template struct FuncTypeHelper { - using Type = std::function(data&)>; + using Type = std::function::Type>(data&)>; }; + } template class function, Encode, Storage, Context> { private: typename impl::FuncTypeHelper::Type func_; + using ResponseDataType = typename impl::FuncReturnTypeHelper::Type; public: template function(Func func): func_{std::move(func)} {} - error_or> call(data& req, Context ctx = {}){ + error_or call(data& req, Context ctx = {}){ if constexpr (std::is_same_v){ (void) ctx; return func_(req); @@ -84,13 +96,13 @@ public: template error_or< - data< + typename impl::FuncReturnTypeHelper< typename parameter_pack_type< parameter_key_pack_index< Lit, Names... >::value , Responses...>::type - , Encode, Storage>> call( + , Encode, Storage>::Type > call( data< typename parameter_pack_type< parameter_key_pack_index< diff --git a/modules/codec/c++/schema.hpp b/modules/codec/c++/schema.hpp index 2ef7c77..6a69425 100644 --- a/modules/codec/c++/schema.hpp +++ b/modules/codec/c++/schema.hpp @@ -11,6 +11,12 @@ template struct is_primitive { namespace schema { // NOLINTBEGIN + +/** + * Void Type used for function schemas + */ +struct Void {}; + template struct Member { static constexpr string_literal name = "Member"; diff --git a/modules/codec/tests/codec.cpp b/modules/codec/tests/codec.cpp index 720b734..1bec214 100644 --- a/modules/codec/tests/codec.cpp +++ b/modules/codec/tests/codec.cpp @@ -40,12 +40,15 @@ using TestInt32Pair = Tuple< Int32 >; +using TestVoidReturnFunction = Function; + using TestCalcFunction = Function; using TestInterface = Interface< Member, Member, - Member + Member, + Member >; } SAW_TEST("One Dimensional Array") { @@ -378,7 +381,7 @@ SAW_TEST("Interface basics"){ data native; auto func_add = - [](data req){ + [](data& req){ data resp; resp.set(req.get<0>().get() + req.get<1>().get()); @@ -386,14 +389,14 @@ SAW_TEST("Interface basics"){ return resp; }; auto func_sub = - [](data req){ + [](data& req){ data resp; resp.set(req.get<0>().get() - req.get<1>().get()); return resp; }; - auto func_multiply = [](data req){ + auto func_multiply = [](data& req){ data resp; resp.set(req.get<0>().get() * req.get<1>().get()); @@ -401,7 +404,12 @@ SAW_TEST("Interface basics"){ return resp; }; - auto iface = interface_factory::create(std::move(func_add), std::move(func_sub), std::move(func_multiply)); + auto func_void = [](data& req) -> error_or { + (void) req; + return void_t{}; + }; + + auto iface = interface_factory::create(std::move(func_add), std::move(func_sub), std::move(func_multiply), std::move(func_void)); { data native; diff --git a/modules/io_codec/c++/transfer.hpp b/modules/io_codec/c++/transfer.hpp new file mode 100644 index 0000000..b6aa977 --- /dev/null +++ b/modules/io_codec/c++/transfer.hpp @@ -0,0 +1,9 @@ +#pragma once + +namespace saw { +template +class data_client; + +template +class data_server; +} diff --git a/modules/remote-sycl/.nix/derivation.nix b/modules/remote-sycl/.nix/derivation.nix index 488b8a8..2247ec0 100644 --- a/modules/remote-sycl/.nix/derivation.nix +++ b/modules/remote-sycl/.nix/derivation.nix @@ -49,7 +49,7 @@ in stdenv.mkDerivation { scons prefix=$out build_examples=${build_examples} install ''; - doCheck = false; + doCheck = true; checkPhase = '' scons test ./bin/tests diff --git a/modules/remote-sycl/c++/remote.hpp b/modules/remote-sycl/c++/remote.hpp index 54b7a7b..24756be 100644 --- a/modules/remote-sycl/c++/remote.hpp +++ b/modules/remote-sycl/c++/remote.hpp @@ -11,11 +11,30 @@ class remote; template class device; +template +class data { +private: + cl::sycl::buffer> data_; +public: + data(data& data__): + data_{&data__, 1u} + {} + + auto& get_handle() { + return data_; + } + + template + auto access(cl::sycl::handler& h){ + return data_.template get_access(h); + } +}; + /** * Remote data class for the Sycl backend. */ template -class remote_data { +class remote_data final { private: /** * An identifier to the data being held on the remote @@ -30,20 +49,15 @@ public: /** * Main constructor */ - remote_data(data& remote_data__, cl::sycl::queue& queue__): - remote_data_{&remote_data__}, + 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(){ - if(remote_data_){ - cl::sycl::free(remote_data_,queue_); - remote_data_ = nullptr; - } - } + ~remote_data(){} SAW_FORBID_COPY(remote_data); SAW_FORBID_MOVE(remote_data); @@ -82,7 +96,7 @@ private: /** * The actual data */ - data* device_data_; + data* device_data_; /** * The sycl queue object */ @@ -91,7 +105,7 @@ public: /** * Main constructor */ - device_data(data& device_data__, cl::sycl::queue& queue__): + device_data(data& device_data__, cl::sycl::queue& queue__): device_data_{&device_data__}, queue_{&queue__} {} @@ -111,6 +125,11 @@ public: }; namespace impl { +template +struct device_id_map { + std::vector> data; +}; + template struct rpc_id_map_helper { static_assert(always_false, "Only supports Interface schema types."); @@ -122,9 +141,12 @@ struct rpc_id_map_helper, Encoding, Storage> { }; } +} +// Maybe a helper impl tmpl file? +namespace saw { + /** * Represents a remote Sycl device. - * */ template<> class device final { @@ -160,11 +182,6 @@ public: } }; -/** - * Device data transport - */ - - /** * Rpc Client class for the Sycl backend. */ @@ -222,7 +239,7 @@ private: /** * Basic storage for response data. */ - // impl::rpc_id_map_helper storage_; + impl::rpc_id_map_helper storage_; public: /** @@ -369,22 +386,4 @@ public: } }; -template -template -error_or, encode::Native, rmt::Sycl>> data, encode::Native, rmt::Sycl>::copy_to_device(const data, encode::Native, Storage>& host_data, device& dev){ - /** - * Retrieve handle - */ - auto& cmd_handle = dev.get_handle(); - - uint64_t* dev_len = cl::sycl::malloc_device(1u, cmd_handle); - uint64_t len = host_data.size(); - cmd_handle.template copy(&len,dev_len, 1u); - - auto dev_dat = cl::sycl::malloc_device>(host_data.size(), cmd_handle); - cmd_handle.copy(&host_data.at(0), dev_dat, host_data.size()); - cmd_handle.wait(); - - return data,encode::Native, rmt::Sycl>{dev_len, dev_dat, cmd_handle}; -} } diff --git a/modules/remote-sycl/tests/calculator.cpp b/modules/remote-sycl/tests/calculator.cpp deleted file mode 100644 index 6d061ad..0000000 --- a/modules/remote-sycl/tests/calculator.cpp +++ /dev/null @@ -1,69 +0,0 @@ -#include - -#include "../c++/remote.hpp" - -namespace { -namespace schema { -using namespace saw::schema; -using Calculator = Interface< - Member< - Function, Int64>, "add" - > -, Member< - Function, Int64>, "multiply" - > ->; -} - -SAW_TEST("Sycl Interface Calculator"){ - using namespace saw; - - cl::sycl::queue cmd_queue; - - interface, cl::sycl::queue*> cl_iface { -[](data>& in, cl::sycl::queue* cmd) -> data { - std::array h_xy{in.get<0>().get(), in.get<1>().get()}; - int64_t res{}; - cl::sycl::buffer d_xy { h_xy.data(), h_xy.size() }; - cl::sycl::buffer d_z { &res, 1u }; - cmd->submit([&](cl::sycl::handler& h){ - auto a_xy = d_xy.get_access(h); - auto a_z = d_z.get_access(h); - - h.parallel_for(cl::sycl::range<1>(1u), [=] (cl::sycl::id<1> it){ - a_z[0] = a_xy[0] + a_xy[1]; - }); - }); - cmd->wait(); - return data{res}; - }, - [](data,encode::Native,rmt::Sycl>& in, cl::sycl::queue* cmd) -> data { - return data{in.get<0>().get() * in.get<1>().get()}; - } - }; - - int64_t x = 1; - int64_t y = -2; - - int64_t sum = x + y; - int64_t mult = x * y; - - - data> input; - input.template get<0>().set(x); - input.template get<1>().set(y); - - { - auto eov = cl_iface.template call<"add">(input, &cmd_queue); - SAW_EXPECT(eov.is_value(), "Returned error on add"); - - SAW_EXPECT(eov.get_value().get() == sum, "Addition was incorrect"); - } - { - auto eov = cl_iface.template call<"multiply">(input, &cmd_queue); - SAW_EXPECT(eov.is_value(), "Returned error on add"); - - SAW_EXPECT(eov.get_value().get() == mult, "Addition was incorrect"); - } -} -} diff --git a/modules/remote-sycl/tests/calculator.foo b/modules/remote-sycl/tests/calculator.foo new file mode 100644 index 0000000..745bd3d --- /dev/null +++ b/modules/remote-sycl/tests/calculator.foo @@ -0,0 +1,69 @@ +#include + +#include "../c++/remote.hpp" + +namespace { +namespace schema { +using namespace saw::schema; +using Calculator = Interface< + Member< + Function, Int64>, "add" + > +, Member< + Function, Int64>, "multiply" + > +>; +} + +SAW_TEST("SYCL Interface Calculator"){ + using namespace saw; + + cl::sycl::queue cmd_queue; + + interface, cl::sycl::queue*> cl_iface { +[](data>& in, cl::sycl::queue* cmd) -> data { + std::array h_xy{in.get<0>().get(), in.get<1>().get()}; + int64_t res{}; + cl::sycl::buffer d_xy { h_xy.data(), h_xy.size() }; + cl::sycl::buffer d_z { &res, 1u }; + cmd->submit([&](cl::sycl::handler& h){ + auto a_xy = d_xy.get_access(h); + auto a_z = d_z.get_access(h); + + h.parallel_for(cl::sycl::range<1>(1u), [=] (cl::sycl::id<1> it){ + a_z[0] = a_xy[0] + a_xy[1]; + }); + }); + cmd->wait(); + return data{res}; + }, + [](data,encode::Native,rmt::Sycl>& in, cl::sycl::queue* cmd) -> data { + return data{in.get<0>().get() * in.get<1>().get()}; + } + }; + + int64_t x = 1; + int64_t y = -2; + + int64_t sum = x + y; + int64_t mult = x * y; + + + data> input; + input.template get<0>().set(x); + input.template get<1>().set(y); + + { + auto eov = cl_iface.template call<"add">(input, &cmd_queue); + SAW_EXPECT(eov.is_value(), "Returned error on add"); + + SAW_EXPECT(eov.get_value().get() == sum, "Addition was incorrect"); + } + { + auto eov = cl_iface.template call<"multiply">(input, &cmd_queue); + SAW_EXPECT(eov.is_value(), "Returned error on add"); + + SAW_EXPECT(eov.get_value().get() == mult, "Addition was incorrect"); + } +} +} diff --git a/modules/remote-sycl/tests/sycl_basics.cpp b/modules/remote-sycl/tests/sycl_basics.cpp new file mode 100644 index 0000000..bf41983 --- /dev/null +++ b/modules/remote-sycl/tests/sycl_basics.cpp @@ -0,0 +1,96 @@ +#include + +#include "../c++/remote.hpp" + +namespace { +namespace schema { +using namespace saw::schema; + +using TestStruct = Struct< + Member, + Member, + Member, "doubles"> +>; + +using Foo = Interface< + Member, "foo"> +>; + +using Calculator = Interface< + Member< + Function, Int64>, "add" + > +, Member< + Function, Int64>, "multiply" + > +>; +} +SAW_TEST("SYCL Test Setup"){ + using namespace saw; + + data host_data; + host_data.template get<"foo">() = 321; + host_data.template get<"bar">() = 50.0; + host_data.template get<"doubles">() = {1024u}; + + saw::event_loop loop; + saw::wait_scope wait{loop}; + + remote rmt; + saw::own> rmt_addr{}; + + rmt.resolve_address().then([&](auto addr){ + rmt_addr = std::move(addr); + }).detach(); + + wait.poll(); + SAW_EXPECT(rmt_addr, "Remote Address class hasn't been filled"); + + auto device = rmt.connect_device(*rmt_addr); + + data device_data{host_data}; + + interface cl_iface { +[&](data& in, cl::sycl::queue* cmd) -> error_or { + + cmd->submit([&](cl::sycl::handler& h){ + + auto acc_buff = in.template access(h); + + uint64_t si = host_data.template get<"doubles">().size(); + + h.parallel_for(cl::sycl::range<1>(si), [=] (cl::sycl::id<1> it){ + acc_buff[0u].template get<"foo">() = acc_buff[0u].template get<"doubles">().size(); + auto& dbls = acc_buff[0u].template get<"doubles">(); + dbls.at(it[0u]) = it[0u] * 2.0; + }); + }); + /* + cmd->submit([&](cl::sycl::handler& h){ + auto acc_buff = in.template access(h); + h.copy(acc_buff, &host_data); + }); + */ + + /** + cl::sycl::host_accessor result{in.get_handle()}; + std::cout<().get()<().get()<()<().get()<().get()<(device_data, &device.get_handle()); + device.get_handle().wait(); + std::cout<().get()<().get()<(); + std::cout<