Skip to content

Commit

Permalink
Merge pull request #39532 from fwyzard/alpaka_framework
Browse files Browse the repository at this point in the history 
Read SoA scalars and Eigen columns from a ROOT streamer
  • Loading branch information
@cmsbuild
cmsbuild committed Oct 1, 2022
2 parents 2d2ca4d + 954edbc commit a89eb4c
Show file tree
Hide file tree
Showing 14 changed files with 186 additions and 53 deletions.
5 changes: 4 additions & 1 deletion CUDADataFormats/PortableTestObjects/interface/TestDeviceCollection.h
View file
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
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
@@ -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
@@ -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
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
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
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
99 changes: 76 additions & 23 deletions HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc
View file
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
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
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
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
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

0 comments on commit a89eb4c

Please sign in to comment.