#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); 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 = 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 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_); } }; } 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); } }; } 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_; } }; } }