Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Read SoA scalars and Eigen columns from a ROOT streamer #39532

Merged
merged 8 commits into from Oct 1, 2022
Merged

Read SoA scalars and Eigen columns from a ROOT streamer #39532

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
e5736f4
Add missing includes
makortel Sep 20, 2022
60ec4f9
Implement the ROOT streamer for reading back SoA scalars
fwyzard Sep 29, 2022
fe1038a
Make the TestAlpakaAnalyzer and CUDA TestPortableAnalyzer more verbose
fwyzard Sep 28, 2022
74a1321
Check that scalars are set and read back correctly
fwyzard Sep 29, 2022
7cae10b
Make the CUDA and Alpaka write/read tests fail on errors
fwyzard Sep 28, 2022
566ce1b
Add ROOT streaming support for Eigen SoA columns
ericcano Sep 29, 2022
f9683f1
Add missing dependency
fwyzard Sep 29, 2022
954edbc
Add missing typename specifier
fwyzard Sep 29, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
5 changes: 4 additions & 1 deletion CUDADataFormats/PortableTestObjects/interface/TestDeviceCollection.h
View file
Open in desktop
Expand Up @@ -6,7 +6,10 @@

namespace cudatest {

// SoA with x, y, z, id fields in device global memory
// Eigen matrix
using Matrix = portabletest::Matrix;

// SoA with x, y, z, id fields, r scalar, m matrix, in device global memory
using TestDeviceCollection = cms::cuda::PortableDeviceCollection<portabletest::TestSoA>;

} // namespace cudatest
Expand Down
5 changes: 4 additions & 1 deletion CUDADataFormats/PortableTestObjects/interface/TestHostCollection.h
View file
Open in desktop
Expand Up @@ -6,7 +6,10 @@

namespace cudatest {

// SoA with x, y, z, id fields in host memory
// Eigen matrix
using Matrix = portabletest::Matrix;

// SoA with x, y, z, id fields, r scalar, m matrix, in host memory
using TestHostCollection = cms::cuda::PortableHostCollection<portabletest::TestSoA>;

} // namespace cudatest
Expand Down
2 changes: 2 additions & 0 deletions DataFormats/PortableTestObjects/BuildFile.xml
View file
Open in desktop
@@ -1,6 +1,8 @@
<use name="rootcore"/>
<use name="eigen"/>
<use name="DataFormats/Common"/>
<use name="DataFormats/Portable"/>
<use name="DataFormats/SoATemplate"/>
<use name="HeterogeneousCore/AlpakaInterface"/>
<flags ALPAKA_BACKENDS="1"/>
<export>
Expand Down
9 changes: 8 additions & 1 deletion DataFormats/PortableTestObjects/interface/TestSoA.h
View file
Open in desktop
@@ -1,12 +1,16 @@
#ifndef DataFormats_PortableTestObjects_interface_TestSoA_h
#define DataFormats_PortableTestObjects_interface_TestSoA_h

#include <Eigen/Core>
#include <Eigen/Dense>

#include "DataFormats/SoATemplate/interface/SoACommon.h"
#include "DataFormats/SoATemplate/interface/SoALayout.h"
#include "DataFormats/SoATemplate/interface/SoAView.h"

namespace portabletest {

using Matrix = Eigen::Matrix<double, 3, 6>;
// SoA layout with x, y, z, id fields
GENERATE_SOA_LAYOUT(TestSoALayout,
// columns: one value per element
Expand All @@ -15,7 +19,10 @@ namespace portabletest {
SOA_COLUMN(double, z),
SOA_COLUMN(int32_t, id),
// scalars: one value for the whole structure
SOA_SCALAR(double, r))
SOA_SCALAR(double, r),
// Eigen columns
// the typedef is needed because commas confuse macros
SOA_EIGEN_COLUMN(Matrix, m))

using TestSoA = TestSoALayout<>;

Expand Down
1 change: 1 addition & 0 deletions DataFormats/PortableTestObjects/src/classes_def.xml
View file
Open in desktop
Expand Up @@ -6,6 +6,7 @@
<field name="y_" comment="[elements_]"/>
<field name="z_" comment="[elements_]"/>
<field name="id_" comment="[elements_]"/>
<field name="m_" comment="[mElementsWithPadding_]"/>
<field name="r_" comment="[scalar_]"/>
</class>
<class name="portabletest::TestSoA::View"/>
Expand Down
4 changes: 2 additions & 2 deletions DataFormats/SoATemplate/interface/SoACommon.h
View file
Open in desktop
Expand Up @@ -390,7 +390,7 @@ namespace cms::soa {
}

using ValueType = typename C::Scalar;
static constexpr auto valueSize = sizeof(C::Scalar);
static constexpr auto valueSize = sizeof(typename C::Scalar);
SOA_HOST_DEVICE SOA_INLINE byte_size_type stride() const { return stride_; }

private:
Expand Down Expand Up @@ -503,7 +503,7 @@ namespace cms::soa {
SOA_HOST_DEVICE SOA_INLINE const C* operator&() const { return &cVal_; }

using ValueType = typename C::Scalar;
static constexpr auto valueSize = sizeof(C::Scalar);
static constexpr auto valueSize = sizeof(typename C::Scalar);

SOA_HOST_DEVICE SOA_INLINE byte_size_type stride() const { return stride_; }

Expand Down
24 changes: 15 additions & 9 deletions DataFormats/SoATemplate/interface/SoALayout.h
View file
Open in desktop
Expand Up @@ -7,7 +7,6 @@
*/

#include <cassert>
#include <iostream>

#include "SoACommon.h"
#include "SoAView.h"
Expand Down Expand Up @@ -173,6 +172,7 @@
/* Column */ \
(BOOST_PP_CAT(NAME, _)(nullptr)), \
/* Eigen column */ \
(BOOST_PP_CAT(NAME, ElementsWithPadding_)(0)) \
(BOOST_PP_CAT(NAME, _)(nullptr)) \
(BOOST_PP_CAT(NAME, Stride_)(0)) \
)
Expand All @@ -189,6 +189,7 @@
/* Column */ \
(BOOST_PP_CAT(NAME, _){other.BOOST_PP_CAT(NAME, _)}), \
/* Eigen column */ \
(BOOST_PP_CAT(NAME, ElementsWithPadding_){other.BOOST_PP_CAT(NAME, ElementsWithPadding_)}) \
(BOOST_PP_CAT(NAME, _){other.BOOST_PP_CAT(NAME, _)}) \
(BOOST_PP_CAT(NAME, Stride_){other.BOOST_PP_CAT(NAME, Stride_)}) \
)
Expand All @@ -201,12 +202,13 @@
#define _DECLARE_MEMBER_ASSIGNMENT_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \
_SWITCH_ON_TYPE(VALUE_TYPE, \
/* Scalar */ \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _);, \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _);, \
/* Column */ \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _);, \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _);, \
/* Eigen column */ \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _); \
BOOST_PP_CAT(NAME, Stride_) = other.BOOST_PP_CAT(NAME, Stride_); \
BOOST_PP_CAT(NAME, ElementsWithPadding_) = other.BOOST_PP_CAT(NAME, ElementsWithPadding_); \
BOOST_PP_CAT(NAME, _) = other.BOOST_PP_CAT(NAME, _); \
BOOST_PP_CAT(NAME, Stride_) = other.BOOST_PP_CAT(NAME, Stride_); \
)
// clang-format on

Expand Down Expand Up @@ -279,11 +281,13 @@
curMem += cms::soa::alignSize(elements_ * sizeof(CPP_TYPE), alignment); \
, \
/* Eigen column */ \
BOOST_PP_CAT(NAME, Stride_) = cms::soa::alignSize(elements_ * sizeof(CPP_TYPE::Scalar), alignment) \
/ sizeof(CPP_TYPE::Scalar); \
BOOST_PP_CAT(NAME, ElementsWithPadding_) = BOOST_PP_CAT(NAME, Stride_) \
* CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime; \
BOOST_PP_CAT(NAME, _) = reinterpret_cast<CPP_TYPE::Scalar*>(curMem); \
curMem += cms::soa::alignSize(elements_ * sizeof(CPP_TYPE::Scalar), alignment) * CPP_TYPE::RowsAtCompileTime \
* CPP_TYPE::ColsAtCompileTime; \
BOOST_PP_CAT(NAME, Stride_) = cms::soa::alignSize(elements_ * sizeof(CPP_TYPE::Scalar), alignment) \
/ sizeof(CPP_TYPE::Scalar); \
) \
if constexpr (alignmentEnforcement == AlignmentEnforcement::enforced) \
if (reinterpret_cast<intptr_t>(BOOST_PP_CAT(NAME, _)) % alignment) \
Expand Down Expand Up @@ -361,13 +365,14 @@
#define _STREAMER_READ_SOA_DATA_MEMBER_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \
_SWITCH_ON_TYPE(VALUE_TYPE, \
/* Scalar */ \
/* TODO: implement*/ \
memcpy(BOOST_PP_CAT(NAME, _), onfile.BOOST_PP_CAT(NAME, _), sizeof(CPP_TYPE)); \
, \
/* Column */ \
memcpy(BOOST_PP_CAT(NAME, _), onfile.BOOST_PP_CAT(NAME, _), sizeof(CPP_TYPE) * onfile.elements_); \
, \
/* Eigen column */ \
/* TODO: implement*/ \
memcpy(BOOST_PP_CAT(NAME, _), onfile.BOOST_PP_CAT(NAME, _), \
sizeof(CPP_TYPE::Scalar) * BOOST_PP_CAT(NAME, ElementsWithPadding_)); \
)
// clang-format on

Expand All @@ -387,6 +392,7 @@
CPP_TYPE * BOOST_PP_CAT(NAME, _) = nullptr; \
, \
/* Eigen column */ \
size_type BOOST_PP_CAT(NAME, ElementsWithPadding_); /* For ROOT serialization, (displikes the default value) */ \
CPP_TYPE::Scalar * BOOST_PP_CAT(NAME, _) = nullptr; \
byte_size_type BOOST_PP_CAT(NAME, Stride_) = 0; \
)
Expand Down
1 change: 1 addition & 0 deletions HeterogeneousCore/AlpakaCore/interface/ContextState.h
View file
Open in desktop
Expand Up @@ -7,6 +7,7 @@

#include <alpaka/alpaka.hpp>

#include "FWCore/Utilities/interface/Exception.h"
#include "HeterogeneousCore/AlpakaInterface/interface/ScopedContextFwd.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"

Expand Down
1 change: 1 addition & 0 deletions HeterogeneousCore/AlpakaCore/interface/chooseDevice.h
View file
Open in desktop
@@ -1,6 +1,7 @@
#ifndef HeterogeneousCore_AlpakaCore_interface_chooseDevice_h
#define HeterogeneousCore_AlpakaCore_interface_chooseDevice_h

#include "FWCore/Utilities/interface/StreamID.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/devices.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"
Expand Down
99 changes: 76 additions & 23 deletions HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc
View file
Open in desktop
Expand Up @@ -12,39 +12,92 @@
#include "FWCore/Utilities/interface/EDGetToken.h"
#include "FWCore/Utilities/interface/InputTag.h"

namespace {

template <typename T>
class Column {
public:
Column(T const* data, size_t size) : data_(data), size_(size) {}

void print(std::ostream& out) const {
std::stringstream buffer;
buffer << "{ ";
if (size_ > 0) {
buffer << data_[0];
}
if (size_ > 1) {
buffer << ", " << data_[1];
}
if (size_ > 2) {
buffer << ", " << data_[2];
}
if (size_ > 3) {
buffer << ", ...";
}
buffer << '}';
out << buffer.str();
}

private:
T const* const data_;
size_t const size_;
};

template <typename T>
std::ostream& operator<<(std::ostream& out, Column<T> const& column) {
column.print(out);
return out;
}
} // namespace

class TestAlpakaAnalyzer : public edm::stream::EDAnalyzer<> {
public:
TestAlpakaAnalyzer(edm::ParameterSet const& config)
: source_{config.getParameter<edm::InputTag>("source")}, token_{consumes(source_)} {}

void analyze(edm::Event const& event, edm::EventSetup const&) override {
portabletest::TestHostCollection const& product = event.get(token_);

auto const& view = product.const_view();
for (int32_t i = 0; i < view.metadata().size(); ++i) {
assert(view[i].id() == i);

{
edm::LogInfo msg("TestAlpakaAnalyzer");
msg << source_.encode() << ".size() = " << view.metadata().size() << '\n';
msg << " data @ " << product.buffer().data() << ",\n"
<< " x @ " << view.metadata().addressOf_x() << " = " << Column(view.x(), view.metadata().size()) << ",\n"
<< " y @ " << view.metadata().addressOf_y() << " = " << Column(view.y(), view.metadata().size()) << ",\n"
<< " z @ " << view.metadata().addressOf_z() << " = " << Column(view.z(), view.metadata().size()) << ",\n"
<< " id @ " << view.metadata().addressOf_id() << " = " << Column(view.id(), view.metadata().size())
<< ",\n"
<< " r @ " << view.metadata().addressOf_r() << " = " << view.r() << '\n'
<< " m @ " << view.metadata().addressOf_m() << " = { ... {" << view[1].m()(1, Eigen::all)
<< " } ... } \n";
msg << std::hex << " [y - x] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_y()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_x())
<< " [z - y] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_z()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_y())
<< " [id - z] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_id()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_z())
<< " [r - id] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_r()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_id())
<< " [m - r] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_m()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_r());
}

edm::LogInfo msg("TestAlpakaAnalyzer");
msg << source_.encode() << ".size() = " << view.metadata().size() << '\n';
msg << " data = " << product.buffer().data() << ",\n"
<< " x = " << view.metadata().addressOf_x() << ",\n"
<< " y = " << view.metadata().addressOf_y() << ",\n"
<< " z = " << view.metadata().addressOf_z() << ",\n"
<< " id = " << view.metadata().addressOf_id() << ",\n"
<< " r = " << view.metadata().addressOf_r() << '\n';
msg << std::hex << " [y - x] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_y()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_x())
<< " [z - y] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_z()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_y())
<< " [id - z] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_id()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_z())
<< " [r - id] = 0x"
<< reinterpret_cast<intptr_t>(view.metadata().addressOf_r()) -
reinterpret_cast<intptr_t>(view.metadata().addressOf_id());
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) {
auto vi = view[i];
assert(vi.x() == 0.);
assert(vi.y() == 0.);
assert(vi.z() == 0.);
assert(vi.id() == i);
assert(vi.m() == matrix * i);
}
}

static void fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
Expand Down
4 changes: 3 additions & 1 deletion HeterogeneousCore/AlpakaTest/plugins/alpaka/TestAlgo.dev.cc
View file
Open in desktop
Expand Up @@ -20,11 +20,13 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
// this example accepts an arbitrary number of blocks and threads, and always uses 1 element per thread
const int32_t thread = alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u];
const int32_t stride = alpaka::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)[0u];
const portabletest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}};

if (thread == 0) {
view.r() = 1.;
}
for (auto i = thread; i < size; i += stride) {
view[i] = {0., 0., 0., i};
view[i] = {0., 0., 0., i, matrix * i};
}
}
};
Expand Down
6 changes: 3 additions & 3 deletions HeterogeneousCore/AlpakaTest/test/testHeterogeneousCoreAlpakaTestWriteRead.sh
View file
Open in desktop
Expand Up @@ -11,14 +11,14 @@ rm -f test.root
echo "--------------------------------------------------------------------------------"
echo "$ cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/writer.py"
echo
cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/writer.py
cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/writer.py || exit $?
echo
echo "--------------------------------------------------------------------------------"
echo "$ cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/reader.py"
echo
cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/reader.py
cmsRun ${LOCALTOP}/src/HeterogeneousCore/AlpakaTest/test/reader.py || exit $?
echo
echo "--------------------------------------------------------------------------------"
echo "$ edmDumpEventContent test.root"
echo
edmDumpEventContent test.root
edmDumpEventContent test.root || exit $?
4 changes: 3 additions & 1 deletion HeterogeneousCore/CUDATest/plugins/TestAlgo.cc
View file
Open in desktop
Expand Up @@ -5,10 +5,12 @@
namespace cudatest {

static void testAlgoKernel(cudatest::TestHostCollection::View view, int32_t size) {
const cudatest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}};

view.r() = 1.;

for (auto i = 0; i < size; ++i) {
view[i] = {0., 0., 0., i};
view[i] = {0., 0., 0., i, matrix * i};
}
}

Expand Down
3 changes: 2 additions & 1 deletion HeterogeneousCore/CUDATest/plugins/TestAlgo.cu
View file
Open in desktop
Expand Up @@ -9,12 +9,13 @@ namespace cudatest {
static __global__ void testAlgoKernel(cudatest::TestDeviceCollection::View view, int32_t size) {
const int32_t thread = blockIdx.x * blockDim.x + threadIdx.x;
const int32_t stride = blockDim.x * gridDim.x;
const cudatest::Matrix matrix{{1, 2, 3, 4, 5, 6}, {2, 4, 6, 8, 10, 12}, {3, 6, 9, 12, 15, 18}};

if (thread == 0) {
view.r() = 1.;
}
for (auto i = thread; i < size; i += stride) {
view[i] = {0., 0., 0., i};
view[i] = {0., 0., 0., i, matrix * i};
}
}

Expand Down