#include "common.hpp" #include namespace kel { namespace lbm { namespace encode { template struct Sycl { }; } /* namespace impl { template struct struct_has_only_equal_dimension_array{}; } */ } } namespace saw { template class data, kel::lbm::encode::Sycl> final { public: using Schema = schema::FixedArray; private: acpp::sycl::queue* q_; data* values_; SAW_FORBID_COPY(data); SAW_DEFAULT_MOVE(data); public: data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} { values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); SAW_ASSERT(values_ and q_); } ~data(){ if(not values_){ return; } SAW_ASSERT(q_); acpp::sycl::free(values_,*q_); values_ = nullptr; } static constexpr data> get_dims() { return saw::data>{{Dims...}}; } constexpr data& at(data>& index){ return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } constexpr data& at(data>& index) const{ return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } constexpr data* flat_data() { return values_; } }; template class data>, kel::lbm::encode::Sycl> final { public: using Schema = schema::Ptr>; private: data* values_; SAW_DEFAULT_COPY(data); SAW_DEFAULT_MOVE(data); public: data(): values_{nullptr} {} data(data* values__): values_{values__} {} static constexpr data> get_dims() { return saw::data>{{Dims...}}; } constexpr data& at(data>& index){ return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } constexpr data& at(data>& index) const{ return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } constexpr data* flat_data() { return values_; } }; template class data,kel::lbm::encode::Sycl> final { public: using Schema = kel::lbm::sch::Chunk; private: using InnerSchema = typename Schema::InnerSchema; using ValueSchema = typename InnerSchema::ValueType; data> values_; public: data(acpp::sycl::queue& q__): values_{q__} {} data>& ghost_at(const data>& index){ return values_.at(index); } data>& ghost_at(const data>& index) const { return values_.at(index); } static constexpr auto get_ghost_dims() { return data>::get_dims(); } data& at(const data>& index){ std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; } return values_.at(ind); } data& at(const data>& index) const { std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; } return values_.at(ind); } static constexpr auto get_dims(){ return data,kel::lbm::encode::Sycl>{{Sides...}}; } auto flat_data(){ return values_.flat_data(); } static constexpr auto flat_size() { return data>::flat_size(); } }; template class data>,kel::lbm::encode::Sycl> final { public: using Schema = schema::Ptr>; private: using InnerSchema = typename kel::lbm::sch::Chunk::InnerSchema; using ValueSchema = typename InnerSchema::ValueType; data,kel::lbm::encode::Sycl> values_; SAW_DEFAULT_MOVE(data); SAW_DEFAULT_COPY(data); public: data(): values_{nullptr} {} data(data* values__): values_{values__} {} data>& ghost_at(const data>& index){ return values_.at(index); } data>& ghost_at(const data>& index) const { return values_.at(index); } static constexpr auto get_ghost_dims() { return data>::get_dims(); } data& at(const data>& index){ std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; } return values_.at(ind); } data& at(const data>& index) const { std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; } return values_.at(ind); } static constexpr auto get_dims(){ return data,kel::lbm::encode::Sycl>{{Sides...}}; } auto flat_data(){ return values_.flat_data(); } static constexpr auto flat_size() { return data>::flat_size(); } }; template class data, kel::lbm::encode::Sycl > final { public: using StorageT = std::tuple>...>; using Schema = schema::Struct; private: /** * @todo Check by static assert that the members all have the same dimensions. Alternatively * Do it here by specializing. */ StorageT members_; template constexpr data(acpp::sycl::queue& q, std::index_sequence): members_{(static_cast(Is), q)...} {} public: data(acpp::sycl::queue& q__): data{q__, std::make_index_sequence{}} { q__.wait(); } template auto& get(){ return std::get::value>(members_); } template const auto& get() const { return std::get::value>(members_); } }; template class data...>>, kel::lbm::encode::Sycl > final { public: using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; using Schema = schema::Struct,Keys>...>; private: /** * @todo Check by static assert that the members all have the same dimensions. Alternatively * Do it here by specializing. */ StorageT members_; public: data() = default; template auto& get(){ return std::get::value>(members_); } template const auto& get() const { return std::get::value>(members_); } }; } namespace kel { namespace lbm { namespace impl { template struct sycl_copy_helper; template struct sycl_copy_helper, Encode> final { using Schema = sch::Struct; template static saw::error_or copy_to_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type::type; auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); auto host_ptr = host_member_data.flat_data(); auto sycl_ptr = sycl_member_data.flat_data(); static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); q.submit([&](acpp::sycl::handler& h){ h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get()); }).wait(); return copy_to_device_member(host_data,sycl_data,q); } return saw::make_void(); } static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ return copy_to_device_member<0u>(host_data, sycl_data, q); } template static saw::error_or copy_to_host_member(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type::type; auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); auto host_ptr = host_member_data.flat_data(); auto sycl_ptr = sycl_member_data.flat_data(); q.submit([&](acpp::sycl::handler& h){ h.copy(sycl_ptr,host_ptr, host_member_data.flat_size().get()); }).wait(); return copy_to_host_member(sycl_data,host_data,q); } return saw::make_void(); } static saw::error_or copy_to_host(saw::data& sycl_data, saw::data>& host_data, sycl::queue& q){ return copy_to_host_member<0u>(sycl_data, host_data, q); } }; template struct make_chunk_struct_view_helper; template struct make_chunk_struct_view_helper...>, encode::Sycl> { private: template static void apply_i( saw::data...>, encode::Sycl>& dat, saw::data...>>, encode::Sycl>& dat_view) { if constexpr (i < sizeof...(Sch)){ using M = typename saw::parameter_pack_type...>::type; dat_view.template get() = {dat.template get().flat_data()}; apply_i(dat,dat_view); } } public: static auto apply(saw::data...>, encode::Sycl>& dat){ saw::data...>>, encode::Sycl> str_view; apply_i<0u>(dat,str_view); return str_view; } }; } // Ptr => Ptr>> => Ptr template auto make_chunk_struct_view(saw::data& dat){ return impl::make_chunk_struct_view_helper::apply(dat); } class device final { private: sycl::queue q_; SAW_FORBID_COPY(device); SAW_FORBID_MOVE(device); public: device() = default; ~device() = default; template saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data){ return impl::sycl_copy_helper::copy_to_device(host_data, sycl_data, q_); } template saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data){ return impl::sycl_copy_helper::copy_to_host(sycl_data, host_data, q_); } auto& get_handle(){ return q_; } }; } }