Skip to content

Commit

Permalink
Merge pull request #40087 from fwyzard/alpaka_framework
Browse files Browse the repository at this point in the history
Improve the SoA accessors and optimisations [12.5.x]
cmsbuild authored Nov 17, 2022

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature. The key has expired.
2 parents 10c050a + 067becf commit 3f6d8f9
Showing 6 changed files with 142 additions and 71 deletions.
5 changes: 3 additions & 2 deletions DataFormats/SoATemplate/README.md
Original file line number Diff line number Diff line change
@@ -66,8 +66,9 @@ Serialization of Eigen data is not yet supported.
The template shared by layouts and parameters are:
- Byte aligment (defaulting to the nVidia GPU cache line size (128 bytes))
- Alignment enforcement (`relaxed` or `enforced`). When enforced, the alignment will be checked at construction
time, and the accesses are done with compiler hinting (using the widely supported `__builtin_assume_aligned`
intrinsic).
time.~~, and the accesses are done with compiler hinting (using the widely supported `__builtin_assume_aligned`
intrinsic).~~ It turned out that hinting `nvcc` for alignement removed the benefit of more important `__restrict__`
hinting. The `__builtin_assume_aligned` is hence currently not use.

In addition, the views also provide access parameters:
- Restrict qualify: add restrict hints to read accesses, so that the compiler knows it can relax accesses to the
91 changes: 46 additions & 45 deletions DataFormats/SoATemplate/interface/SoACommon.h
Original file line number Diff line number Diff line change
@@ -287,19 +287,19 @@ namespace cms::soa {

SOA_HOST_DEVICE SOA_INLINE Ref operator()() {
// Ptr type will add the restrict qualifyer if needed
Ptr col = alignedCol();
Ptr col = col_;
return col[idx_];
}

SOA_HOST_DEVICE SOA_INLINE RefToConst operator()() const {
// PtrToConst type will add the restrict qualifyer if needed
PtrToConst col = alignedCol();
PtrToConst col = col_();
return col[idx_];
}

SOA_HOST_DEVICE SOA_INLINE Ptr operator&() { return &alignedCol()[idx_]; }
SOA_HOST_DEVICE SOA_INLINE Ptr operator&() { return &col_[idx_]; }

SOA_HOST_DEVICE SOA_INLINE PtrToConst operator&() const { return &alignedCol()[idx_]; }
SOA_HOST_DEVICE SOA_INLINE PtrToConst operator&() const { return &col_[idx_]; }

/* This was an attempt to implement the syntax
*
@@ -318,7 +318,7 @@ namespace cms::soa {
template <typename T2>
SOA_HOST_DEVICE SOA_INLINE Ref operator=(const T2& v) {
return alignedCol()[idx_] = v;
return col_[idx_] = v;
}
*/

@@ -327,13 +327,6 @@ namespace cms::soa {
static constexpr auto valueSize = sizeof(T);

private:
SOA_HOST_DEVICE SOA_INLINE Ptr alignedCol() const {
if constexpr (ALIGNMENT) {
return reinterpret_cast<Ptr>(__builtin_assume_aligned(col_, ALIGNMENT));
}
return reinterpret_cast<Ptr>(col_);
}

size_type idx_;
T* col_;
};
@@ -437,11 +430,11 @@ namespace cms::soa {

SOA_HOST_DEVICE SOA_INLINE RefToConst operator()() const {
// Ptr type will add the restrict qualifyer if needed
PtrToConst col = alignedCol();
PtrToConst col = col_;
return col[idx_];
}

SOA_HOST_DEVICE SOA_INLINE const T* operator&() const { return &alignedCol()[idx_]; }
SOA_HOST_DEVICE SOA_INLINE const T* operator&() const { return &col_[idx_]; }

/* This was an attempt to implement the syntax
*
@@ -461,13 +454,6 @@ namespace cms::soa {
static constexpr auto valueSize = sizeof(T);

private:
SOA_HOST_DEVICE SOA_INLINE PtrToConst alignedCol() const {
if constexpr (ALIGNMENT) {
return reinterpret_cast<PtrToConst>(__builtin_assume_aligned(col_, ALIGNMENT));
}
return reinterpret_cast<PtrToConst>(col_);
}

size_type idx_;
const T* col_;
};
@@ -584,47 +570,49 @@ namespace cms::soa {
/* Column accessors: templates implementing the global accesors (soa::x() and soa::x(index) */
enum class SoAAccessType : bool { mutableAccess, constAccess };

template <typename, SoAColumnType, SoAAccessType>
template <typename, SoAColumnType, SoAAccessType, byte_size_type, bool>
struct SoAColumnAccessorsImpl {};

// TODO from Eric Cano:
// - add alignment support
// - SFINAE-based const/non const variants

// Column
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::column, SoAAccessType::mutableAccess> {
//SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {}
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::column, SoAAccessType::mutableAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl<SoAColumnType::column, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE T* operator()() { return params_.addr_; }
using NoParamReturnType = T*;
using ParamReturnType = T&;
SOA_HOST_DEVICE SOA_INLINE T& operator()(size_type index) { return params_.addr_[index]; }

private:
SoAParametersImpl<SoAColumnType::column, T> params_;
};

// Const column
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::column, SoAAccessType::constAccess> {
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::column, SoAAccessType::constAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl<SoAColumnType::column, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE const T* operator()() const { return params_.addr_; }
using NoParamReturnType = const T*;
SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; }
using ParamReturnType = const T&;
SOA_HOST_DEVICE SOA_INLINE T const& operator()(size_type index) const { return params_.addr_[index]; }

private:
SoAConstParametersImpl<SoAColumnType::column, T> params_;
};

// Scalar
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::scalar, SoAAccessType::mutableAccess> {
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::scalar, SoAAccessType::mutableAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl<SoAColumnType::scalar, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE T& operator()() { return *params_.addr_; }
using NoParamReturnType = T&;
using ParamReturnType = void;
SOA_HOST_DEVICE SOA_INLINE void operator()(size_type index) const {
assert(false && "Indexed access impossible for SoA scalars.");
}
@@ -634,12 +622,13 @@ namespace cms::soa {
};

// Const scalar
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::scalar, SoAAccessType::constAccess> {
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::scalar, SoAAccessType::constAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl<SoAColumnType::scalar, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE T operator()() const { return *params_.addr_; }
using NoParamReturnType = T;
SOA_HOST_DEVICE SOA_INLINE T const& operator()() const { return *params_.addr_; }
using NoParamReturnType = T const&;
using ParamReturnType = void;
SOA_HOST_DEVICE SOA_INLINE void operator()(size_type index) const {
assert(false && "Indexed access impossible for SoA scalars.");
}
@@ -649,27 +638,32 @@ namespace cms::soa {
};

// Eigen-type
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::eigen, SoAAccessType::mutableAccess> {
//SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {}
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::eigen, SoAAccessType::mutableAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl<SoAColumnType::eigen, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE typename T::Scalar* operator()() { return params_.addr_; }
using NoParamReturnType = typename T::Scalar*;
//SOA_HOST_DEVICE SOA_INLINE T& operator()(size_type index) { return params_.addr_[index]; }
using ParamReturnType = typename SoAValue<SoAColumnType::eigen, T, alignment, restrictQualify>::MapType;
SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) {
return SoAValue<SoAColumnType::eigen, T, alignment, restrictQualify>(index, params_)();
}

private:
SoAParametersImpl<SoAColumnType::eigen, T> params_;
};

// Const Eigen-type
template <typename T>
struct SoAColumnAccessorsImpl<T, SoAColumnType::eigen, SoAAccessType::constAccess> {
template <typename T, byte_size_type alignment, bool restrictQualify>
struct SoAColumnAccessorsImpl<T, SoAColumnType::eigen, SoAAccessType::constAccess, alignment, restrictQualify> {
SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl<SoAColumnType::eigen, T>& params)
: params_(params) {}
SOA_HOST_DEVICE SOA_INLINE const typename T::Scalar* operator()() const { return params_.addr_; }
using NoParamReturnType = typename T::Scalar*;
//SOA_HOST_DEVICE SOA_INLINE T operator()(size_type index) const { return params_.addr_[index]; }
SOA_HOST_DEVICE SOA_INLINE typename T::Scalar const* operator()() const { return params_.addr_; }
using NoParamReturnType = typename T::Scalar const*;
using ParamReturnType = typename SoAValue<SoAColumnType::eigen, T, alignment, restrictQualify>::CMapType;
SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) const {
return SoAConstValue<SoAColumnType::eigen, T, alignment, restrictQualify>(index, params_)();
}

private:
SoAConstParametersImpl<SoAColumnType::eigen, T> params_;
@@ -681,8 +675,15 @@ namespace cms::soa {
template <auto columnType>
struct ColumnType {
template <auto accessType>
struct AccessType : public SoAColumnAccessorsImpl<T, columnType, accessType> {
using SoAColumnAccessorsImpl<T, columnType, accessType>::SoAColumnAccessorsImpl;
struct AccessType {
template <auto alignment>
struct Alignment {
template <auto restrictQualify>
struct RestrictQualifier
: public SoAColumnAccessorsImpl<T, columnType, accessType, alignment, restrictQualify> {
using SoAColumnAccessorsImpl<T, columnType, accessType, alignment, restrictQualify>::SoAColumnAccessorsImpl;
};
};
};
};
};
11 changes: 0 additions & 11 deletions DataFormats/SoATemplate/interface/SoALayout.h
Original file line number Diff line number Diff line change
@@ -581,17 +581,6 @@
throw std::runtime_error("In " #CLASS "::" #CLASS ": unexpected end pointer."); \
} \
\
/* Range checker conditional to the macro _DO_RANGECHECK */ \
SOA_HOST_DEVICE SOA_INLINE \
void rangeCheck(size_type index) const { \
if constexpr (_DO_RANGECHECK) { \
if (index >= elements_) { \
printf("In " #CLASS "::rangeCheck(): index out of range: %zu with elements: %zu\n", index, elements_); \
assert(false); \
} \
} \
} \
\
/* Data members */ \
std::byte* mem_; \
size_type elements_; \
58 changes: 48 additions & 10 deletions DataFormats/SoATemplate/interface/SoAView.h
Original file line number Diff line number Diff line change
@@ -92,6 +92,16 @@ namespace cms::soa {
typename BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::Metadata::BOOST_PP_CAT(ParametersTypeOf_, LAYOUT_MEMBER); \
constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) = \
BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::Metadata::BOOST_PP_CAT(ColumnTypeOf_, LAYOUT_MEMBER); \
using BOOST_PP_CAT(ConstAccessorOf_, LOCAL_NAME) = \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::constAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify> ; \
using BOOST_PP_CAT(MutableAccessorOf_, LOCAL_NAME) = \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::mutableAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify> ; \
SOA_HOST_DEVICE SOA_INLINE \
const auto BOOST_PP_CAT(parametersOf_, LOCAL_NAME)() const { \
return CAST(parent_.BOOST_PP_CAT(LOCAL_NAME, Parameters_)); \
@@ -364,18 +374,30 @@ namespace cms::soa {
SOA_HOST_DEVICE SOA_INLINE \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::mutableAccess>::NoParamReturnType \
cms::soa::SoAAccessType::mutableAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>::NoParamReturnType \
LOCAL_NAME() { \
return typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::mutableAccess>(const_cast_SoAParametersImpl( \
base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(); \
cms::soa::SoAAccessType::mutableAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>(const_cast_SoAParametersImpl( \
base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(); \
} \
SOA_HOST_DEVICE SOA_INLINE auto& LOCAL_NAME(size_type index) { \
SOA_HOST_DEVICE SOA_INLINE \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::mutableAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>::ParamReturnType \
LOCAL_NAME(size_type index) { \
if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \
if (index >= base_type::elements_) \
SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #LOCAL_NAME "(size_type index)") \
} \
return typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::mutableAccess>(const_cast_SoAParametersImpl( \
base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(index); \
cms::soa::SoAAccessType::mutableAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>(const_cast_SoAParametersImpl( \
base_type:: BOOST_PP_CAT(LOCAL_NAME, Parameters_)))(index); \
}
// clang-format on

@@ -388,15 +410,31 @@ namespace cms::soa {
// clang-format off
#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \
/* Column or scalar */ \
SOA_HOST_DEVICE SOA_INLINE auto LOCAL_NAME() const { \
SOA_HOST_DEVICE SOA_INLINE \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::constAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>::NoParamReturnType \
LOCAL_NAME() const { \
return typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \
cms::soa::SoAAccessType::constAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(); \
} \
SOA_HOST_DEVICE SOA_INLINE auto LOCAL_NAME(size_type index) const { \
SOA_HOST_DEVICE SOA_INLINE \
typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::constAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>::ParamReturnType \
LOCAL_NAME(size_type index) const { \
if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \
if (index >= elements_) \
SOA_THROW_OUT_OF_RANGE("Out of range index in const " #LOCAL_NAME "(size_type index)") \
} \
return typename cms::soa::SoAAccessors<typename BOOST_PP_CAT(Metadata::TypeOf_, LOCAL_NAME)>:: \
template ColumnType<BOOST_PP_CAT(Metadata::ColumnTypeOf_, LOCAL_NAME)>::template AccessType< \
cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \
cms::soa::SoAAccessType::constAccess>::template Alignment<conditionalAlignment>:: \
template RestrictQualifier<restrictQualify>(BOOST_PP_CAT(LOCAL_NAME, Parameters_))(index); \
}
// clang-format on

19 changes: 16 additions & 3 deletions DataFormats/SoATemplate/test/SoALayoutAndView_t.cu
Original file line number Diff line number Diff line change
@@ -101,7 +101,7 @@ using RangeCheckingHostDeviceView =

// We expect to just run one thread.
__global__ void rangeCheckKernel(RangeCheckingHostDeviceView soa) {
printf("About to fail range-check in CUDA thread: %d\n", threadIdx.x);
printf("About to fail range-check (operator[]) in CUDA thread: %d\n", threadIdx.x);
[[maybe_unused]] auto si = soa[soa.metadata().size()];
printf("Fail: range-check failure should have stopped the kernel.\n");
}
@@ -250,10 +250,23 @@ int main(void) {
soa1viewRangeChecking(h_soahdLayout);
// This should throw an exception
[[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()];
std::cout << "Fail: expected range-check exception not caught on the host." << std::endl;
std::cout << "Fail: expected range-check exception (operator[]) not caught on the host." << std::endl;
assert(false);
} catch (const std::out_of_range&) {
std::cout << "Pass: expected range-check exception successfully caught on the host." << std::endl;
std::cout << "Pass: expected range-check exception (operator[]) successfully caught on the host." << std::endl;
}

try {
// Get a view like the default, except for range checking
SoAHostDeviceLayout::ViewTemplate<SoAHostDeviceView::restrictQualify, cms::soa::RangeChecking::enabled>
soa1viewRangeChecking(h_soahdLayout);
// This should throw an exception
[[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.metadata().size()];
std::cout << "Fail: expected range-check exception (view-level index access) not caught on the host." << std::endl;
assert(false);
} catch (const std::out_of_range&) {
std::cout << "Pass: expected range-check exception (view-level index access) successfully caught on the host."
<< std::endl;
}

// Validation of range checking in a kernel
29 changes: 29 additions & 0 deletions HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc
Original file line number Diff line number Diff line change
@@ -48,6 +48,29 @@ namespace {
column.print(out);
return out;
}

template <typename T>
void checkViewAddresses(T const& view) {
assert(view.metadata().addressOf_x() == view.x());
assert(view.metadata().addressOf_x() == &view.x(0));
assert(view.metadata().addressOf_x() == &view[0].x());
assert(view.metadata().addressOf_y() == view.y());
assert(view.metadata().addressOf_y() == &view.y(0));
assert(view.metadata().addressOf_y() == &view[0].y());
assert(view.metadata().addressOf_z() == view.z());
assert(view.metadata().addressOf_z() == &view.z(0));
assert(view.metadata().addressOf_z() == &view[0].z());
assert(view.metadata().addressOf_id() == view.id());
assert(view.metadata().addressOf_id() == &view.id(0));
assert(view.metadata().addressOf_id() == &view[0].id());
assert(view.metadata().addressOf_m() == view.m());
assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0, 0));
assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0));
assert(view.metadata().addressOf_r() == &view.r());
//assert(view.metadata().addressOf_r() == &view.r(0)); // cannot access a scalar with an index
//assert(view.metadata().addressOf_r() == &view[0].r()); // cannot access a scalar via a SoA row-like accessor
}

} // namespace

class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> {
@@ -58,6 +81,8 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> {
void analyze(edm::Event const& event, edm::EventSetup const&) override {
portabletest::TestHostCollection const& product = event.get(token_);
auto const& view = product.const_view();
auto& mview = product.view();
auto const& cmview = product.view();

{
edm::LogInfo msg("TestAlpakaAnalyzer");
@@ -88,6 +113,10 @@ class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> {
reinterpret_cast<intptr_t>(view.metadata().addressOf_r());
}

checkViewAddresses(view);
checkViewAddresses(mview);
checkViewAddresses(cmview);

const portabletest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}};
assert(view.r() == 1.);
for (int32_t i = 0; i < view.metadata().size(); ++i) {

0 comments on commit 3f6d8f9

Please sign in to comment.