Skip to content

Commit

Permalink
Merge pull request #42481 from ericcano/NoColumnSoA
Browse files Browse the repository at this point in the history
Fixes bug when SoA has not columns, only scalars [13.2.x]
  • Loading branch information
cmsbuild authored Aug 16, 2023
2 parents cfc01ed + 792466d commit d03bf03
Show file tree
Hide file tree
Showing 3 changed files with 21 additions and 4 deletions.
15 changes: 11 additions & 4 deletions DataFormats/SoATemplate/interface/SoALayout.h
Original file line number Diff line number Diff line change
Expand Up @@ -493,10 +493,17 @@
_ITERATE_ON_ALL(_DEFINE_METADATA_MEMBERS, ~, __VA_ARGS__) \
\
struct value_element { \
SOA_HOST_DEVICE SOA_INLINE value_element( \
_ITERATE_ON_ALL_COMMA(_VALUE_ELEMENT_CTOR_ARGS, ~, __VA_ARGS__) \
) : \
_ITERATE_ON_ALL_COMMA(_VALUE_ELEMENT_INITIALIZERS, ~, __VA_ARGS__) \
SOA_HOST_DEVICE SOA_INLINE value_element \
BOOST_PP_IF( \
BOOST_PP_SEQ_SIZE(_ITERATE_ON_ALL(_VALUE_ELEMENT_CTOR_ARGS, ~, __VA_ARGS__) ), \
(_ITERATE_ON_ALL_COMMA(_VALUE_ELEMENT_CTOR_ARGS, ~, __VA_ARGS__)):, \
()) \
BOOST_PP_TUPLE_ENUM(BOOST_PP_IF( \
BOOST_PP_SEQ_SIZE(_ITERATE_ON_ALL(_VALUE_ELEMENT_CTOR_ARGS, ~, __VA_ARGS__)), \
BOOST_PP_SEQ_TO_TUPLE(_ITERATE_ON_ALL(_VALUE_ELEMENT_INITIALIZERS, ~, __VA_ARGS__)), \
() \
) \
) \
{} \
\
_ITERATE_ON_ALL(_DEFINE_VALUE_ELEMENT_MEMBERS, ~, __VA_ARGS__) \
Expand Down
5 changes: 5 additions & 0 deletions DataFormats/SoATemplate/test/SoALayoutAndView_t.cu
Original file line number Diff line number Diff line change
Expand Up @@ -67,6 +67,11 @@ GENERATE_SOA_VIEW(SoAFullDeviceConstViewTemplate,
using SoAFullDeviceView =
SoAFullDeviceViewTemplate<cms::soa::CacheLineSize::NvidiaGPU, cms::soa::AlignmentEnforcement::enforced>;

// These SoAs validate that the generating macros do not get confused in the special case where there are
// no columns and only scalar elements in the SoA.
GENERATE_SOA_LAYOUT(TestSoALayoutNoColumn, SOA_SCALAR(double, r))
GENERATE_SOA_LAYOUT(TestSoALayoutNoColumn2, SOA_SCALAR(double, r), SOA_SCALAR(double, r2))

// Eigen cross product kernel (on store)
__global__ void crossProduct(SoAHostDeviceView soa, const unsigned int numElements) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
Expand Down
5 changes: 5 additions & 0 deletions DataFormats/SoATemplate/test/SoALayoutAndView_t.hip.cc
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,11 @@ GENERATE_SOA_VIEW(SoAFullDeviceConstViewTemplate,
using SoAFullDeviceView =
SoAFullDeviceViewTemplate<cms::soa::CacheLineSize::NvidiaGPU, cms::soa::AlignmentEnforcement::enforced>;

// These SoAs validate that the generating macros do not get confused in the special case where there are
// no columns and only scalar elements in the SoA.
GENERATE_SOA_LAYOUT(TestSoALayoutNoColumn, SOA_SCALAR(double, r))
GENERATE_SOA_LAYOUT(TestSoALayoutNoColumn2, SOA_SCALAR(double, r), SOA_SCALAR(double, r2))

// Eigen cross product kernel (on store)
__global__ void crossProduct(SoAHostDeviceView soa, const int numElements) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
Expand Down

0 comments on commit d03bf03

Please sign in to comment.