From 1d578450dc82843bd4b24f3a6aad2c1a82bbda5e Mon Sep 17 00:00:00 2001 From: Claudius 'keldu' Holeksa Date: Tue, 17 Sep 2024 11:21:05 +0200 Subject: Managed to get hip to compile --- modules/codec/c++/data.hpp | 4 +- modules/codec/c++/data_raw.hpp | 114 +++++++++++----------- modules/core/c++/common.hpp | 6 +- modules/core/c++/error.hpp | 8 +- modules/remote-hip/c++/data.hpp | 8 +- modules/remote-hip/c++/device.tmpl.hpp | 2 +- modules/remote-hip/examples/hip_transfer_data.cpp | 75 ++++++++++---- 7 files changed, 126 insertions(+), 91 deletions(-) diff --git a/modules/codec/c++/data.hpp b/modules/codec/c++/data.hpp index b14d785..33f301f 100644 --- a/modules/codec/c++/data.hpp +++ b/modules/codec/c++/data.hpp @@ -675,11 +675,11 @@ class data, encode::Native> { static_assert(sizeof...(Dims)==Dim, "Argument size must be equal to the Dimension"); } - data& at(const std::array& ind){ + constexpr data& at(const std::array& ind){ return value_.at(this->get_flat_index(ind)); } - const data& at(const std::array& ind) const { + constexpr const data& at(const std::array& ind) const { return value_.at(this->get_flat_index(ind)); } diff --git a/modules/codec/c++/data_raw.hpp b/modules/codec/c++/data_raw.hpp index db57af3..2da49b4 100644 --- a/modules/codec/c++/data_raw.hpp +++ b/modules/codec/c++/data_raw.hpp @@ -62,54 +62,54 @@ public: private: typename native_data_type::type value_; public: - data():value_{}{} - data(data):value_{}{} + constexpr data():value_{}{} + constexpr data(data):value_{}{} SAW_DEFAULT_COPY(data); SAW_DEFAULT_MOVE(data); - data(typename native_data_type::type value__): + constexpr data(typename native_data_type::type value__): value_{std::move(value__)}{} - void set(typename native_data_type::type val){ + constexpr void set(typename native_data_type::type val){ value_ = val; } - typename native_data_type::type get() const {return value_;} + constexpr typename native_data_type::type get() const {return value_;} - data operator*(const data& rhs)const{ + constexpr data operator*(const data& rhs)const{ return {get() * rhs.get()}; } - data operator/(const data& rhs)const{ + constexpr data operator/(const data& rhs)const{ return {get() / rhs.get()}; } - data operator+(const data& rhs)const{ + constexpr data operator+(const data& rhs)const{ return {get() + rhs.get()}; } - data& operator+=(const data& rhs)const{ + constexpr data& operator+=(const data& rhs)const{ value_ += rhs.get(); return *this; } - data& operator-=(const data& rhs)const{ + constexpr data& operator-=(const data& rhs)const{ value_ -= rhs.get(); return *this; } - data operator-(const data& rhs)const{ + constexpr data operator-(const data& rhs)const{ return {get() - rhs.get()}; } template - bool operator==(const data& rhs)const{ + constexpr bool operator==(const data& rhs)const{ return get() == rhs.get(); } template - bool operator<(const data& rhs) const { + constexpr bool operator<(const data& rhs) const { return get() < rhs.get(); } @@ -117,7 +117,7 @@ public: * Casts */ template - data cast_to() const { + constexpr data cast_to() const { auto raw_to = static_cast::type>(value_); return {raw_to}; } @@ -132,64 +132,64 @@ public: private: ref::type > value_; public: - data(ref::type > value__): + constexpr data(ref::type > value__): value_{value__} {} SAW_DEFAULT_COPY(data); SAW_FORBID_MOVE(data); - void set(typename native_data_type::type val){ + constexpr void set(typename native_data_type::type val){ value_() = val; } - typename native_data_type::type get() const { + constexpr typename native_data_type::type get() const { return value_(); } - data operator*(const data& rhs)const{ + constexpr data operator*(const data& rhs)const{ return {get() * rhs.get()}; } - data operator/(const data& rhs)const{ + constexpr data operator/(const data& rhs)const{ return {get() / rhs.get()}; } - data operator+(const data& rhs)const{ + constexpr data operator+(const data& rhs)const{ return {get() + rhs.get()}; } - data& operator+=(const data& rhs)const{ + constexpr data& operator+=(const data& rhs)const{ value_ += rhs.get(); return *this; } - data& operator-=(const data& rhs)const{ + constexpr data& operator-=(const data& rhs)const{ value_ -= rhs.get(); return *this; } - data operator-(const data& rhs)const{ + constexpr data operator-(const data& rhs)const{ return {get() - rhs.get()}; } template - bool operator==(const data& rhs)const{ + constexpr bool operator==(const data& rhs)const{ return get() == rhs.get(); } template - bool operator==(const data& rhs)const{ + constexpr bool operator==(const data& rhs)const{ return get() == rhs.get(); } template - bool operator<(const data& rhs) const { + constexpr bool operator<(const data& rhs) const { return get() < rhs.get(); } template - bool operator<(const data& rhs) const { + constexpr bool operator<(const data& rhs) const { return get() < rhs.get(); } }; @@ -210,70 +210,70 @@ public: data(typename saw::native_data_type::type val__):value_{static_cast::type>(val__)}{} - typename saw::native_data_type::type get() const { + constexpr typename saw::native_data_type::type get() const { return value_.template cast_to().get(); } - data(const saw::data& val){ + constexpr data(const saw::data& val){ value_ = val.template cast_to(); } - void set(typename saw::native_data_type::type val){ + constexpr void set(typename saw::native_data_type::type val){ value_.set(static_cast(val)); } - data operator*(const data& rhs) const { + constexpr data operator*(const data& rhs) const { using CalcType = typename native_data_type::type; CalcType left = static_cast(value_.get()); CalcType right = static_cast(rhs.get()); return {left * right}; } - data operator*(const data& rhs) const { + constexpr data operator*(const data& rhs) const { using CalcType = typename native_data_type::type; CalcType left = static_cast(value_.get()); CalcType right = rhs.get(); return {left * right}; } - data operator/(const data& rhs)const{ + constexpr data operator/(const data& rhs)const{ using CalcType = typename native_data_type::type; CalcType left = static_cast(value_.get()); CalcType right = static_cast(rhs.get()); return {left / right}; } - data operator+(const data& rhs)const{ + constexpr data operator+(const data& rhs)const{ using CalcType = typename native_data_type::type; CalcType left = static_cast(value_.get()); CalcType right = static_cast(rhs.get()); return {left + right}; } - data& operator+=(const data& rhs)const{ + constexpr data& operator+=(const data& rhs)const{ *this = *this + rhs.get(); return *this; } - data operator-(const data& rhs)const{ + constexpr data operator-(const data& rhs)const{ using CalcType = typename native_data_type::type; CalcType left = static_cast(value_.get()); CalcType right = static_cast(rhs.get()); return {left - right}; } - data& operator-=(const data& rhs) const { + constexpr data& operator-=(const data& rhs) const { *this = *this - rhs.get(); return *this; } template - bool operator==(const data& rhs)const{ + constexpr bool operator==(const data& rhs)const{ return get() == rhs.get(); } template - bool operator<(const data& rhs) const { + constexpr bool operator<(const data& rhs) const { return get() < rhs.get(); } }; @@ -415,14 +415,14 @@ class data, encode::NativeRaw> { typename raw_native_array_type_helper::Type* value_; uint64_t value_size_; - uint64_t get_full_size() const { + constexpr uint64_t get_full_size() const { uint64_t s = 1; for(uint64_t iter = 0; iter < Dim; ++iter){ - assert(dims_.at(iter) > 0); - s *= dims_.at(iter); + s *= dims_[iter]; } + return s; } public: @@ -438,7 +438,7 @@ class data, encode::NativeRaw> { /** * Purely internal function for working C interfacing */ - typename raw_native_array_type_helper::Type* get_raw_data() { + constexpr typename raw_native_array_type_helper::Type* get_raw_data() { if(value_size_ == 0u){ return nullptr; } @@ -477,54 +477,53 @@ class data, encode::NativeRaw> { } template - data(Dims... size_): + constexpr data(Dims... size_): data{{static_cast(size_)...}} { static_assert(sizeof...(Dims)==Dim, "Argument size must be equal to the Dimension"); } - data, encode::NativeRaw> at(const std::array& ind){ + constexpr data, encode::NativeRaw> at(const std::array& ind){ return {value_[this->get_flat_index(ind)]}; } - const data, encode::NativeRaw> at(const std::array& ind) const { + constexpr const data, encode::NativeRaw> at(const std::array& ind) const { return {value_[this->get_flat_index(ind)]}; } template - data, encode::NativeRaw> at(Dims... i){ + constexpr data, encode::NativeRaw> at(Dims... i){ return {value_[this->get_flat_index(std::array{static_cast(i)...})]}; } template - const data, encode::NativeRaw> at(Dims... i) const { + constexpr const data, encode::NativeRaw> at(Dims... i) const { return {value_[this->get_flat_index(std::array{static_cast(i)...})]}; } template - data,encode::NativeRaw> at(const data, Encoding>& i){ + constexpr data,encode::NativeRaw> at(const data, Encoding>& i){ return {value_[this->get_flat_index(i)]}; } template - const data,encode::NativeRaw> at(const data, Encoding>& i)const{ + constexpr const data,encode::NativeRaw> at(const data, Encoding>& i)const{ return {value_[this->get_flat_index(i)]}; } - std::size_t get_dim_size(uint64_t i) const { + constexpr std::size_t get_dim_size(uint64_t i) const { return dims_.at(i); } - uint64_t size() const { return value_size_;} + constexpr uint64_t size() const { return value_size_;} - data, encode::NativeRaw> get_dims() const { + constexpr data, encode::NativeRaw> get_dims() const { return {dims_}; } private: template - uint64_t get_flat_index(const U& i) const { - assert(value_size_ == get_full_size()); + constexpr uint64_t get_flat_index(const U& i) const { uint64_t s = 0; uint64_t stride = 1; @@ -539,10 +538,9 @@ private: }else{ static_assert(always_false, "Cases exhausted"); } - }(i.at(iter)); - assert(ind < dims_.at(iter)); + }(i[iter]); s += ind * stride; - stride *= dims_.at(iter); + stride *= dims_[iter]; } return s; diff --git a/modules/core/c++/common.hpp b/modules/core/c++/common.hpp index f63c531..de464b8 100644 --- a/modules/core/c++/common.hpp +++ b/modules/core/c++/common.hpp @@ -90,7 +90,7 @@ public: /** * Main constructor. */ - ref(T& ref__): + constexpr ref(T& ref__): ref_{&ref__} {} @@ -99,14 +99,14 @@ public: /** * Operator retrieving the itself. */ - T& operator()(){ + constexpr T& operator()(){ return *ref_; } /** * Operator retrieving the itself. */ - const T& operator()() const { + constexpr const T& operator()() const { return *ref_; } }; diff --git a/modules/core/c++/error.hpp b/modules/core/c++/error.hpp index dab297e..ac45fbc 100644 --- a/modules/core/c++/error.hpp +++ b/modules/core/c++/error.hpp @@ -18,7 +18,7 @@ namespace saw { * critical and recoverable errors. Additional code ids can be provided to the * constructor if additional distinctions are necessary. */ -class error { +class error final { public: /** * Type alias @@ -101,7 +101,7 @@ namespace impl { /** * Internal registry for all error types. Is generated dynamically. */ -class error_registry { +class error_registry final { private: struct error_info { error_info() = delete; @@ -301,6 +301,10 @@ public: } }; +/** + * This tries to catch cases where error starts including itself as a type which can happen in more complicated cases. + * So this acts as a type safe guard. + */ template class error_or> { private: error_or() = delete; diff --git a/modules/remote-hip/c++/data.hpp b/modules/remote-hip/c++/data.hpp index 3e7c3ed..5d3635f 100644 --- a/modules/remote-hip/c++/data.hpp +++ b/modules/remote-hip/c++/data.hpp @@ -8,16 +8,16 @@ 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> { +template +class data> { private: - data* data_; + data* data_; public: data(): data_{nullptr} {} - data** get_device_data() { + data** get_device_data() { return &data_; } }; diff --git a/modules/remote-hip/c++/device.tmpl.hpp b/modules/remote-hip/c++/device.tmpl.hpp index 6edf431..0517f67 100644 --- a/modules/remote-hip/c++/device.tmpl.hpp +++ b/modules/remote-hip/c++/device.tmpl.hpp @@ -33,7 +33,7 @@ struct hip_copy_to_device, Encoding> { static error_or apply(data& from, data** to){ typename native_data_type::type* dat{}; hipError_t data_malloc_err = hipMalloc(&dat,sizeof(typename native_data_type::type) * from.size()); - hipError_t data_copy_err = hipMemcpy(&dat, &(from.get_raw_data()),sizeof(typename native_data_type::type) * from.size(), hipMemcpyHostToDevice); + hipError_t data_copy_err = hipMemcpy(&dat, (from.get_raw_data()),sizeof(typename native_data_type::type) * from.size(), hipMemcpyHostToDevice); if(from.size() == 0u){ // Everything is fine. We just don't want to allocate data which doesn't exist. diff --git a/modules/remote-hip/examples/hip_transfer_data.cpp b/modules/remote-hip/examples/hip_transfer_data.cpp index 18c82df..3f02d87 100644 --- a/modules/remote-hip/examples/hip_transfer_data.cpp +++ b/modules/remote-hip/examples/hip_transfer_data.cpp @@ -1,20 +1,22 @@ +#include + #include "../c++/remote.hpp" #include "../c++/transfer.hpp" #include -__global__ void print_value(saw::data* val){ +__global__ void print_value(saw::data* val){ int v = val->get(); printf("Hello world: %d\n", v); } -__global__ void print_array_vals(saw::data* val){ +__global__ void print_array_vals(saw::data, saw::encode::NativeRaw>* val){ uint64_t orig_len = val->size(); long len = (long) orig_len; printf("Array size: %ld\n", len); - for(uint64_t i = 0; i < orig_len; +i){ - int v = val->at(i); + for(uint64_t i = 0; i < orig_len; ++i){ + int v = val->at(i).get(); printf("%d ", v); } printf("\n"); @@ -35,28 +37,59 @@ saw::error_or real_main(){ } auto& addr = eo_addr.get_value(); - auto eo_dat_srv = rmt.data_listen(*addr); - if(eo_dat_srv.is_error()){ - return std::move(eo_dat_srv.get_error()); - } - auto& dat_srv = eo_dat_srv.get_value(); + { + auto eo_dat_srv = rmt.data_listen(*addr); + if(eo_dat_srv.is_error()){ + return std::move(eo_dat_srv.get_error()); + } + auto& dat_srv = eo_dat_srv.get_value(); - data val{42}; - id id_val{0u}; - auto eo_send = dat_srv->send(val, id_val); - if(eo_send.is_error()){ - return std::move(eo_send.get_error()); - } + data val{42}; + id id_val{0u}; + auto eo_send = dat_srv->send(val, id_val); + if(eo_send.is_error()){ + return std::move(eo_send.get_error()); + } + + auto eo_dfind = dat_srv->find(id_val); + if(eo_dfind.is_error()){ + return std::move(eo_dfind.get_error()); + } + auto dfind = eo_dfind.get_value(); - auto eo_dfind = dat_srv->find(id_val); - if(eo_dfind.is_error()){ - return std::move(eo_dfind.get_error()); + auto& v = dfind(); + + print_value<<>>(*(v.get_device_data())); } - auto dfind = eo_dfind.get_value(); - auto& v = dfind(); + { + auto eo_dat_srv = rmt.data_listen, encode::NativeRaw>(*addr); + if(eo_dat_srv.is_error()){ + return std::move(eo_dat_srv.get_error()); + } + auto& dat_srv = eo_dat_srv.get_value(); + + data,encode::NativeRaw> val{4}; + val.at(0u).set(5); + val.at(1u).set(3); + val.at(2u).set(-6); + val.at(3u).set(1); + id> id_val{0u}; + auto eo_send = dat_srv->send(val, id_val); + if(eo_send.is_error()){ + return std::move(eo_send.get_error()); + } - print_value<<>>(*(v.get_device_data())); + auto eo_dfind = dat_srv->find(id_val); + if(eo_dfind.is_error()){ + return std::move(eo_dfind.get_error()); + } + auto dfind = eo_dfind.get_value(); + + auto& v = dfind(); + + print_array_vals<<>>(*(v.get_device_data())); + } return make_void(); } -- cgit v1.2.3