diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index 89a70369fa08f..847da45226c6f 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -49,9 +49,9 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection(4 * nHits(), stream); size_t rowSize = sizeof(float) * nHits(); - size_t srcPitch = ptrdiff_t(view().yLocal()) - ptrdiff_t(view().xLocal()); - cudaCheck( - cudaMemcpy2DAsync(ret.get(), rowSize, view().xLocal(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream)); + size_t srcPitch = ptrdiff_t(view().yLocal().data()) - ptrdiff_t(view().xLocal().data()); + cudaCheck(cudaMemcpy2DAsync( + ret.get(), rowSize, view().xLocal().data(), srcPitch, rowSize, 4, cudaMemcpyDeviceToHost, stream)); return ret; } //move to utilities diff --git a/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu b/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu index 48e8dea96911e..7609e2ca12eb1 100644 --- a/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu +++ b/CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu @@ -45,11 +45,11 @@ namespace testTrackingRecHitSoA { cudaCheck(cudaDeviceSynchronize()); cms::cuda::fillManyFromVector(&(hits.view().phiBinner()), 10, - hits.view().iphi(), + hits.view().iphi().data(), hits.view().hitsLayerStart().data(), 2000, 256, - hits.view().phiBinnerStorage(), + hits.view().phiBinnerStorage().data(), stream); cudaCheck(cudaDeviceSynchronize()); show<<<10, 1000, 0, stream>>>(hits.view()); diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc index f3635d6df45da..3dc60de8416c6 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc @@ -157,8 +157,8 @@ void SiPixelCompareTrackSoA::analyze(const edm::Event& iEvent, const edm::Eve auto const& tsoaGPU = *tsoaHandleGPU; auto maxTracksCPU = tsoaCPU.view().metadata().size(); //this should be same for both? auto maxTracksGPU = tsoaGPU.view().metadata().size(); //this should be same for both? - auto const* qualityCPU = tsoaCPU.view().quality(); - auto const* qualityGPU = tsoaGPU.view().quality(); + auto const qualityCPU = tsoaCPU.view().quality(); + auto const qualityGPU = tsoaGPU.view().quality(); int32_t nTracksCPU = 0; int32_t nTracksGPU = 0; int32_t nLooseAndAboveTracksCPU = 0; diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTracks.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTracks.cc index 9e279f249d810..5f5480421a9ad 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTracks.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTracks.cc @@ -166,8 +166,8 @@ void SiPixelCompareTracks::analyzeSeparate(U tokenRef, V tokenTar, const edm::Ev auto maxTracksRef = tsoaRef.view().metadata().size(); //this should be same for both? auto maxTracksTar = tsoaTar.view().metadata().size(); //this should be same for both? - auto const* qualityRef = tsoaRef.view().quality(); - auto const* qualityTar = tsoaTar.view().quality(); + auto const qualityRef = tsoaRef.view().quality(); + auto const qualityTar = tsoaTar.view().quality(); int32_t nTracksRef = 0; int32_t nTracksTar = 0; diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc index f3ccb74bc3fea..59565f4932ecc 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc @@ -86,7 +86,7 @@ void SiPixelMonitorTrackSoA::analyze(const edm::Event& iEvent, const edm::Eve using helper = TracksUtilities; auto const& tsoa = *tsoaHandle.product(); auto maxTracks = tsoa.view().metadata().size(); - auto const* quality = tsoa.view().quality(); + auto const quality = tsoa.view().quality(); int32_t nTracks = 0; int32_t nLooseAndAboveTracks = 0; diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoAAlpaka.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoAAlpaka.cc index 9b519116e149d..877ecb483691e 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoAAlpaka.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoAAlpaka.cc @@ -84,7 +84,7 @@ void SiPixelMonitorTrackSoAAlpaka::analyze(const edm::Event& iEvent, const edm:: auto const& tsoa = *tsoaHandle.product(); auto maxTracks = tsoa.view().metadata().size(); - auto const* quality = tsoa.view().quality(); + auto const quality = tsoa.view().quality(); int32_t nTracks = 0; int32_t nLooseAndAboveTracks = 0; diff --git a/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h b/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h index a97dfadea52c4..c3c02be163745 100644 --- a/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h +++ b/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h @@ -3,6 +3,7 @@ #include #include +#include #include namespace legacy { @@ -17,12 +18,15 @@ namespace legacy { class SiPixelDigisSoA { public: SiPixelDigisSoA() = default; - explicit SiPixelDigisSoA( - size_t nDigis, const uint32_t* pdigi, const uint32_t* rawIdArr, const uint16_t* adc, const int32_t* clus) - : pdigi_(pdigi, pdigi + nDigis), - rawIdArr_(rawIdArr, rawIdArr + nDigis), - adc_(adc, adc + nDigis), - clus_(clus, clus + nDigis) {} + explicit SiPixelDigisSoA(size_t nDigis, + std::span pdigi, + std::span rawIdArr, + std::span adc, + std::span clus) + : pdigi_(pdigi.data(), pdigi.data() + nDigis), + rawIdArr_(rawIdArr.data(), rawIdArr.data() + nDigis), + adc_(adc.data(), adc.data() + nDigis), + clus_(clus.data(), clus.data() + nDigis) {} ~SiPixelDigisSoA() = default; diff --git a/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsDevice.h b/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsDevice.h index fa87863dc7b16..0c04544bd1f23 100644 --- a/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsDevice.h +++ b/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsDevice.h @@ -26,7 +26,6 @@ class SiPixelDigiErrorsDevice : public PortableDeviceCollection(maxFedWords, device) {} - auto& error_data() const { return (*this->view().pixelErrors()); } auto maxFedWords() const { return maxFedWords_; } private: diff --git a/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h b/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h index c8974156d29a3..1150a82046428 100644 --- a/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h +++ b/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h @@ -22,9 +22,6 @@ class SiPixelDigiErrorsHost : public PortableHostCollection struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params) - : params_(params) {} - SOA_HOST_DEVICE SOA_INLINE T* operator()() { return params_.addr_; } + SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& params, + size_type size) + : params_(params), size_(size) {} + SOA_HOST_DEVICE SOA_INLINE std::span operator()() { return std::span(params_.addr_, size_); } - using NoParamReturnType = T*; + using NoParamReturnType = std::span; using ParamReturnType = T&; SOA_HOST_DEVICE SOA_INLINE T& operator()(size_type index) { return params_.addr_[index]; } private: SoAParametersImpl params_; + size_type size_ = 0; }; // Const column template struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) - : params_(params) {} - SOA_HOST_DEVICE SOA_INLINE const T* operator()() const { return params_.addr_; } - using NoParamReturnType = const T*; + SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params, + size_type size) + : params_(params), size_(size) {} + SOA_HOST_DEVICE SOA_INLINE std::span operator()() const { + return std::span(params_.addr_, size_); + } + using NoParamReturnType = std::span; using ParamReturnType = const T&; SOA_HOST_DEVICE SOA_INLINE T const& operator()(size_type index) const { return params_.addr_[index]; } private: SoAConstParametersImpl params_; + const size_type size_ = 0; }; // Scalar @@ -738,10 +744,13 @@ namespace cms::soa { // Eigen-type template struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAParametersImpl& 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 SoAColumnAccessorsImpl(const SoAParametersImpl& params, + size_type size) + : params_(params), size_(size) {} + SOA_HOST_DEVICE SOA_INLINE std::span operator()() { + return std::span(params_.addr_, size_); + } + using NoParamReturnType = std::span; using ParamReturnType = typename SoAValue::MapType; SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) { return SoAValue(index, params_)(); @@ -749,15 +758,19 @@ namespace cms::soa { private: SoAParametersImpl params_; + size_type size_ = 0; }; // Const Eigen-type template struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params) - : params_(params) {} - SOA_HOST_DEVICE SOA_INLINE typename T::Scalar const* operator()() const { return params_.addr_; } - using NoParamReturnType = typename T::Scalar const*; + SOA_HOST_DEVICE SOA_INLINE SoAColumnAccessorsImpl(const SoAConstParametersImpl& params, + size_type size) + : params_(params), size_(size) {} + SOA_HOST_DEVICE SOA_INLINE std::span operator()() const { + return std::span(params_.addr_, size_); + } + using NoParamReturnType = std::span; using ParamReturnType = typename SoAValue::CMapType; SOA_HOST_DEVICE SOA_INLINE ParamReturnType operator()(size_type index) const { return SoAConstValue(index, params_)(); @@ -765,6 +778,7 @@ namespace cms::soa { private: SoAConstParametersImpl params_; + const size_type size_ = 0; }; /* A helper template stager to avoid commas inside macros */ diff --git a/DataFormats/SoATemplate/interface/SoALayout.h b/DataFormats/SoATemplate/interface/SoALayout.h index 01fda57d53099..5a7eb1aa2a75c 100644 --- a/DataFormats/SoATemplate/interface/SoALayout.h +++ b/DataFormats/SoATemplate/interface/SoALayout.h @@ -792,36 +792,103 @@ _SWITCH_ON_TYPE(VALUE_TYPE, * Direct access to column pointer (const) and indexed access. */ // clang-format off -#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(VALUE_TYPE, CPP_TYPE, NAME, ARGS) \ - /* Column or scalar */ \ - SOA_HOST_DEVICE SOA_INLINE \ - typename cms::soa::SoAAccessors:: \ - template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>::template Alignment:: \ - template RestrictQualifier::NoParamReturnType \ - NAME() const { \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>::template Alignment:: \ - template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_))(); \ - } \ - SOA_HOST_DEVICE SOA_INLINE \ - typename cms::soa::SoAAccessors:: \ - template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>::template Alignment:: \ - template RestrictQualifier::ParamReturnType \ - NAME(size_type _soa_impl_index) const { \ - if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ - if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ - SOA_THROW_OUT_OF_RANGE("Out of range index in const " #NAME "(size_type index)", \ - _soa_impl_index, elements_) \ - } \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType::template AccessType< \ - cms::soa::SoAAccessType::constAccess>::template Alignment:: \ - template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_))(_soa_impl_index); \ - } \ - // clang-format on +#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(VALUE_TYPE, CPP_TYPE, NAME, ARGS) \ +_SWITCH_ON_TYPE(VALUE_TYPE, \ + /* Scalar */ \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ + NAME() const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_))(); \ + } \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + NAME(size_type _soa_impl_index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #NAME "(size_type index)", \ + _soa_impl_index, elements_) \ + } \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_))(_soa_impl_index); \ + } \ + , \ + /* Column */ \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ + NAME() const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_), elements_)(); \ + } \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + NAME(size_type _soa_impl_index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #NAME "(size_type index)", \ + _soa_impl_index, elements_) \ + } \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_), \ + elements_)(_soa_impl_index); \ + } \ + , \ + /* Eigen column */ \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ + NAME() const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_), \ + cms::soa::alignSize(elements_ * sizeof(CPP_TYPE::Scalar), alignment) / \ + sizeof(CPP_TYPE::Scalar) * CPP_TYPE::RowsAtCompileTime * \ + CPP_TYPE::ColsAtCompileTime)(); \ + } \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + NAME(size_type _soa_impl_index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (_soa_impl_index >= elements_ or _soa_impl_index < 0) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #NAME "(size_type index)", \ + _soa_impl_index, elements_) \ + } \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>::template Alignment:: \ + template RestrictQualifier(BOOST_PP_CAT(NAME, Parameters_), \ + cms::soa::alignSize(elements_ * sizeof(CPP_TYPE::Scalar), alignment) / \ + sizeof(CPP_TYPE::Scalar) * CPP_TYPE::RowsAtCompileTime * \ + CPP_TYPE::ColsAtCompileTime)(_soa_impl_index); \ + } \ +) +// clang-format on #define _DECLARE_VIEW_SOA_CONST_ACCESSOR(R, DATA, TYPE_NAME) \ BOOST_PP_IF(BOOST_PP_GREATER(BOOST_PP_TUPLE_ELEM(0, TYPE_NAME), _VALUE_LAST_COLUMN_TYPE), \ @@ -1055,7 +1122,8 @@ _SWITCH_ON_TYPE(VALUE_TYPE, */ // clang-format off #define _DECLARE_VIEW_SOA_ACCESSOR_IMPL(VALUE_TYPE, CPP_TYPE, NAME, ARGS) \ - /* Column or scalar */ \ + _SWITCH_ON_TYPE(VALUE_TYPE, \ + /* Scalar */ \ SOA_HOST_DEVICE SOA_INLINE \ typename cms::soa::SoAAccessors:: \ template ColumnType::template AccessType< \ @@ -1084,7 +1152,76 @@ _SWITCH_ON_TYPE(VALUE_TYPE, cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ template RestrictQualifier(cms::soa::const_cast_SoAParametersImpl( \ base_type:: BOOST_PP_CAT(NAME, Parameters_)))(_soa_impl_index); \ - } + } \ + , \ + /* Column */ \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ + NAME() { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(cms::soa::const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(NAME, Parameters_)), base_type::elements_)(); \ + } \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + NAME(size_type _soa_impl_index) { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (_soa_impl_index >= base_type::elements_ or _soa_impl_index < 0) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #NAME "(size_type index)", \ + _soa_impl_index, base_type::elements_) \ + } \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(cms::soa::const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(NAME, Parameters_)), base_type::elements_)(_soa_impl_index); \ + } \ + , \ + /* Eigen column */ \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::NoParamReturnType \ + NAME() { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(cms::soa::const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(NAME, Parameters_)), \ + cms::soa::alignSize(base_type::elements_ * sizeof(CPP_TYPE::Scalar), alignment) / \ + sizeof(CPP_TYPE::Scalar) * CPP_TYPE::RowsAtCompileTime * \ + CPP_TYPE::ColsAtCompileTime)(); \ + } \ + SOA_HOST_DEVICE SOA_INLINE \ + typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier::ParamReturnType \ + NAME(size_type _soa_impl_index) { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::enabled) { \ + if (_soa_impl_index >= base_type::elements_ or _soa_impl_index < 0) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in mutable " #NAME "(size_type index)", \ + _soa_impl_index, base_type::elements_) \ + } \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>::template Alignment:: \ + template RestrictQualifier(cms::soa::const_cast_SoAParametersImpl( \ + base_type:: BOOST_PP_CAT(NAME, Parameters_)), \ + cms::soa::alignSize(base_type::elements_ * sizeof(CPP_TYPE::Scalar), alignment) / \ + sizeof(CPP_TYPE::Scalar) * CPP_TYPE::RowsAtCompileTime * \ + CPP_TYPE::ColsAtCompileTime)(_soa_impl_index); \ + } \ +) // clang-format on #define _DECLARE_VIEW_SOA_ACCESSOR(R, DATA, TYPE_NAME) \ diff --git a/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.cc b/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.cc index eaf499345388d..dc3508979e4dc 100644 --- a/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.cc +++ b/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.cc @@ -52,17 +52,18 @@ int main() { hitsX[i] = float(i) * 2; } - auto moduleStartD = cms::alpakatools::make_device_view(queue, clusters.view().clusModuleStart(), nHits); + auto moduleStartD = + cms::alpakatools::make_device_view(queue, clusters.view().clusModuleStart().data(), nHits); alpaka::memcpy(queue, moduleStartD, moduleStartH); TrackingRecHitsSoACollection tkhit(queue, clusters); // exercise the copy of a full column (on device) - auto hitXD = cms::alpakatools::make_device_view(queue, tkhit.view().xLocal(), nHits); + auto hitXD = cms::alpakatools::make_device_view(queue, tkhit.view().xLocal().data(), nHits); alpaka::memcpy(queue, hitXD, hitsX); // exercise the memset of a colum (on device) - auto hitYD = cms::alpakatools::make_device_view(queue, tkhit.view().yGlobal(), nHits); + auto hitYD = cms::alpakatools::make_device_view(queue, tkhit.view().yGlobal().data(), nHits); constexpr float constYG = -14.0458; std::vector constYV(nHits, constYG); auto constYGV_v = cms::alpakatools::make_host_view(constYV.data(), nHits); @@ -86,7 +87,7 @@ int main() { ::reco::TrackingRecHitHost host_collection_2(cms::alpakatools::host(), nHits, nModules); // exercise the memset of a colum (on host) - auto hitLYH = cms::alpakatools::make_host_view(host_collection_2.view().yLocal(), nHits); + auto hitLYH = cms::alpakatools::make_host_view(host_collection_2.view().yLocal().data(), nHits); constexpr float constYL = -27.0855; std::vector constYLV(nHits, constYL); auto constYL_v = cms::alpakatools::make_host_view(constYLV.data(), nHits); diff --git a/EventFilter/EcalRawToDigi/plugins/EcalDigisFromPortableProducer.cc b/EventFilter/EcalRawToDigi/plugins/EcalDigisFromPortableProducer.cc index e3918c465c4b4..fba76df5490b4 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalDigisFromPortableProducer.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalDigisFromPortableProducer.cc @@ -151,12 +151,15 @@ void EcalDigisFromPortableProducer::produce(edm::Event& event, edm::EventSetup c auto const digisEEDataSize = digisEESize * ecalPh1::sampleSize; // Intermediate containers because the DigiCollection containers are accessible only as const - EBDigiCollection::IdContainer digisIdsEB(digisEBSoAView.id(), digisEBSoAView.id() + digisEBSize); - EEDigiCollection::IdContainer digisIdsEE(digisEESoAView.id(), digisEESoAView.id() + digisEESize); - EBDigiCollection::DataContainer digisDataEB(digisEBSoAView.data()->data(), - digisEBSoAView.data()->data() + digisEBDataSize); - EEDigiCollection::DataContainer digisDataEE(digisEESoAView.data()->data(), - digisEESoAView.data()->data() + digisEEDataSize); + EBDigiCollection::IdContainer digisIdsEB(digisEBSoAView.id().data(), digisEBSoAView.id().data() + digisEBSize); + EEDigiCollection::IdContainer digisIdsEE(digisEESoAView.id().data(), digisEESoAView.id().data() + digisEESize); + // digisEBSoAView.data() returns a span, where EcalDataArray is an array of uint16_t + // digisEBSoAView.data().data() returns a pointer to the first EcalDataArray of the data column + // digisEBSoAView.data().data()->data() returns a pointer to the first uint16_t of the first EcalDataArray of the data column + EBDigiCollection::DataContainer digisDataEB(digisEBSoAView.data().data()->data(), + digisEBSoAView.data().data()->data() + digisEBDataSize); + EEDigiCollection::DataContainer digisDataEE(digisEESoAView.data().data()->data(), + digisEESoAView.data().data()->data() + digisEEDataSize); digisEB->swap(digisIdsEB, digisDataEB); digisEE->swap(digisIdsEE, digisDataEE); diff --git a/EventFilter/EcalRawToDigi/plugins/alpaka/UnpackPortable.dev.cc b/EventFilter/EcalRawToDigi/plugins/alpaka/UnpackPortable.dev.cc index c85dea4cbb915..cd51cc9ffd804 100644 --- a/EventFilter/EcalRawToDigi/plugins/alpaka/UnpackPortable.dev.cc +++ b/EventFilter/EcalRawToDigi/plugins/alpaka/UnpackPortable.dev.cc @@ -37,8 +37,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw { // size auto const gridDim = alpaka::getWorkDiv(acc)[0u]; auto const size = ifed == gridDim - 1 ? nbytesTotal - offset : offsets[ifed + 1] - offset; - auto* samples = isBarrel ? digisDevEB.data()->data() : digisDevEE.data()->data(); - auto* ids = isBarrel ? digisDevEB.id() : digisDevEE.id(); + // digisDevEB.data() returns a span, where EcalDataArray is an array of uint16_t + // digisDevEB.data().data() returns a pointer to the first EcalDataArray of the data column + // digisDevEB.data().data()->data() returns a pointer to the first uint16_t of the first EcalDataArray of the data column + auto* samples = isBarrel ? digisDevEB.data().data()->data() : digisDevEE.data().data()->data(); + auto* ids = isBarrel ? digisDevEB.id().data() : digisDevEE.id().data(); auto* pChannelsCounter = isBarrel ? &digisDevEB.size() : &digisDevEE.size(); // offset to the right raw buffer diff --git a/HeterogeneousCore/AlpakaInterface/README.md b/HeterogeneousCore/AlpakaInterface/README.md index 21dfd2cf8a319..0fd1d3f219f61 100644 --- a/HeterogeneousCore/AlpakaInterface/README.md +++ b/HeterogeneousCore/AlpakaInterface/README.md @@ -361,6 +361,11 @@ These helper functions instantiate zero-dimensional (scalars) and one-dimensiona instantiates a one-dimensional view over an array starting at `span.data()` with `span.size()` elements in host memory; + - `auto make_host_view(std::span span, size)` + instantiates a one-dimensional view over an array starting at `span.data()` with + `size` elements in host memory. If `size` is greater than `span.size()`, an exception + is launched. + ## Device memory views @@ -395,11 +400,16 @@ These helper functions instantiate zero-dimensional (scalars) and one-dimensiona - `auto make_device_view(device, T* data, size)` instantiates a one-dimensional view over an array starting at `data` with - `size` elements in device global memory. + `size` elements in device global memory; - `auto make_device_view(device, std::span span)` instantiates a one-dimensional view over an array starting at `span.data()` with - `span.size()` elements in device global memory. + `span.size()` elements in device global memory; + + - `auto make_device_view(device, std::span span, size)` + instantiates a one-dimensional view over an array starting at `span.data()` with + `size` elements in host memory; if `size` is greater than `span.size()`, an exception + is launched. The `make_device_view` functions can also accept as a first argument a `queue` instead of a `device`: @@ -414,11 +424,16 @@ instead of a `device`: - `auto make_device_view(queue, T* data, size)` instantiates a one-dimensional view over an array starting at `data` with - `size` elements in device global memory. + `size` elements in device global memory; - `auto make_device_view(queue, std::span span)` instantiates a one-dimensional view over an array starting at `span.data()` with - `span.size()` elements in device global memory. + `span.size()` elements in device global memory; + + - `auto make_device_view(queue, std::span span, size)` + instantiates a one-dimensional view over an array starting at `span.data()` with + `size` elements in host memory; if `size` is greater than `span.size()`, an exception + is launched. These functions use the device associated to the given `queue`. These operations are otherwise identical to those that take a `device` as their first argument. diff --git a/HeterogeneousCore/AlpakaInterface/interface/memory.h b/HeterogeneousCore/AlpakaInterface/interface/memory.h index 66db519e2c4ec..3a2d5caf1281f 100644 --- a/HeterogeneousCore/AlpakaInterface/interface/memory.h +++ b/HeterogeneousCore/AlpakaInterface/interface/memory.h @@ -164,6 +164,14 @@ namespace cms::alpakatools { return alpaka::ViewPlainPtr(span.data(), host(), Vec1D{span.size()}); } + template + host_view make_host_view(std::span span, Extent extent) { + if (extent > span.size()) { + throw std::runtime_error("make_host_view: span size is smaller than the specified extent"); + } + return alpaka::ViewPlainPtr(span.data(), host(), Vec1D{extent}); + } + template std::enable_if_t and not std::is_array_v>, host_view> make_host_view(T& data, Extent extent) { @@ -280,6 +288,16 @@ namespace cms::alpakatools { return alpaka::ViewPlainPtr(span.data(), device, Vec1D{span.size()}); } + template + std::enable_if_t, device_view> make_device_view(TDev const& device, + std::span span, + Extent extent) { + if (extent > span.size()) { + throw std::runtime_error("make_device_view: span size is smaller than the specified extent"); + } + return alpaka::ViewPlainPtr(span.data(), device, Vec1D{extent}); + } + template std::enable_if_t and cms::is_unbounded_array_v and not std::is_array_v>, @@ -316,6 +334,16 @@ namespace cms::alpakatools { span.data(), alpaka::getDev(queue), Vec1D{span.size()}); } + template + std::enable_if_t, device_view, T[]>> make_device_view(TQueue const& queue, + std::span span, + Extent extent) { + if (extent > span.size()) { + throw std::runtime_error("make_device_view: span size is smaller than the specified extent"); + } + return alpaka::ViewPlainPtr, T, Dim1D, Idx>(span.data(), alpaka::getDev(queue), Vec1D{extent}); + } + template std::enable_if_t and cms::is_unbounded_array_v and not std::is_array_v>, diff --git a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc index cbeae5e4fef81..2a5437bf30ce6 100644 --- a/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc +++ b/HeterogeneousCore/AlpakaTest/plugins/TestAlpakaAnalyzer.cc @@ -19,6 +19,8 @@ namespace { public: Column(T const* data, size_t size) : data_(data), size_(size) {} + Column(std::span span) : data_(span.data()), size_(span.size()) {} + void print(std::ostream& out) const { std::stringstream buffer; buffer << "{ "; @@ -52,16 +54,16 @@ namespace { template void checkViewAddresses(T const& view) { // columns - assert(view.metadata().addressOf_x() == view.x()); + assert(view.metadata().addressOf_x() == view.x().data()); 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().data()); 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().data()); 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().data()); assert(view.metadata().addressOf_id() == &view.id(0)); assert(view.metadata().addressOf_id() == &view[0].id()); // scalars @@ -69,30 +71,30 @@ namespace { //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 // columns of arrays - assert(view.metadata().addressOf_flags() == view.flags()); + assert(view.metadata().addressOf_flags() == view.flags().data()); assert(view.metadata().addressOf_flags() == &view.flags(0)); assert(view.metadata().addressOf_flags() == &view[0].flags()); // columns of Eigen matrices - assert(view.metadata().addressOf_m() == view.m()); + assert(view.metadata().addressOf_m() == view.m().data()); assert(view.metadata().addressOf_m() == &view.m(0).coeffRef(0, 0)); assert(view.metadata().addressOf_m() == &view[0].m().coeffRef(0, 0)); } template void checkViewAddresses2(T const& view) { - assert(view.metadata().addressOf_x2() == view.x2()); + assert(view.metadata().addressOf_x2() == view.x2().data()); assert(view.metadata().addressOf_x2() == &view.x2(0)); assert(view.metadata().addressOf_x2() == &view[0].x2()); - assert(view.metadata().addressOf_y2() == view.y2()); + assert(view.metadata().addressOf_y2() == view.y2().data()); assert(view.metadata().addressOf_y2() == &view.y2(0)); assert(view.metadata().addressOf_y2() == &view[0].y2()); - assert(view.metadata().addressOf_z2() == view.z2()); + assert(view.metadata().addressOf_z2() == view.z2().data()); assert(view.metadata().addressOf_z2() == &view.z2(0)); assert(view.metadata().addressOf_z2() == &view[0].z2()); - assert(view.metadata().addressOf_id2() == view.id2()); + assert(view.metadata().addressOf_id2() == view.id2().data()); assert(view.metadata().addressOf_id2() == &view.id2(0)); assert(view.metadata().addressOf_id2() == &view[0].id2()); - assert(view.metadata().addressOf_m2() == view.m2()); + assert(view.metadata().addressOf_m2() == view.m2().data()); assert(view.metadata().addressOf_m2() == &view.m2(0).coeffRef(0, 0)); assert(view.metadata().addressOf_m2() == &view[0].m2().coeffRef(0, 0)); assert(view.metadata().addressOf_r2() == &view.r2()); @@ -102,19 +104,19 @@ namespace { template void checkViewAddresses3(T const& view) { - assert(view.metadata().addressOf_x3() == view.x3()); + assert(view.metadata().addressOf_x3() == view.x3().data()); assert(view.metadata().addressOf_x3() == &view.x3(0)); assert(view.metadata().addressOf_x3() == &view[0].x3()); - assert(view.metadata().addressOf_y3() == view.y3()); + assert(view.metadata().addressOf_y3() == view.y3().data()); assert(view.metadata().addressOf_y3() == &view.y3(0)); assert(view.metadata().addressOf_y3() == &view[0].y3()); - assert(view.metadata().addressOf_z3() == view.z3()); + assert(view.metadata().addressOf_z3() == view.z3().data()); assert(view.metadata().addressOf_z3() == &view.z3(0)); assert(view.metadata().addressOf_z3() == &view[0].z3()); - assert(view.metadata().addressOf_id3() == view.id3()); + assert(view.metadata().addressOf_id3() == view.id3().data()); assert(view.metadata().addressOf_id3() == &view.id3(0)); assert(view.metadata().addressOf_id3() == &view[0].id3()); - assert(view.metadata().addressOf_m3() == view.m3()); + assert(view.metadata().addressOf_m3() == view.m3().data()); assert(view.metadata().addressOf_m3() == &view.m3(0).coeffRef(0, 0)); assert(view.metadata().addressOf_m3() == &view[0].m3().coeffRef(0, 0)); assert(view.metadata().addressOf_r3() == &view.r3()); @@ -155,14 +157,12 @@ class TestAlpakaAnalyzer : public edm::global::EDAnalyzer<> { 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" + << " x @ " << view.metadata().addressOf_x() << " = " << Column(view.x()) << ",\n" + << " y @ " << view.metadata().addressOf_y() << " = " << Column(view.y()) << ",\n" + << " z @ " << view.metadata().addressOf_z() << " = " << Column(view.z()) << ",\n" + << " id @ " << view.metadata().addressOf_id() << " = " << Column(view.id()) << ",\n" << " r @ " << view.metadata().addressOf_r() << " = " << view.r() << '\n' - << " flags @ " << view.metadata().addressOf_flags() << " = " << Column(view.flags(), view.metadata().size()) - << ",\n" + << " flags @ " << view.metadata().addressOf_flags() << " = " << Column(view.flags()) << ",\n" << " m @ " << view.metadata().addressOf_m() << " = { ... {" << view[1].m()(1, Eigen::indexing::all) << " } ... } \n"; msg << std::hex << " [y - x] = 0x" diff --git a/HeterogeneousCore/CUDATest/plugins/TestPortableAnalyzer.cc b/HeterogeneousCore/CUDATest/plugins/TestPortableAnalyzer.cc index 68d6774ea915e..1d3bd585ccdfe 100644 --- a/HeterogeneousCore/CUDATest/plugins/TestPortableAnalyzer.cc +++ b/HeterogeneousCore/CUDATest/plugins/TestPortableAnalyzer.cc @@ -19,6 +19,8 @@ namespace { public: Column(T const* data, size_t size) : data_(data), size_(size) {} + Column(std::span span) : data_(span.data()), size_(span.size()) {} + void print(std::ostream& out) const { std::stringstream buffer; buffer << "{ "; @@ -66,13 +68,12 @@ class TestPortableAnalyzer : public edm::stream::EDAnalyzer<> { edm::LogInfo msg("TestPortableAnalyzer"); msg << source_.encode() << ".size() = " << view.metadata().size() << '\n'; msg << " data @ " << product.buffer().get() << ",\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" + << " x @ " << view.metadata().addressOf_x() << " = " << Column(view.x()) << ",\n" + << " y @ " << view.metadata().addressOf_y() << " = " << Column(view.y()) << ",\n" + << " z @ " << view.metadata().addressOf_z() << " = " << Column(view.z()) << ",\n" + << " id @ " << view.metadata().addressOf_id() << " = " << Column(view.id()) << ",\n" << " r @ " << view.metadata().addressOf_r() << " = " << view.r() << '\n' - << " flags @ " << view.metadata().addressOf_flags() << " = " << Column(view.flags(), view.metadata().size()) - << ",\n" + << " flags @ " << view.metadata().addressOf_flags() << " = " << Column(view.flags()) << ",\n" << " m @ " << view.metadata().addressOf_m() << " = { ... {" << view[1].m()(1, Eigen::indexing::all) << " } ... } \n"; msg << std::hex << " [y - x] = 0x" diff --git a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationCommonKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationCommonKernels.h index f71136b4f5989..ed94911a9ed5f 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationCommonKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationCommonKernels.h @@ -75,8 +75,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { auto const ch = idx.global / nsamples; // for accessing input arrays int const inputTx = ch >= nchannelsEB ? idx.global - nchannelsEB * nsamples : idx.global; + // digisDevEE.data() returns a span, where EcalDataArray is an array of uint16_t + // digisDevEE.data().data() returns a pointer to the first EcalDataArray of the data column + // digisDevEE.data().data()->data() returns a pointer to the first uint16_t of the first EcalDataArray of the data column // eb is first and then ee - auto const* digis_in = ch >= nchannelsEB ? digisDevEE.data()->data() : digisDevEB.data()->data(); + auto const* digis_in = + ch >= nchannelsEB ? digisDevEE.data().data()->data() : digisDevEB.data().data()->data(); auto const gainId = ecalMGPA::gainId(digis_in[inputTx]); // store into shared mem for initialization @@ -170,7 +174,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { int const inputCh = ch >= nchannelsEB ? ch - nchannelsEB : ch; int const inputTx = ch >= nchannelsEB ? idx.global - nchannelsEB * nsamples : idx.global; - auto const* dids = ch >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); + auto const dids = ch >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); auto const did = DetId{dids[inputCh]}; auto const isBarrel = did.subdetId() == EcalBarrel; // TODO offset for ee, 0 for eb @@ -178,16 +182,17 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { : offsetForHashes + reconstruction::hashedIndexEE(did.rawId()); // eb is first and then ee - auto const* digis_in = ch >= nchannelsEB ? digisDevEE.data()->data() : digisDevEB.data()->data(); + auto const* digis_in = + ch >= nchannelsEB ? digisDevEE.data().data()->data() : digisDevEB.data().data()->data(); auto* amplitudesForMinimization = reinterpret_cast<::ecal::multifit::SampleVector*>( - ch >= nchannelsEB ? uncalibRecHitsEE.outOfTimeAmplitudes()->data() - : uncalibRecHitsEB.outOfTimeAmplitudes()->data()); - auto* energies = ch >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); - auto* chi2 = ch >= nchannelsEB ? uncalibRecHitsEE.chi2() : uncalibRecHitsEB.chi2(); - auto* g_pedestal = ch >= nchannelsEB ? uncalibRecHitsEE.pedestal() : uncalibRecHitsEB.pedestal(); - auto* dids_out = ch >= nchannelsEB ? uncalibRecHitsEE.id() : uncalibRecHitsEB.id(); - auto* flags = ch >= nchannelsEB ? uncalibRecHitsEE.flags() : uncalibRecHitsEB.flags(); + ch >= nchannelsEB ? uncalibRecHitsEE.outOfTimeAmplitudes().data()->data() + : uncalibRecHitsEB.outOfTimeAmplitudes().data()->data()); + auto energies = ch >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); + auto chi2 = ch >= nchannelsEB ? uncalibRecHitsEE.chi2() : uncalibRecHitsEB.chi2(); + auto g_pedestal = ch >= nchannelsEB ? uncalibRecHitsEE.pedestal() : uncalibRecHitsEB.pedestal(); + auto dids_out = ch >= nchannelsEB ? uncalibRecHitsEE.id() : uncalibRecHitsEB.id(); + auto flags = ch >= nchannelsEB ? uncalibRecHitsEE.flags() : uncalibRecHitsEB.flags(); auto const adc = ecalMGPA::adc(digis_in[inputTx]); auto const gainId = ecalMGPA::gainId(digis_in[inputTx]); @@ -345,7 +350,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { constexpr bool simplifiedNoiseModelForGainSwitch = true; //---- default is true // pulse matrix - auto const* pulse_shapes = reinterpret_cast(conditionsDev.pulseShapes()->data()); + auto const* pulse_shapes = reinterpret_cast(conditionsDev.pulseShapes().data()->data()); auto const blockDimX = alpaka::getWorkDiv(acc)[1u]; auto const elemsPerBlockX = alpaka::getWorkDiv(acc)[1u]; @@ -359,17 +364,17 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { // to access input arrays (ids and digis only) int const inputCh = ch >= nchannelsEB ? ch - nchannelsEB : ch; - auto const* dids = ch >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); + auto const dids = ch >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); auto const did = DetId{dids[inputCh]}; auto const isBarrel = did.subdetId() == EcalBarrel; auto const hashedId = isBarrel ? ecal::reconstruction::hashedIndexEB(did.rawId()) : offsetForHashes + ecal::reconstruction::hashedIndexEE(did.rawId()); - auto const* G12SamplesCorrelation = isBarrel ? conditionsDev.sampleCorrelation_EB_G12().data() - : conditionsDev.sampleCorrelation_EE_G12().data(); - auto const* G6SamplesCorrelation = + auto const G12SamplesCorrelation = isBarrel ? conditionsDev.sampleCorrelation_EB_G12().data() + : conditionsDev.sampleCorrelation_EE_G12().data(); + auto const G6SamplesCorrelation = isBarrel ? conditionsDev.sampleCorrelation_EB_G6().data() : conditionsDev.sampleCorrelation_EE_G6().data(); - auto const* G1SamplesCorrelation = + auto const G1SamplesCorrelation = isBarrel ? conditionsDev.sampleCorrelation_EB_G1().data() : conditionsDev.sampleCorrelation_EE_G1().data(); auto const hasGainSwitch = hasSwitchToGain6[ch] || hasSwitchToGain1[ch] || isSaturated[ch]; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationKernels.dev.cc b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationKernels.dev.cc index 64aefd0109db8..7eef3c0de1f37 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationKernels.dev.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/AmplitudeComputationKernels.dev.cc @@ -84,7 +84,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { auto const nchannels = nchannelsEB + digisDevEE.size(); auto const offsetForHashes = conditionsDev.offsetEE(); - auto const* pulse_covariance = reinterpret_cast(conditionsDev.pulseCovariance()); + auto const* pulse_covariance = + reinterpret_cast(conditionsDev.pulseCovariance().data()); // shared memory DataType* shrmem = alpaka::getDynSharedMem(acc); @@ -101,15 +102,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { DataType* shrAtAStorage = shrmem + calo::multifit::MapSymM::total * (elemIdx + elemsPerBlock); + // uncalibRecHitsEE.outOfTimeAmplitudes() returns a span, where EcalOotAmpArray is an array of float + // uncalibRecHitsEE.outOfTimeAmplitudes().data() returns a pointer to the first EcalOotAmpArray of the data column + // uncalibRecHitsEE.outOfTimeAmplitudes().data()->data() returns a pointer to the first float of the first EcalOotAmpArray of the data column auto* amplitudes = - reinterpret_cast(idx >= nchannelsEB ? uncalibRecHitsEE.outOfTimeAmplitudes()->data() - : uncalibRecHitsEB.outOfTimeAmplitudes()->data()); - auto* energies = idx >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); - auto* chi2s = idx >= nchannelsEB ? uncalibRecHitsEE.chi2() : uncalibRecHitsEB.chi2(); + reinterpret_cast(idx >= nchannelsEB ? uncalibRecHitsEE.outOfTimeAmplitudes().data()->data() + : uncalibRecHitsEB.outOfTimeAmplitudes().data()->data()); + auto energies = idx >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); + auto chi2s = idx >= nchannelsEB ? uncalibRecHitsEE.chi2() : uncalibRecHitsEB.chi2(); // get the hash int const inputCh = idx >= nchannelsEB ? idx - nchannelsEB : idx; - auto const* dids = idx >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); + auto const dids = idx >= nchannelsEB ? digisDevEE.id() : digisDevEB.id(); auto const did = DetId{dids[inputCh]}; auto const isBarrel = did.subdetId() == EcalBarrel; auto const hashedId = isBarrel ? ecal::reconstruction::hashedIndexEB(did.rawId()) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/EnergyComputationKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/EnergyComputationKernels.h index 59ea4f43ccff4..e793a553a5da6 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/EnergyComputationKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/EnergyComputationKernels.h @@ -74,18 +74,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::rechit { } makeRecHit(ch, - uncalibRecHits.id(), - uncalibRecHits.amplitude(), - uncalibRecHits.amplitudeError(), - uncalibRecHits.jitter(), - uncalibRecHits.aux(), - uncalibRecHits.chi2(), - uncalibRecHits.flags(), - recHits.id(), - recHits.energy(), - recHits.time(), - recHits.flagBits(), - recHits.extra(), + uncalibRecHits.id().data(), + uncalibRecHits.amplitude().data(), + uncalibRecHits.amplitudeError().data(), + uncalibRecHits.jitter().data(), + uncalibRecHits.aux().data(), + uncalibRecHits.chi2().data(), + uncalibRecHits.flags().data(), + recHits.id().data(), + recHits.energy().data(), + recHits.time().data(), + recHits.flagBits().data(), + recHits.extra().data(), conditionsDev, parametersDev, eventTime, @@ -130,20 +130,20 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::rechit { int const inputCh = isEndcap ? ch - nchannelsEB : ch; // inputs - auto const* didCh = isEndcap ? eeUncalibRecHits.id() : ebUncalibRecHits.id(); - auto const* amplitude = isEndcap ? eeUncalibRecHits.amplitude() : ebUncalibRecHits.amplitude(); - auto const* amplitudeError = isEndcap ? eeUncalibRecHits.amplitudeError() : ebUncalibRecHits.amplitudeError(); - auto const* jitter = isEndcap ? eeUncalibRecHits.jitter() : ebUncalibRecHits.jitter(); - auto const* aux = isEndcap ? eeUncalibRecHits.aux() : ebUncalibRecHits.aux(); - auto const* chi2_in = isEndcap ? eeUncalibRecHits.chi2() : ebUncalibRecHits.chi2(); - auto const* flags_in = isEndcap ? eeUncalibRecHits.flags() : ebUncalibRecHits.flags(); + auto const didCh = isEndcap ? eeUncalibRecHits.id() : ebUncalibRecHits.id(); + auto const amplitude = isEndcap ? eeUncalibRecHits.amplitude() : ebUncalibRecHits.amplitude(); + auto const amplitudeError = isEndcap ? eeUncalibRecHits.amplitudeError() : ebUncalibRecHits.amplitudeError(); + auto const jitter = isEndcap ? eeUncalibRecHits.jitter() : ebUncalibRecHits.jitter(); + auto const aux = isEndcap ? eeUncalibRecHits.aux() : ebUncalibRecHits.aux(); + auto const chi2_in = isEndcap ? eeUncalibRecHits.chi2() : ebUncalibRecHits.chi2(); + auto const flags_in = isEndcap ? eeUncalibRecHits.flags() : ebUncalibRecHits.flags(); // outputs - auto* did = isEndcap ? eeRecHits.id() : ebRecHits.id(); - auto* energy = isEndcap ? eeRecHits.energy() : ebRecHits.energy(); - auto* time = isEndcap ? eeRecHits.time() : ebRecHits.time(); - auto* flagBits = isEndcap ? eeRecHits.flagBits() : ebRecHits.flagBits(); - auto* extra = isEndcap ? eeRecHits.extra() : ebRecHits.extra(); + auto did = isEndcap ? eeRecHits.id() : ebRecHits.id(); + auto energy = isEndcap ? eeRecHits.energy() : ebRecHits.energy(); + auto time = isEndcap ? eeRecHits.time() : ebRecHits.time(); + auto flagBits = isEndcap ? eeRecHits.flagBits() : ebRecHits.flagBits(); + auto extra = isEndcap ? eeRecHits.extra() : ebRecHits.extra(); bool const recoverIsolatedChannels = isEndcap ? configParams.recoverEEIsolatedChannels : configParams.recoverEBIsolatedChannels; @@ -153,18 +153,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::rechit { float const laserMAX = isEndcap ? configParams.EELaserMAX : configParams.EBLaserMAX; makeRecHit(inputCh, - didCh, - amplitude, - amplitudeError, - jitter, - aux, - chi2_in, - flags_in, - did, - energy, - time, - flagBits, - extra, + didCh.data(), + amplitude.data(), + amplitudeError.data(), + jitter.data(), + aux.data(), + chi2_in.data(), + flags_in.data(), + did.data(), + energy.data(), + time.data(), + flagBits.data(), + extra.data(), conditionsDev, parametersDev, eventTime, diff --git a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h index a8bd81cef2391..30bd93d599556 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h @@ -168,13 +168,13 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { auto const ch_start = ch * nsamples; auto const inputCh = ch >= offsetForInputs ? ch - offsetForInputs : ch; - auto const* dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); + auto const dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); auto const did = DetId{dids[inputCh]}; auto const isBarrel = did.subdetId() == EcalBarrel; - auto* const amplitudeFitParameters = + auto const amplitudeFitParameters = isBarrel ? paramsDev->amplitudeFitParamsEB.data() : paramsDev->amplitudeFitParamsEE.data(); - auto* const timeFitParameters = + auto const timeFitParameters = isBarrel ? paramsDev->timeFitParamsEB.data() : paramsDev->timeFitParamsEE.data(); auto const timeFitParameters_size = isBarrel ? paramsDev->timeFitParamsEB.size() : paramsDev->timeFitParamsEE.size(); @@ -565,13 +565,13 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { auto const elemIdx = gtx % elemsPerBlock; auto const sample = elemIdx % nsamples; - auto const* dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); + auto const dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); auto const inputCh = ch >= offsetForInputs ? ch - offsetForInputs : ch; auto state = g_state[ch]; auto const did = DetId{dids[inputCh]}; - auto* const amplitudeFitParameters = did.subdetId() == EcalBarrel ? paramsDev->amplitudeFitParamsEB.data() - : paramsDev->amplitudeFitParamsEE.data(); + auto const amplitudeFitParameters = did.subdetId() == EcalBarrel ? paramsDev->amplitudeFitParamsEB.data() + : paramsDev->amplitudeFitParamsEE.data(); // TODO is that better than storing into global and launching another kernel // for the first 10 threads @@ -755,7 +755,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { continue; int const inputGtx = ch >= offsetForInputs ? gtx - offsetForInputs * nsamples : gtx; - auto const* digis = ch >= offsetForInputs ? digisDevEE.data()->data() : digisDevEB.data()->data(); + // digisDevEE.data() returns a span, where EcalDataArray is an array of uint16_t + // digisDevEE.data().data() returns a pointer to the first EcalDataArray of the data column + // digisDevEE.data().data()->data() returns a pointer to the first uint16_t of the first EcalDataArray of the data column + auto const* digis = ch >= offsetForInputs ? digisDevEE.data().data()->data() : digisDevEB.data().data()->data(); auto const gainIdPrev = ecalMGPA::gainId(digis[inputGtx - 1]); auto const gainIdNext = ecalMGPA::gainId(digis[inputGtx]); @@ -805,8 +808,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { int const inputTx = ch >= offsetForInputs ? tx - offsetForInputs * nsamples : tx; int const inputCh = ch >= offsetForInputs ? ch - offsetForInputs : ch; - auto const* digis = ch >= offsetForInputs ? digisDevEE.data()->data() : digisDevEB.data()->data(); - auto const* dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); + auto const* digis = ch >= offsetForInputs ? digisDevEE.data().data()->data() : digisDevEB.data().data()->data(); + auto const dids = ch >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); // indices/inits auto const sample = tx % nsamples; @@ -986,23 +989,24 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit { for (auto gtx : cms::alpakatools::uniform_elements(acc, nchannels)) { const int inputGtx = gtx >= offsetForInputs ? gtx - offsetForInputs : gtx; - auto const* dids = gtx >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); - auto const* digis = gtx >= offsetForInputs ? digisDevEE.data()->data() : digisDevEB.data()->data(); + auto const dids = gtx >= offsetForInputs ? digisDevEE.id() : digisDevEB.id(); + auto const* digis = + gtx >= offsetForInputs ? digisDevEE.data().data()->data() : digisDevEB.data().data()->data(); - auto* g_amplitude = gtx >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); - auto* g_jitter = gtx >= nchannelsEB ? uncalibRecHitsEE.jitter() : uncalibRecHitsEB.jitter(); - auto* g_jitterError = gtx >= nchannelsEB ? uncalibRecHitsEE.jitterError() : uncalibRecHitsEB.jitterError(); - auto* flags = gtx >= nchannelsEB ? uncalibRecHitsEE.flags() : uncalibRecHitsEB.flags(); + auto g_amplitude = gtx >= nchannelsEB ? uncalibRecHitsEE.amplitude() : uncalibRecHitsEB.amplitude(); + auto g_jitter = gtx >= nchannelsEB ? uncalibRecHitsEE.jitter() : uncalibRecHitsEB.jitter(); + auto g_jitterError = gtx >= nchannelsEB ? uncalibRecHitsEE.jitterError() : uncalibRecHitsEB.jitterError(); + auto flags = gtx >= nchannelsEB ? uncalibRecHitsEE.flags() : uncalibRecHitsEB.flags(); auto const did = DetId{dids[inputGtx]}; auto const isBarrel = did.subdetId() == EcalBarrel; auto const hashedId = isBarrel ? ecal::reconstruction::hashedIndexEB(did.rawId()) : offsetForHashes + ecal::reconstruction::hashedIndexEE(did.rawId()); // need to access the underlying data directly here because the std::arrays have different size for EB and EE, which is not compatible with the ? operator - auto* const amplitudeBins = isBarrel ? conditionsDev.timeBiasCorrections_amplitude_EB().data() - : conditionsDev.timeBiasCorrections_amplitude_EE().data(); - auto* const shiftBins = isBarrel ? conditionsDev.timeBiasCorrections_shift_EB().data() - : conditionsDev.timeBiasCorrections_shift_EE().data(); + auto const amplitudeBins = isBarrel ? conditionsDev.timeBiasCorrections_amplitude_EB().data() + : conditionsDev.timeBiasCorrections_amplitude_EE().data(); + auto const shiftBins = isBarrel ? conditionsDev.timeBiasCorrections_shift_EB().data() + : conditionsDev.timeBiasCorrections_shift_EE().data(); auto const amplitudeBinsSize = isBarrel ? conditionsDev.timeBiasCorrectionSizeEB() : conditionsDev.timeBiasCorrectionSizeEE(); auto const timeConstantTerm = isBarrel ? timeConstantTermEB : timeConstantTermEE; diff --git a/RecoLocalCalo/HGCalRecAlgos/interface/HGCalESProducerTools.h b/RecoLocalCalo/HGCalRecAlgos/interface/HGCalESProducerTools.h index 89ef2fc1411ba..dadc9066ebf86 100644 --- a/RecoLocalCalo/HGCalRecAlgos/interface/HGCalESProducerTools.h +++ b/RecoLocalCalo/HGCalRecAlgos/interface/HGCalESProducerTools.h @@ -25,7 +25,7 @@ namespace hgcal { // @short fill SoA column with data from vector for any type with some offset template void fill_SoA_column( - T* column_SoA, const std::vector& values, const int offset, const int nrows, int arr_offset = 0) { + std::span column_SoA, const std::vector& values, const int offset, const int nrows, int arr_offset = 0) { const int nrows_vals = values.size(); if (arr_offset < 0) { arr_offset = 0; diff --git a/RecoLocalCalo/HGCalRecAlgos/plugins/alpaka/HGCalRecHitCalibrationESProducer.cc b/RecoLocalCalo/HGCalRecAlgos/plugins/alpaka/HGCalRecHitCalibrationESProducer.cc index 978953ed0bff9..bf16aeff44fac 100644 --- a/RecoLocalCalo/HGCalRecAlgos/plugins/alpaka/HGCalRecHitCalibrationESProducer.cc +++ b/RecoLocalCalo/HGCalRecAlgos/plugins/alpaka/HGCalRecHitCalibrationESProducer.cc @@ -191,7 +191,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { << "layer=" << layer << ", celltype=" << celltype << ", isSiPM=" << isSiPM << ", dEdx=" << dEdx << ", sf=" << sf << std::endl; dEdx *= sf * 1e3; // apply correction and convert from MeV to GeV - fill_SoA_column_single(product.view().EM_scale(), dEdx, offset, nrows); + fill_SoA_column_single(product.view().EM_scale().data(), dEdx, offset, nrows); } // end of loop over modules diff --git a/RecoLocalCalo/HGCalRecProducers/plugins/alpaka/HGCalLayerClustersAlgoWrapper.dev.cc b/RecoLocalCalo/HGCalRecProducers/plugins/alpaka/HGCalLayerClustersAlgoWrapper.dev.cc index bd672655308a2..b63f56fc27ef0 100644 --- a/RecoLocalCalo/HGCalRecProducers/plugins/alpaka/HGCalLayerClustersAlgoWrapper.dev.cc +++ b/RecoLocalCalo/HGCalRecProducers/plugins/alpaka/HGCalLayerClustersAlgoWrapper.dev.cc @@ -38,17 +38,17 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { alpaka::memset(queue, isSeed, 0x0); algoStandalone.makeClustersCMSSW(size, - inputs.dim1(), - inputs.dim2(), - inputs.layer(), - inputs.energy(), - inputs.sigmaNoise(), - inputs.detid(), - outputs.rho(), - outputs.delta(), - outputs.nearestHigher(), - outputs.clusterIndex(), - outputs.isSeed(), + inputs.dim1().data(), + inputs.dim2().data(), + inputs.layer().data(), + inputs.energy().data(), + inputs.sigmaNoise().data(), + inputs.detid().data(), + outputs.rho().data(), + outputs.delta().data(), + outputs.nearestHigher().data(), + outputs.clusterIndex().data(), + outputs.isSeed().data(), &outputs.numberOfClustersScalar()); } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 1e2e3ad235b79..af979c7c29452 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -71,7 +71,7 @@ namespace pixelgpudetails { slopeCol = 1; rowOffset = 2 * pixelgpudetails::numRowsInRoc - 1; colOffset = (rocIdInDetUnit - 8) * pixelgpudetails::numColsInRoc; - } // if roc + } // if roc } else { // +Z side: 4 non-flipped modules oriented like 'pppp', but all 8 in layer1 if (rocIdInDetUnit < 8) { slopeRow = -1; @@ -583,15 +583,15 @@ namespace pixelgpudetails { (std::max(int(wordCounter), int(TrackerTraits::numberOfModules)) + threadsPerBlock - 1) / threadsPerBlock; gpuCalibPixel::calibDigis<<>>(clusterThresholds, - digis_d.view().moduleId(), - digis_d.view().xx(), - digis_d.view().yy(), - digis_d.view().adc(), + digis_d.view().moduleId().data(), + digis_d.view().xx().data(), + digis_d.view().yy().data(), + digis_d.view().adc().data(), gains, wordCounter, - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->clusModuleStart()); + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->clusModuleStart().data()); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG @@ -604,7 +604,7 @@ namespace pixelgpudetails { #endif countModules<<>>( - digis_d->moduleId(), clusters_d->moduleStart(), digis_d->clus(), wordCounter); + digis_d->moduleId().data(), clusters_d->moduleStart().data(), digis_d->clus().data(), wordCounter); cudaCheck(cudaGetLastError()); // should be larger than maxPixInModule/16 aka (maxPixInModule/maxiter in the kernel) @@ -614,14 +614,14 @@ namespace pixelgpudetails { std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - findClus<<>>(digis_d->rawIdArr(), - digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), + findClus<<>>(digis_d->rawIdArr().data(), + digis_d->moduleId().data(), + digis_d->xx().data(), + digis_d->yy().data(), + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->moduleId().data(), + digis_d->clus().data(), wordCounter); cudaCheck(cudaGetLastError()); @@ -631,12 +631,12 @@ namespace pixelgpudetails { // apply charge cut clusterChargeCut<<>>(clusterThresholds, - digis_d->moduleId(), - digis_d->adc(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), + digis_d->moduleId().data(), + digis_d->adc().data(), + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->moduleId().data(), + digis_d->clus().data(), wordCounter); cudaCheck(cudaGetLastError()); @@ -647,9 +647,9 @@ namespace pixelgpudetails { // synchronization/ExternalWork auto nModules_Clusters_d = cms::cuda::make_device_unique(3, stream); // MUST be ONE block - fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), - clusters_d->clusModuleStart(), - clusters_d->moduleStart(), + fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule().data(), + clusters_d->clusModuleStart().data(), + clusters_d->moduleStart().data(), nModules_Clusters_d.get()); // copy to host @@ -679,12 +679,15 @@ namespace pixelgpudetails { nDigis = numDigis; digis_d = SiPixelDigisCUDA(numDigis, stream); - cudaCheck(cudaMemcpyAsync(digis_d->moduleId(), moduleIds, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(digis_d->xx(), xDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(digis_d->yy(), yDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(digis_d->adc(), adcDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(digis_d->pdigi(), packedData, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(digis_d->rawIdArr(), rawIds, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck( + cudaMemcpyAsync(digis_d->moduleId().data(), moduleIds, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck(cudaMemcpyAsync(digis_d->xx().data(), xDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck(cudaMemcpyAsync(digis_d->yy().data(), yDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck(cudaMemcpyAsync(digis_d->adc().data(), adcDigis, sizeof(uint16_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck( + cudaMemcpyAsync(digis_d->pdigi().data(), packedData, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); + cudaCheck( + cudaMemcpyAsync(digis_d->rawIdArr().data(), rawIds, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); clusters_d = SiPixelClustersCUDA(TrackerTraits::numberOfModules, stream); @@ -694,12 +697,12 @@ namespace pixelgpudetails { int blocks = (int(numDigis) + threadsPerBlock - 1) / threadsPerBlock; gpuCalibPixel::calibDigisPhase2<<>>(clusterThresholds, - digis_d->moduleId(), - digis_d->adc(), + digis_d->moduleId().data(), + digis_d->adc().data(), numDigis, - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->clusModuleStart()); + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->clusModuleStart().data()); cudaCheck(cudaGetLastError()); @@ -709,12 +712,12 @@ namespace pixelgpudetails { #endif countModules<<>>( - digis_d->moduleId(), clusters_d->moduleStart(), digis_d->clus(), numDigis); + digis_d->moduleId().data(), clusters_d->moduleStart().data(), digis_d->clus().data(), numDigis); cudaCheck(cudaGetLastError()); // read the number of modules into a data member, used by getProduct()) cudaCheck(cudaMemcpyAsync( - &(nModules_Clusters_h[0]), clusters_d->moduleStart(), sizeof(uint32_t), cudaMemcpyDefault, stream)); + &(nModules_Clusters_h[0]), clusters_d->moduleStart().data(), sizeof(uint32_t), cudaMemcpyDefault, stream)); threadsPerBlock = 256; blocks = TrackerTraits::numberOfModules; @@ -723,14 +726,14 @@ namespace pixelgpudetails { cudaCheck(cudaStreamSynchronize(stream)); std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - findClus<<>>(digis_d->rawIdArr(), - digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), + findClus<<>>(digis_d->rawIdArr().data(), + digis_d->moduleId().data(), + digis_d->xx().data(), + digis_d->yy().data(), + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->moduleId().data(), + digis_d->clus().data(), numDigis); cudaCheck(cudaGetLastError()); @@ -742,12 +745,12 @@ namespace pixelgpudetails { // apply charge cut clusterChargeCut<<>>(clusterThresholds, - digis_d->moduleId(), - digis_d->adc(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), + digis_d->moduleId().data(), + digis_d->adc().data(), + clusters_d->moduleStart().data(), + clusters_d->clusInModule().data(), + clusters_d->moduleId().data(), + digis_d->clus().data(), numDigis); cudaCheck(cudaGetLastError()); @@ -759,9 +762,9 @@ namespace pixelgpudetails { #endif // MUST be ONE block - fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), - clusters_d->clusModuleStart(), - clusters_d->moduleStart(), + fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule().data(), + clusters_d->clusModuleStart().data(), + clusters_d->moduleStart().data(), nModules_Clusters_d.get()); nModules_Clusters_h = cms::cuda::make_host_unique(3, stream); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc index 6f204a142ca6b..811d22df124eb 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc @@ -308,7 +308,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { *(regions_->modulesToUnpack()), cabling_.get(), fedIds_, iEvent.queue()); modulesToUnpack = modulesToUnpackRegional->data(); } else { - modulesToUnpack = hMap->modToUnpDefault(); + modulesToUnpack = hMap->modToUnpDefault().data(); } const auto& buffers = iEvent.get(rawGetToken_); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc index 1a500bb2c7883..7925176668d97 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc @@ -622,20 +622,20 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { alpaka::exec(queue, workDivPrefixScan, cms::alpakatools::multiBlockPrefixScan(), - clusters_d->view().clusInModule(), - clusters_d->view().clusModuleStart() + 1, + clusters_d->view().clusInModule().data(), + clusters_d->view().clusModuleStart().data() + 1, TrackerTraits::numberOfModules, blocksPrefixScan, bCounter.data(), alpaka::getPreferredWarpSize(alpaka::getDev(queue))); // last element holds the number of all clusters - const auto clusModuleStartLastElement = - cms::alpakatools::make_device_view(queue, clusters_d->const_view().clusModuleStart() + numberOfModules, 1u); + const auto clusModuleStartLastElement = cms::alpakatools::make_device_view( + queue, clusters_d->const_view().clusModuleStart().data() + numberOfModules, 1u); constexpr int startBPIX2 = TrackerTraits::layerStart[1]; // element startBPIX2 hold the number of clusters until BPIX2 - const auto bpix2ClusterStart = - cms::alpakatools::make_device_view(queue, clusters_d->const_view().clusModuleStart() + startBPIX2, 1u); + const auto bpix2ClusterStart = cms::alpakatools::make_device_view( + queue, clusters_d->const_view().clusModuleStart().data() + startBPIX2, 1u); auto nModules_Clusters_h_1 = cms::alpakatools::make_host_view(nModules_Clusters_h.data() + 1, 1u); alpaka::memcpy(queue, nModules_Clusters_h_1, clusModuleStartLastElement); @@ -731,20 +731,20 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { alpaka::exec(queue, workDivPrefixScan, cms::alpakatools::multiBlockPrefixScan(), - clusters_d->view().clusInModule(), - clusters_d->view().clusModuleStart() + 1, + clusters_d->view().clusInModule().data(), + clusters_d->view().clusModuleStart().data() + 1, TrackerTraits::numberOfModules, blocksPrefixScan, bCounter.data(), alpaka::getPreferredWarpSize(alpaka::getDev(queue))); // last element holds the number of all clusters - const auto clusModuleStartLastElement = - cms::alpakatools::make_device_view(queue, clusters_d->const_view().clusModuleStart() + numberOfModules, 1u); + const auto clusModuleStartLastElement = cms::alpakatools::make_device_view( + queue, clusters_d->const_view().clusModuleStart().data() + numberOfModules, 1u); constexpr int startBPIX2 = pixelTopology::Phase2::layerStart[1]; // element startBPIX2 hold the number of clusters until BPIX2 const auto bpix2ClusterStart = - cms::alpakatools::make_device_view(queue, clusters_d->const_view().clusModuleStart() + startBPIX2, 1u); + cms::alpakatools::make_device_view(queue, clusters_d->const_view().clusModuleStart().data() + startBPIX2, 1u); auto nModules_Clusters_h_1 = cms::alpakatools::make_host_view(nModules_Clusters_h.data() + 1, 1u); alpaka::memcpy(queue, nModules_Clusters_h_1, clusModuleStartLastElement); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu index b1e5e1c3c90e9..e0086bb32f72d 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu @@ -53,7 +53,7 @@ namespace pixelgpudetails { auto nHits = clusters_d.nClusters(); TrackingRecHitSoADevice hits_d( - nHits, clusters_d.offsetBPIX2(), cpeParams, clusters_d->clusModuleStart(), stream); + nHits, clusters_d.offsetBPIX2(), cpeParams, clusters_d->clusModuleStart().data(), stream); int activeModulesWithDigis = digis_d.nModules(); // protect from empty events @@ -73,17 +73,17 @@ namespace pixelgpudetails { // assuming full warp of threads is better than a smaller number... if (nHits) { - setHitsLayerStart - <<<1, 32, 0, stream>>>(clusters_d->clusModuleStart(), cpeParams, hits_d.view().hitsLayerStart().data()); + setHitsLayerStart<<<1, 32, 0, stream>>>( + clusters_d->clusModuleStart().data(), cpeParams, hits_d.view().hitsLayerStart().data()); cudaCheck(cudaGetLastError()); constexpr auto nLayers = TrackerTraits::numberOfLayers; cms::cuda::fillManyFromVector(&(hits_d.view().phiBinner()), nLayers, - hits_d.view().iphi(), + hits_d.view().iphi().data(), hits_d.view().hitsLayerStart().data(), nHits, 256, - hits_d.view().phiBinnerStorage(), + hits_d.view().phiBinnerStorage().data(), stream); cudaCheck(cudaGetLastError()); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 21da864c1c348..22b3daf7defa1 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -121,10 +121,10 @@ void SiPixelRecHitSoAFromLegacyT::produce(edm::StreamID streamID, cms::cuda::PortableHostCollection> clusters_h(nModules + 1); - memset(clusters_h.view().clusInModule(), 0, (nModules + 1) * sizeof(uint32_t)); // needed?? - memset(clusters_h.view().moduleStart(), 0, (nModules + 1) * sizeof(uint32_t)); - memset(clusters_h.view().moduleId(), 0, (nModules + 1) * sizeof(uint32_t)); - memset(clusters_h.view().clusModuleStart(), 0, (nModules + 1) * sizeof(uint32_t)); + memset(clusters_h.view().clusInModule().data(), 0, (nModules + 1) * sizeof(uint32_t)); // needed?? + memset(clusters_h.view().moduleStart().data(), 0, (nModules + 1) * sizeof(uint32_t)); + memset(clusters_h.view().moduleId().data(), 0, (nModules + 1) * sizeof(uint32_t)); + memset(clusters_h.view().clusModuleStart().data(), 0, (nModules + 1) * sizeof(uint32_t)); assert(0 == clusters_h.view()[nModules].clusInModule()); clusters_h.view()[1].moduleStart() = 0; @@ -151,8 +151,10 @@ void SiPixelRecHitSoAFromLegacyT::produce(edm::StreamID streamID, assert((uint32_t)numberOfClusters == clusters_h.view()[nModules].clusModuleStart()); // output SoA // element 96 is the start of BPIX2 (i.e. the number of clusters in BPIX1) - HitsOnHost output( - numberOfClusters, clusters_h.view()[startBPIX2].clusModuleStart(), &cpeView, clusters_h.view().clusModuleStart()); + HitsOnHost output(numberOfClusters, + clusters_h.view()[startBPIX2].clusModuleStart(), + &cpeView, + clusters_h.view().clusModuleStart().data()); if (0 == numberOfClusters) { iEvent.emplace(tokenHit_, std::move(output)); @@ -268,18 +270,18 @@ void SiPixelRecHitSoAFromLegacyT::produce(edm::StreamID streamID, cms::cuda::fillManyFromVector(&(output.view().phiBinner()), nLayers, - output.view().iphi(), + output.view().iphi().data(), output.view().hitsLayerStart().data(), output.view().nHits(), 256, - output.view().phiBinnerStorage()); + output.view().phiBinnerStorage().data()); LogDebug("SiPixelRecHitSoAFromLegacy") << "created HitSoa for " << numberOfClusters << " clusters in " << numberOfDetUnits << " Dets" << "\n"; // copy pointer to data (SoA view) to allocated buffer - memcpy(hitsModuleStart, clusters_h.view().clusModuleStart(), nModules * sizeof(uint32_t)); + memcpy(hitsModuleStart, clusters_h.view().clusModuleStart().data(), nModules * sizeof(uint32_t)); iEvent.emplace(tokenHit_, std::move(output)); if (convert2Legacy_) diff --git a/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc b/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc index 7778fb7b9bd71..e144f88f2a208 100644 --- a/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc +++ b/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc @@ -582,7 +582,7 @@ void L2TauNNProducer::selectGoodTracksAndVertices(const ZVertexSoAHost& patavtx_ trkGood.reserve(maxTracks); vtxGood.clear(); vtxGood.reserve(nv); - auto const* quality = patatracks_tsoa.view().quality(); + auto const quality = patatracks_tsoa.view().quality(); // No need to sort either as the algorithms is just using the max (not even the location, just the max value of pt2sum). std::vector pTSquaredSum(nv, 0); diff --git a/RecoTauTag/HLTProducers/src/L2TauTagNNProducerAlpaka.cc b/RecoTauTag/HLTProducers/src/L2TauTagNNProducerAlpaka.cc index af99e53b96934..d9df55a52fca0 100644 --- a/RecoTauTag/HLTProducers/src/L2TauTagNNProducerAlpaka.cc +++ b/RecoTauTag/HLTProducers/src/L2TauTagNNProducerAlpaka.cc @@ -581,7 +581,7 @@ void L2TauNNProducerAlpaka::selectGoodTracksAndVertices(const ZVertexHost& patav trkGood.reserve(maxTracks); vtxGood.clear(); vtxGood.reserve(nv); - auto const* quality = patatracks_tsoa.view().quality(); + auto const quality = patatracks_tsoa.view().quality(); // No need to sort either as the algorithms is just using the max (not even the location, just the max value of pt2sum). std::vector pTSquaredSum(nv, 0); diff --git a/RecoTracker/LSTCore/interface/LSTPrepareInput.h b/RecoTracker/LSTCore/interface/LSTPrepareInput.h index 423272c76ecc9..92ecb14505499 100644 --- a/RecoTracker/LSTCore/interface/LSTPrepareInput.h +++ b/RecoTracker/LSTCore/interface/LSTPrepareInput.h @@ -210,40 +210,40 @@ namespace lst { LSTInputHostCollection lstInputHC(soa_sizes, queue); auto hits = lstInputHC.view(); - std::memcpy(hits.xs(), ph2_x.data(), nHitsOT * sizeof(float)); - std::memcpy(hits.ys(), ph2_y.data(), nHitsOT * sizeof(float)); - std::memcpy(hits.zs(), ph2_z.data(), nHitsOT * sizeof(float)); - std::memcpy(hits.detid(), ph2_detId.data(), nHitsOT * sizeof(unsigned int)); + std::memcpy(hits.xs().data(), ph2_x.data(), nHitsOT * sizeof(float)); + std::memcpy(hits.ys().data(), ph2_y.data(), nHitsOT * sizeof(float)); + std::memcpy(hits.zs().data(), ph2_z.data(), nHitsOT * sizeof(float)); + std::memcpy(hits.detid().data(), ph2_detId.data(), nHitsOT * sizeof(unsigned int)); #ifndef LST_STANDALONE - std::memcpy(hits.hits(), ph2_hits.data(), nHitsOT * sizeof(TrackingRecHit const*)); + std::memcpy(hits.hits().data(), ph2_hits.data(), nHitsOT * sizeof(TrackingRecHit const*)); #endif - std::memcpy(hits.xs() + nHitsOT, trkX.data(), nHitsIT * sizeof(float)); - std::memcpy(hits.ys() + nHitsOT, trkY.data(), nHitsIT * sizeof(float)); - std::memcpy(hits.zs() + nHitsOT, trkZ.data(), nHitsIT * sizeof(float)); - std::memcpy(hits.detid() + nHitsOT, hitId.data(), nHitsIT * sizeof(unsigned int)); + std::memcpy(hits.xs().data() + nHitsOT, trkX.data(), nHitsIT * sizeof(float)); + std::memcpy(hits.ys().data() + nHitsOT, trkY.data(), nHitsIT * sizeof(float)); + std::memcpy(hits.zs().data() + nHitsOT, trkZ.data(), nHitsIT * sizeof(float)); + std::memcpy(hits.detid().data() + nHitsOT, hitId.data(), nHitsIT * sizeof(unsigned int)); #ifndef LST_STANDALONE - std::memset(hits.hits() + nHitsOT, 0, nHitsIT * sizeof(TrackingRecHit const*)); + std::memset(hits.hits().data() + nHitsOT, 0, nHitsIT * sizeof(TrackingRecHit const*)); #endif - std::memcpy(hits.idxs(), hitIdxs.data(), (nHitsIT + nHitsOT) * sizeof(unsigned int)); + std::memcpy(hits.idxs().data(), hitIdxs.data(), (nHitsIT + nHitsOT) * sizeof(unsigned int)); auto pixelSeeds = lstInputHC.view(); - std::memcpy(pixelSeeds.hitIndices(), hitIndices_vec.data(), nPixelSeeds * sizeof(Params_pLS::ArrayUxHits)); - std::memcpy(pixelSeeds.deltaPhi(), deltaPhi_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.ptIn(), ptIn_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.ptErr(), ptErr_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.px(), px_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.py(), py_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.pz(), pz_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.etaErr(), etaErr_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.isQuad(), isQuad_vec.data(), nPixelSeeds * sizeof(char)); - std::memcpy(pixelSeeds.eta(), eta_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.phi(), phi_vec.data(), nPixelSeeds * sizeof(float)); - std::memcpy(pixelSeeds.charge(), charge_vec.data(), nPixelSeeds * sizeof(int)); - std::memcpy(pixelSeeds.seedIdx(), seedIdx_vec.data(), nPixelSeeds * sizeof(unsigned int)); - std::memcpy(pixelSeeds.superbin(), superbin_vec.data(), nPixelSeeds * sizeof(int)); - std::memcpy(pixelSeeds.pixelType(), pixelType_vec.data(), nPixelSeeds * sizeof(PixelType)); + std::memcpy(pixelSeeds.hitIndices().data(), hitIndices_vec.data(), nPixelSeeds * sizeof(Params_pLS::ArrayUxHits)); + std::memcpy(pixelSeeds.deltaPhi().data(), deltaPhi_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.ptIn().data(), ptIn_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.ptErr().data(), ptErr_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.px().data(), px_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.py().data(), py_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.pz().data(), pz_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.etaErr().data(), etaErr_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.isQuad().data(), isQuad_vec.data(), nPixelSeeds * sizeof(char)); + std::memcpy(pixelSeeds.eta().data(), eta_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.phi().data(), phi_vec.data(), nPixelSeeds * sizeof(float)); + std::memcpy(pixelSeeds.charge().data(), charge_vec.data(), nPixelSeeds * sizeof(int)); + std::memcpy(pixelSeeds.seedIdx().data(), seedIdx_vec.data(), nPixelSeeds * sizeof(unsigned int)); + std::memcpy(pixelSeeds.superbin().data(), superbin_vec.data(), nPixelSeeds * sizeof(int)); + std::memcpy(pixelSeeds.pixelType().data(), pixelType_vec.data(), nPixelSeeds * sizeof(PixelType)); return lstInputHC; } diff --git a/RecoTracker/LSTCore/src/LSTESData.cc b/RecoTracker/LSTCore/src/LSTESData.cc index dad8522bbe2cd..4f74ac486e1b1 100644 --- a/RecoTracker/LSTCore/src/LSTESData.cc +++ b/RecoTracker/LSTCore/src/LSTESData.cc @@ -92,10 +92,10 @@ std::unique_ptr> lst::loadAndFillESHost(s auto endcapGeometryDev = std::make_shared(endcapGeometry.nEndCapMap, cms::alpakatools::host()); - std::memcpy(endcapGeometryDev->view().geoMapDetId(), + std::memcpy(endcapGeometryDev->view().geoMapDetId().data(), endcapGeometry.geoMapDetId_buf.data(), endcapGeometry.nEndCapMap * sizeof(unsigned int)); - std::memcpy(endcapGeometryDev->view().geoMapPhi(), + std::memcpy(endcapGeometryDev->view().geoMapPhi().data(), endcapGeometry.geoMapPhi_buf.data(), endcapGeometry.nEndCapMap * sizeof(float)); diff --git a/RecoTracker/LSTCore/src/ModuleMethods.h b/RecoTracker/LSTCore/src/ModuleMethods.h index 13421dbd7e011..107e0ce0724ca 100644 --- a/RecoTracker/LSTCore/src/ModuleMethods.h +++ b/RecoTracker/LSTCore/src/ModuleMethods.h @@ -25,13 +25,13 @@ namespace lst { // https://github.com/cms-sw/cmssw/blob/5e809e8e0a625578aa265dc4b128a93830cb5429/Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h#L29 }; - bool parseIsLower(bool isInvertedx, unsigned int detId) { return (isInvertedx) ? !(detId & 1) : (detId & 1); } + inline bool parseIsLower(bool isInvertedx, unsigned int detId) { return (isInvertedx) ? !(detId & 1) : (detId & 1); } - unsigned int parsePartnerModuleId(unsigned int detId, bool isLowerx, bool isInvertedx) { + inline unsigned int parsePartnerModuleId(unsigned int detId, bool isLowerx, bool isInvertedx) { return isLowerx ? (isInvertedx ? detId - 1 : detId + 1) : (isInvertedx ? detId + 1 : detId - 1); } - bool parseIsInverted(short subdet, short side, short module, short layer) { + inline bool parseIsInverted(short subdet, short side, short module, short layer) { if (subdet == Endcap) { if (side == NegZ) { return module % 2 == 1; @@ -130,8 +130,8 @@ namespace lst { inline void fillConnectedModuleArrayExplicit(Modules modules, ModuleMetaData const& mmd, ModuleConnectionMap const& moduleConnectionMap) { - Params_Modules::ArrayU16xMaxConnected* moduleMap = modules.moduleMap(); - uint16_t* nConnectedModules = modules.nConnectedModules(); + std::span moduleMap = modules.moduleMap(); + std::span nConnectedModules = modules.nConnectedModules(); for (auto it = mmd.detIdToIndex.begin(); it != mmd.detIdToIndex.end(); ++it) { unsigned int detId = it->first; @@ -145,8 +145,8 @@ namespace lst { } inline void fillMapArraysExplicit(Modules modules, ModuleMetaData const& mmd) { - uint16_t* mapIdx = modules.mapIdx(); - unsigned int* mapdetId = modules.mapdetId(); + std::span mapIdx = modules.mapIdx(); + std::span mapdetId = modules.mapdetId(); unsigned int counter = 0; for (auto it = mmd.detIdToIndex.begin(); it != mmd.detIdToIndex.end(); ++it) { @@ -246,26 +246,26 @@ namespace lst { auto modules_view = modulesHC->view(); // Getting the underlying data pointers - unsigned int* host_detIds = modules_view.detIds(); - short* host_layers = modules_view.layers(); - short* host_rings = modules_view.rings(); - short* host_rods = modules_view.rods(); - short* host_modules = modules_view.modules(); - short* host_subdets = modules_view.subdets(); - short* host_sides = modules_view.sides(); - float* host_eta = modules_view.eta(); - float* host_r = modules_view.r(); - bool* host_isInverted = modules_view.isInverted(); - bool* host_isLower = modules_view.isLower(); - bool* host_isAnchor = modules_view.isAnchor(); - ModuleType* host_moduleType = modules_view.moduleType(); - ModuleLayerType* host_moduleLayerType = modules_view.moduleLayerType(); - float* host_dxdys = modules_view.dxdys(); - float* host_drdzs = modules_view.drdzs(); + std::span host_detIds = modules_view.detIds(); + std::span host_layers = modules_view.layers(); + std::span host_rings = modules_view.rings(); + std::span host_rods = modules_view.rods(); + std::span host_modules = modules_view.modules(); + std::span host_subdets = modules_view.subdets(); + std::span host_sides = modules_view.sides(); + std::span host_eta = modules_view.eta(); + std::span host_r = modules_view.r(); + std::span host_isInverted = modules_view.isInverted(); + std::span host_isLower = modules_view.isLower(); + std::span host_isAnchor = modules_view.isAnchor(); + std::span host_moduleType = modules_view.moduleType(); + std::span host_moduleLayerType = modules_view.moduleLayerType(); + std::span host_dxdys = modules_view.dxdys(); + std::span host_drdzs = modules_view.drdzs(); uint16_t* host_nModules = &modules_view.nModules(); uint16_t* host_nLowerModules = &modules_view.nLowerModules(); - uint16_t* host_partnerModuleIndices = modules_view.partnerModuleIndices(); - int* host_lstLayers = modules_view.lstLayers(); + std::span host_partnerModuleIndices = modules_view.partnerModuleIndices(); + std::span host_lstLayers = modules_view.lstLayers(); //reassign detIdToIndex indices here nLowerModules = (nModules - 1) / 2; diff --git a/RecoTracker/LSTCore/src/alpaka/Hit.h b/RecoTracker/LSTCore/src/alpaka/Hit.h index 6ef8e4330fdda..40a6af2a24b64 100644 --- a/RecoTracker/LSTCore/src/alpaka/Hit.h +++ b/RecoTracker/LSTCore/src/alpaka/Hit.h @@ -64,17 +64,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { alpaka::math::acosh(acc, alpaka::math::sqrt(acc, ihit_x * ihit_x + ihit_y * ihit_y + ihit_z * ihit_z) / hitsExtended.rts()[ihit]); - auto found_pointer = alpaka_std::lower_bound(modules.mapdetId(), modules.mapdetId() + nModules, iDetId); - ALPAKA_ASSERT_ACC(found_pointer != modules.mapdetId() + nModules); - int found_index = std::distance(modules.mapdetId(), found_pointer); + auto found_pointer = + alpaka_std::lower_bound(modules.mapdetId().data(), modules.mapdetId().data() + nModules, iDetId); + ALPAKA_ASSERT_ACC(found_pointer != modules.mapdetId().data() + nModules); + int found_index = std::distance(modules.mapdetId().data(), found_pointer); uint16_t lastModuleIndex = modules.mapIdx()[found_index]; hitsExtended.moduleIndices()[ihit] = lastModuleIndex; if (modules.subdets()[lastModuleIndex] == Endcap && modules.moduleType()[lastModuleIndex] == TwoS) { - found_pointer = alpaka_std::lower_bound(geoMapDetId, geoMapDetId + nEndCapMap, iDetId); - ALPAKA_ASSERT_ACC(found_pointer != geoMapDetId + nEndCapMap); - found_index = std::distance(geoMapDetId, found_pointer); + found_pointer = alpaka_std::lower_bound(geoMapDetId.data(), geoMapDetId.data() + nEndCapMap, iDetId); + ALPAKA_ASSERT_ACC(found_pointer != geoMapDetId.data() + nEndCapMap); + found_index = std::distance(geoMapDetId.data(), found_pointer); float phi = geoMapPhi[found_index]; float cos_phi = alpaka::math::cos(acc, phi); hitsExtended.highEdgeXs()[ihit] = ihit_x + 2.5f * cos_phi; diff --git a/RecoTracker/LSTCore/src/alpaka/LSTEvent.dev.cc b/RecoTracker/LSTCore/src/alpaka/LSTEvent.dev.cc index dc3da6962df89..66def3403b016 100644 --- a/RecoTracker/LSTCore/src/alpaka/LSTEvent.dev.cc +++ b/RecoTracker/LSTCore/src/alpaka/LSTEvent.dev.cc @@ -202,15 +202,14 @@ void LSTEvent::createMiniDoublets() { miniDoubletsDC_.emplace(mds_sizes, queue_); auto mdsOccupancy = miniDoubletsDC_->view(); - auto nMDs_view = cms::alpakatools::make_device_view(queue_, mdsOccupancy.nMDs(), mdsOccupancy.metadata().size()); - auto totOccupancyMDs_view = - cms::alpakatools::make_device_view(queue_, mdsOccupancy.totOccupancyMDs(), mdsOccupancy.metadata().size()); + auto nMDs_view = cms::alpakatools::make_device_view(queue_, mdsOccupancy.nMDs()); + auto totOccupancyMDs_view = cms::alpakatools::make_device_view(queue_, mdsOccupancy.totOccupancyMDs()); alpaka::memset(queue_, nMDs_view, 0u); alpaka::memset(queue_, totOccupancyMDs_view, 0u); } auto mdView = miniDoubletsDC_->view(); - auto connView = cms::alpakatools::make_device_view(queue_, mdView.connectedMax(), mdView.metadata().size()); + auto connView = cms::alpakatools::make_device_view(queue_, mdView.connectedMax()); alpaka::memset(queue_, connView, 0u); unsigned int mdSize = pixelSize_ * 2; @@ -293,13 +292,12 @@ void LSTEvent::createSegmentsWithModuleMap() { auto segmentsOccupancy = segmentsDC_->view(); auto segments = segmentsDC_->view(); - auto nSegments_view = - cms::alpakatools::make_device_view(queue_, segmentsOccupancy.nSegments(), segmentsOccupancy.metadata().size()); - auto totOccupancySegments_view = cms::alpakatools::make_device_view( - queue_, segmentsOccupancy.totOccupancySegments(), segmentsOccupancy.metadata().size()); + auto nSegments_view = cms::alpakatools::make_device_view(queue_, segmentsOccupancy.nSegments()); + auto totOccupancySegments_view = + cms::alpakatools::make_device_view(queue_, segmentsOccupancy.totOccupancySegments()); alpaka::memset(queue_, nSegments_view, 0u); alpaka::memset(queue_, totOccupancySegments_view, 0u); - auto conn_view = cms::alpakatools::make_device_view(queue_, segments.connectedMax(), segments.metadata().size()); + auto conn_view = cms::alpakatools::make_device_view(queue_, segments.connectedMax()); alpaka::memset(queue_, conn_view, 0u); auto src_view_size = cms::alpakatools::make_host_view(pixelSize_); @@ -375,21 +373,19 @@ void LSTEvent::createTriplets() { tripletsDC_.emplace(triplets_sizes, queue_); auto tripletsOccupancy = tripletsDC_->view(); - auto nTriplets_view = - cms::alpakatools::make_device_view(queue_, tripletsOccupancy.nTriplets(), tripletsOccupancy.metadata().size()); + auto nTriplets_view = cms::alpakatools::make_device_view(queue_, tripletsOccupancy.nTriplets()); alpaka::memset(queue_, nTriplets_view, 0u); - auto totOccupancyTriplets_view = cms::alpakatools::make_device_view( - queue_, tripletsOccupancy.totOccupancyTriplets(), tripletsOccupancy.metadata().size()); + auto totOccupancyTriplets_view = + cms::alpakatools::make_device_view(queue_, tripletsOccupancy.totOccupancyTriplets()); alpaka::memset(queue_, totOccupancyTriplets_view, 0u); auto triplets = tripletsDC_->view(); - auto partOfPT5_view = cms::alpakatools::make_device_view(queue_, triplets.partOfPT5(), triplets.metadata().size()); + auto partOfPT5_view = cms::alpakatools::make_device_view(queue_, triplets.partOfPT5()); alpaka::memset(queue_, partOfPT5_view, 0u); - auto partOfT5_view = cms::alpakatools::make_device_view(queue_, triplets.partOfT5(), triplets.metadata().size()); + auto partOfT5_view = cms::alpakatools::make_device_view(queue_, triplets.partOfT5()); alpaka::memset(queue_, partOfT5_view, 0u); - auto partOfPT3_view = cms::alpakatools::make_device_view(queue_, triplets.partOfPT3(), triplets.metadata().size()); + auto partOfPT3_view = cms::alpakatools::make_device_view(queue_, triplets.partOfPT3()); alpaka::memset(queue_, partOfPT3_view, 0u); - auto connectedMax_view = - cms::alpakatools::make_device_view(queue_, triplets.connectedMax(), triplets.metadata().size()); + auto connectedMax_view = cms::alpakatools::make_device_view(queue_, triplets.connectedMax()); alpaka::memset(queue_, connectedMax_view, 0u); } @@ -781,20 +777,17 @@ void LSTEvent::createQuintuplets() { std::array const quintuplets_sizes{{static_cast(nTotalQuintuplets), static_cast(nLowerModules_)}}; quintupletsDC_.emplace(quintuplets_sizes, queue_); auto quintupletsOccupancy = quintupletsDC_->view(); - auto nQuintuplets_view = cms::alpakatools::make_device_view( - queue_, quintupletsOccupancy.nQuintuplets(), quintupletsOccupancy.metadata().size()); + auto nQuintuplets_view = cms::alpakatools::make_device_view(queue_, quintupletsOccupancy.nQuintuplets()); alpaka::memset(queue_, nQuintuplets_view, 0u); - auto totOccupancyQuintuplets_view = cms::alpakatools::make_device_view( - queue_, quintupletsOccupancy.totOccupancyQuintuplets(), quintupletsOccupancy.metadata().size()); + auto totOccupancyQuintuplets_view = + cms::alpakatools::make_device_view(queue_, quintupletsOccupancy.totOccupancyQuintuplets()); alpaka::memset(queue_, totOccupancyQuintuplets_view, 0u); auto quintuplets = quintupletsDC_->view(); - auto isDup_view = cms::alpakatools::make_device_view(queue_, quintuplets.isDup(), quintuplets.metadata().size()); + auto isDup_view = cms::alpakatools::make_device_view(queue_, quintuplets.isDup()); alpaka::memset(queue_, isDup_view, 0u); - auto tightCutFlag_view = - cms::alpakatools::make_device_view(queue_, quintuplets.tightCutFlag(), quintuplets.metadata().size()); + auto tightCutFlag_view = cms::alpakatools::make_device_view(queue_, quintuplets.tightCutFlag()); alpaka::memset(queue_, tightCutFlag_view, 0u); - auto partOfPT5_view = - cms::alpakatools::make_device_view(queue_, quintuplets.partOfPT5(), quintuplets.metadata().size()); + auto partOfPT5_view = cms::alpakatools::make_device_view(queue_, quintuplets.partOfPT5()); alpaka::memset(queue_, partOfPT5_view, 0u); } diff --git a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.h b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.h index 7a3e938fae3ec..e88d8dd13f27e 100644 --- a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.h +++ b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.h @@ -173,9 +173,9 @@ __global__ void kernel_BLFit(TupleMultiplicity const *__restrict_ double *__restrict__ phits, float *__restrict__ phits_ge, double *__restrict__ pfast_fit) { - assert(results_view.pt()); - assert(results_view.eta()); - assert(results_view.chi2()); + assert(results_view.pt().data()); + assert(results_view.eta().data()); + assert(results_view.chi2().data()); assert(pfast_fit); constexpr auto invalidTkId = std::numeric_limits::max(); diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h index 4642e794ac4f9..f5f57dc24a87d 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h @@ -359,7 +359,7 @@ namespace caHitNtupletGeneratorKernels { oc.setStatusBits(Cell::StatusBit::kUsed); } } // loop on inner cells - } // loop on outer cells + } // loop on outer cells } template @@ -405,7 +405,7 @@ namespace caHitNtupletGeneratorKernels { *cellTracks, tracks_view.hitIndices(), *apc, - tracks_view.quality(), + tracks_view.quality().data(), stack, params.minHitsPerNtuplet_, bpix1Start); diff --git a/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc b/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc index c36ed924911f0..4ca868aacf6bb 100644 --- a/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc @@ -10,8 +10,8 @@ void HelixFitOnGPU::allocateOnGPU(TupleMultiplicity const *tupleM assert(tuples_); assert(tupleMultiplicity_); - assert(outputSoa_.chi2()); - assert(outputSoa_.pt()); + assert(outputSoa_.chi2().data()); + assert(outputSoa_.pt().data()); } template diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc index 60317b3da3fd6..feb11ab3f2eba 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc @@ -180,9 +180,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { // workaround for #47808 debug::do_not_optimise(results_view); - ALPAKA_ASSERT_ACC(results_view.pt()); - ALPAKA_ASSERT_ACC(results_view.eta()); - ALPAKA_ASSERT_ACC(results_view.chi2()); + ALPAKA_ASSERT_ACC(results_view.pt().data()); + ALPAKA_ASSERT_ACC(results_view.eta().data()); + ALPAKA_ASSERT_ACC(results_view.chi2().data()); ALPAKA_ASSERT_ACC(pfast_fit); constexpr auto invalidTkId = std::numeric_limits::max(); diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc index c7558009fe5b2..bfaaeb419eef5 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.dev.cc @@ -199,7 +199,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { cms::alpakatools::fillManyFromVector(device_hitPhiHist_->data(), device_hitPhiView_, TrackerTraits::numberOfLayers, // could be ll.metadata().size() - 1 - hh.iphi(), + hh.iphi().data(), this->device_layerStarts_->data(), hh.metadata().size(), (uint32_t)256, diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h index ca6e3eefc83d9..5158f42fba51f 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernelsImpl.h @@ -528,7 +528,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels { nCellTracks, ct, *apc, - tracks_view.quality(), + tracks_view.quality().data(), stack, params.minHitsPerNtuplet_); ALPAKA_ASSERT_ACC(stack.empty()); diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/HelixFit.cc b/RecoTracker/PixelSeeding/plugins/alpaka/HelixFit.cc index 9185bd8b2fb94..e8507dd92a838 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/HelixFit.cc +++ b/RecoTracker/PixelSeeding/plugins/alpaka/HelixFit.cc @@ -12,7 +12,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { ALPAKA_ASSERT_ACC(tuples_); ALPAKA_ASSERT_ACC(tupleMultiplicity_); - ALPAKA_ASSERT_ACC(helix_fit_results.pt()); + ALPAKA_ASSERT_ACC(helix_fit_results.pt().data()); } template diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpAlpaka.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpAlpaka.cc index 2bd9204fb977c..6f4cc62edb662 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpAlpaka.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpAlpaka.cc @@ -49,23 +49,23 @@ void PixelTrackDumpAlpakaT::analyze(edm::StreamID streamID, edm::Event const& iEvent, const edm::EventSetup& iSetup) const { auto const& tracks = iEvent.get(tokenSoATrack_); - assert(tracks.view().quality()); - assert(tracks.view().chi2()); - assert(tracks.view().nLayers()); - assert(tracks.view().eta()); - assert(tracks.view().pt()); - assert(tracks.view().state()); - assert(tracks.view().covariance()); + assert(tracks.view().quality().data()); + assert(tracks.view().chi2().data()); + assert(tracks.view().nLayers().data()); + assert(tracks.view().eta().data()); + assert(tracks.view().pt().data()); + assert(tracks.view().state().data()); + assert(tracks.view().covariance().data()); assert(tracks.view().nTracks()); auto const& vertices = iEvent.get(tokenSoAVertex_); - assert(vertices.view().idv()); - assert(vertices.view().zv()); - assert(vertices.view().wv()); - assert(vertices.view().chi2()); - assert(vertices.view().ptv2()); - assert(vertices.view().ndof()); - assert(vertices.view().sortInd()); + assert(vertices.view().idv().data()); + assert(vertices.view().zv().data()); + assert(vertices.view().wv().data()); + assert(vertices.view().chi2().data()); + assert(vertices.view().ptv2().data()); + assert(vertices.view().ndof().data()); + assert(vertices.view().sortInd().data()); assert(vertices.view().nvFinal()); } diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index ef65feb4f8b28..4eff8bec4dc2f 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -167,7 +167,7 @@ void PixelTrackProducerFromSoAT::produce(edm::StreamID streamID, hits.reserve(5); auto const &tsoa = iEvent.get(tokenTrack_); - auto const *quality = tsoa.view().quality(); + auto const quality = tsoa.view().quality(); auto const &hitIndices = tsoa.view().hitIndices(); auto nTracks = tsoa.view().nTracks(); diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoAAlpaka.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoAAlpaka.cc index acb780850b638..d368ed6b552a2 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoAAlpaka.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoAAlpaka.cc @@ -163,9 +163,9 @@ void PixelTrackProducerFromSoAAlpaka::produce(edm::StreamID streamID, hits.reserve(5); //TODO move to a configurable parameter? auto const &tsoa = iEvent.get(tokenTrack_); - auto const *quality = tsoa.view().quality(); - auto const *hitOffs = tsoa.view().hitOffsets(); - auto const *hitIdxs = tsoa.template view().id(); + auto const quality = tsoa.view().quality(); + auto const hitOffs = tsoa.view().hitOffsets(); + auto const hitIdxs = tsoa.template view().id(); // auto const &hitIndices = tsoa.view().hitIndices(); auto nTracks = tsoa.view().nTracks(); diff --git a/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc b/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc index 620d9973396d7..8122ccbac3130 100644 --- a/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc +++ b/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc @@ -99,7 +99,7 @@ void SeedProducerFromSoAT::produce(edm::StreamID streamID, auto const& tsoa = iEvent.get(tokenTrack_); - auto const* quality = tsoa.view().quality(); + auto const quality = tsoa.view().quality(); auto const& detIndices = tsoa.view().detIndices(); auto maxTracks = tsoa.view().metadata().size(); diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h index 412bdc4dd183a..c2368d10ce639 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h @@ -46,11 +46,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder { ALPAKA_ASSERT_ACC(static_cast(nt) <= ws.metadata().size()); ALPAKA_ASSERT_ACC(static_cast(nt) <= trkdata.metadata().size()); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ iv = ws.iv(); - int32_t* __restrict__ nn = trkdata.ndof(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ iv = ws.iv().data(); + int32_t* __restrict__ nn = trkdata.ndof().data(); ALPAKA_ASSERT_ACC(zt); ALPAKA_ASSERT_ACC(ezt2); ALPAKA_ASSERT_ACC(izt); diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksDBSCAN.h b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksDBSCAN.h index 875db26ff7f0b..f2c579a220ca3 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksDBSCAN.h +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksDBSCAN.h @@ -45,11 +45,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder { ALPAKA_ASSERT_ACC(static_cast(nt) <= ws.metadata().size()); ALPAKA_ASSERT_ACC(static_cast(nt) <= trkdata.metadata().size()); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ iv = ws.iv(); - int32_t* __restrict__ nn = trkdata.ndof(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ iv = ws.iv().data(); + int32_t* __restrict__ nn = trkdata.ndof().data(); ALPAKA_ASSERT_ACC(zt); ALPAKA_ASSERT_ACC(ezt2); ALPAKA_ASSERT_ACC(izt); diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksIterative.h b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksIterative.h index 23004ec98e1ca..8c739fc101651 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksIterative.h +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksIterative.h @@ -41,11 +41,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder { ALPAKA_ASSERT_ACC(static_cast(nt) <= ws.metadata().size()); ALPAKA_ASSERT_ACC(static_cast(nt) <= trkdata.metadata().size()); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ iv = ws.iv(); - int32_t* __restrict__ nn = trkdata.ndof(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ iv = ws.iv().data(); + int32_t* __restrict__ nn = trkdata.ndof().data(); ALPAKA_ASSERT_ACC(zt); ALPAKA_ASSERT_ACC(ezt2); ALPAKA_ASSERT_ACC(izt); diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/fitVertices.h b/RecoVertex/PixelVertexFinding/plugins/alpaka/fitVertices.h index 1b019a1ba811a..198e7d052dd4d 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/fitVertices.h +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/fitVertices.h @@ -27,16 +27,16 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder { auto& __restrict__ trkdata = ptrkdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - float* __restrict__ zv = data.zv(); - float* __restrict__ wv = data.wv(); - float* __restrict__ chi2 = data.chi2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + float* __restrict__ zv = data.zv().data(); + float* __restrict__ wv = data.wv().data(); + float* __restrict__ chi2 = data.chi2().data(); uint32_t& nvFinal = data.nvFinal(); uint32_t& nvIntermediate = ws.nvIntermediate(); - int32_t* __restrict__ nn = trkdata.ndof(); - int32_t* __restrict__ iv = ws.iv(); + int32_t* __restrict__ nn = trkdata.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); ALPAKA_ASSERT_ACC(nvFinal <= nvIntermediate); nvFinal = nvIntermediate; diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/sortByPt2.h b/RecoVertex/PixelVertexFinding/plugins/alpaka/sortByPt2.h index 95242d84c5895..36d65161fb26d 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/sortByPt2.h +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/sortByPt2.h @@ -26,12 +26,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder { ALPAKA_FN_ACC ALPAKA_FN_INLINE void sortByPt2(Acc1D const& acc, VtxSoAView& data, TrkSoAView& trkdata, WsSoAView& ws) { auto nt = ws.ntrks(); - float const* __restrict__ ptt2 = ws.ptt2(); + float const* __restrict__ ptt2 = ws.ptt2().data(); uint32_t const& nvFinal = data.nvFinal(); - int32_t const* __restrict__ iv = ws.iv(); - float* __restrict__ ptv2 = data.ptv2(); - uint16_t* __restrict__ sortInd = data.sortInd(); + int32_t const* __restrict__ iv = ws.iv().data(); + float* __restrict__ ptv2 = data.ptv2().data(); + uint16_t* __restrict__ sortInd = data.sortInd().data(); if (nvFinal < 1) return; diff --git a/RecoVertex/PixelVertexFinding/plugins/alpaka/vertexFinder.dev.cc b/RecoVertex/PixelVertexFinding/plugins/alpaka/vertexFinder.dev.cc index a9fc2a51808fb..64a146a3eb400 100644 --- a/RecoVertex/PixelVertexFinding/plugins/alpaka/vertexFinder.dev.cc +++ b/RecoVertex/PixelVertexFinding/plugins/alpaka/vertexFinder.dev.cc @@ -35,7 +35,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { WsSoAView ws, float ptMin, float ptMax) const { - auto const* quality = tracks_view.quality(); + auto const quality = tracks_view.quality(); for (auto idx : cms::alpakatools::uniform_elements(acc, tracks_view.nTracks())) { [[maybe_unused]] auto nHits = reco::nHits(tracks_view, idx); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h index aaca1c5cc268f..e38c4bd495899 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h @@ -35,15 +35,15 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); uint32_t& nvFinal = data.nvFinal(); uint32_t& nvIntermediate = ws.nvIntermediate(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ nn = data.ndof(); - int32_t* __restrict__ iv = ws.iv(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ nn = data.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); assert(zt); assert(ezt2); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h index d3fc38baf48b8..588fabd3d0d77 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h @@ -31,15 +31,15 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); uint32_t& nvFinal = data.nvFinal(); uint32_t& nvIntermediate = ws.nvIntermediate(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ nn = data.ndof(); - int32_t* __restrict__ iv = ws.iv(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ nn = data.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); assert(zt); assert(iv); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksIterative.h b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksIterative.h index 9a8ef262db767..5a56e23eedeee 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksIterative.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuClusterTracksIterative.h @@ -31,15 +31,15 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); uint32_t& nvFinal = data.nvFinal(); uint32_t& nvIntermediate = ws.nvIntermediate(); - uint8_t* __restrict__ izt = ws.izt(); - int32_t* __restrict__ nn = data.ndof(); - int32_t* __restrict__ iv = ws.iv(); + uint8_t* __restrict__ izt = ws.izt().data(); + int32_t* __restrict__ nn = data.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); assert(zt); assert(nn); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuFitVertices.h b/RecoVertex/PixelVertexFinding/plugins/gpuFitVertices.h index e2a5a82fd6802..bf2652b8acd5f 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuFitVertices.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuFitVertices.h @@ -21,16 +21,16 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - float* __restrict__ zv = data.zv(); - float* __restrict__ wv = data.wv(); - float* __restrict__ chi2 = data.chi2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + float* __restrict__ zv = data.zv().data(); + float* __restrict__ wv = data.wv().data(); + float* __restrict__ chi2 = data.chi2().data(); uint32_t& nvFinal = data.nvFinal(); uint32_t& nvIntermediate = ws.nvIntermediate(); - int32_t* __restrict__ nn = data.ndof(); - int32_t* __restrict__ iv = ws.iv(); + int32_t* __restrict__ nn = data.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); assert(nvFinal <= nvIntermediate); nvFinal = nvIntermediate; diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuSortByPt2.h b/RecoVertex/PixelVertexFinding/plugins/gpuSortByPt2.h index c5a0b5cfc1776..4d9f2c42068d2 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuSortByPt2.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuSortByPt2.h @@ -19,12 +19,12 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ ptt2 = ws.ptt2(); + float const* __restrict__ ptt2 = ws.ptt2().data(); uint32_t const& nvFinal = data.nvFinal(); - int32_t const* __restrict__ iv = ws.iv(); - float* __restrict__ ptv2 = data.ptv2(); - uint16_t* __restrict__ sortInd = data.sortInd(); + int32_t const* __restrict__ iv = ws.iv().data(); + float* __restrict__ ptv2 = data.ptv2().data(); + uint16_t* __restrict__ sortInd = data.sortInd().data(); assert(ptv2); assert(sortInd); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuSplitVertices.h b/RecoVertex/PixelVertexFinding/plugins/gpuSplitVertices.h index 6538bb216737b..c5b14b8bdcdd0 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuSplitVertices.h +++ b/RecoVertex/PixelVertexFinding/plugins/gpuSplitVertices.h @@ -18,15 +18,15 @@ namespace gpuVertexFinder { auto& __restrict__ data = pdata; auto& __restrict__ ws = pws; auto nt = ws.ntrks(); - float const* __restrict__ zt = ws.zt(); - float const* __restrict__ ezt2 = ws.ezt2(); - float* __restrict__ zv = data.zv(); - float* __restrict__ wv = data.wv(); - float const* __restrict__ chi2 = data.chi2(); + float const* __restrict__ zt = ws.zt().data(); + float const* __restrict__ ezt2 = ws.ezt2().data(); + float* __restrict__ zv = data.zv().data(); + float* __restrict__ wv = data.wv().data(); + float const* __restrict__ chi2 = data.chi2().data(); uint32_t& nvFinal = data.nvFinal(); - int32_t const* __restrict__ nn = data.ndof(); - int32_t* __restrict__ iv = ws.iv(); + int32_t const* __restrict__ nn = data.ndof().data(); + int32_t* __restrict__ iv = ws.iv().data(); assert(zt); assert(wv); diff --git a/RecoVertex/PixelVertexFinding/plugins/gpuVertexFinder.cc b/RecoVertex/PixelVertexFinding/plugins/gpuVertexFinder.cc index 73fa1408a0aab..dc638187977ac 100644 --- a/RecoVertex/PixelVertexFinding/plugins/gpuVertexFinder.cc +++ b/RecoVertex/PixelVertexFinding/plugins/gpuVertexFinder.cc @@ -28,7 +28,7 @@ namespace gpuVertexFinder { template __global__ void loadTracks( TrackSoAConstView tracks_view, VtxSoAView soa, WsSoAView pws, float ptMin, float ptMax) { - auto const* quality = tracks_view.quality(); + auto const quality = tracks_view.quality(); using helper = TracksUtilities; auto first = blockIdx.x * blockDim.x + threadIdx.x; for (int idx = first, nt = tracks_view.nTracks(); idx < nt; idx += gridDim.x * blockDim.x) { diff --git a/RecoVertex/PixelVertexFinding/test/VertexFinder_t.h b/RecoVertex/PixelVertexFinding/test/VertexFinder_t.h index 93534f82e06c8..c2e165ddb1e98 100644 --- a/RecoVertex/PixelVertexFinding/test/VertexFinder_t.h +++ b/RecoVertex/PixelVertexFinding/test/VertexFinder_t.h @@ -153,17 +153,17 @@ int main() { auto nt = ev.ztrack.size(); #ifdef __CUDACC__ cudaCheck(cudaMemcpy(&ws_d.view().ntrks(), &nt, sizeof(uint32_t), cudaMemcpyHostToDevice)); - cudaCheck( - cudaMemcpy(ws_d.view().zt(), ev.ztrack.data(), sizeof(float) * ev.ztrack.size(), cudaMemcpyHostToDevice)); - cudaCheck( - cudaMemcpy(ws_d.view().ezt2(), ev.eztrack.data(), sizeof(float) * ev.eztrack.size(), cudaMemcpyHostToDevice)); - cudaCheck( - cudaMemcpy(ws_d.view().ptt2(), ev.pttrack.data(), sizeof(float) * ev.eztrack.size(), cudaMemcpyHostToDevice)); + cudaCheck(cudaMemcpy( + ws_d.view().zt().data(), ev.ztrack.data(), sizeof(float) * ev.ztrack.size(), cudaMemcpyHostToDevice)); + cudaCheck(cudaMemcpy( + ws_d.view().ezt2().data(), ev.eztrack.data(), sizeof(float) * ev.eztrack.size(), cudaMemcpyHostToDevice)); + cudaCheck(cudaMemcpy( + ws_d.view().ptt2().data(), ev.pttrack.data(), sizeof(float) * ev.eztrack.size(), cudaMemcpyHostToDevice)); #else ::memcpy(&ws_d.view().ntrks(), &nt, sizeof(uint32_t)); - ::memcpy(ws_d.view().zt(), ev.ztrack.data(), sizeof(float) * ev.ztrack.size()); - ::memcpy(ws_d.view().ezt2(), ev.eztrack.data(), sizeof(float) * ev.eztrack.size()); - ::memcpy(ws_d.view().ptt2(), ev.pttrack.data(), sizeof(float) * ev.eztrack.size()); + ::memcpy(ws_d.view().zt().data(), ev.ztrack.data(), sizeof(float) * ev.ztrack.size()); + ::memcpy(ws_d.view().ezt2().data(), ev.eztrack.data(), sizeof(float) * ev.eztrack.size()); + ::memcpy(ws_d.view().ptt2().data(), ev.pttrack.data(), sizeof(float) * ev.eztrack.size()); #endif std::cout << "M eps, pset " << kk << ' ' << eps << ' ' << (i % 4) << std::endl; @@ -232,18 +232,18 @@ int main() { nn = hnn; ind = hind; #else - zv = onGPU_d.view().zv(); - wv = onGPU_d.view().wv(); - ptv2 = onGPU_d.view().ptv2(); - nn = onGPU_d.view().ndof(); - ind = onGPU_d.view().sortInd(); + zv = onGPU_d.view().zv().data(); + wv = onGPU_d.view().wv().data(); + ptv2 = onGPU_d.view().ptv2().data(); + nn = onGPU_d.view().ndof().data(); + ind = onGPU_d.view().sortInd().data(); #endif #ifdef __CUDACC__ - cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof().data(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); #else - memcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float)); + memcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float)); #endif for (auto j = 0U; j < nv; ++j) @@ -257,12 +257,12 @@ int main() { #ifdef __CUDACC__ cms::cuda::launch(gpuVertexFinder::fitVerticesKernel, {1, 1024 - 256}, onGPU_d.view(), ws_d.view(), 50.f); cudaCheck(cudaMemcpy(&nv, &onGPU_d.view().nvFinal(), sizeof(uint32_t), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof().data(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); #else gpuVertexFinder::fitVertices(onGPU_d.view(), ws_d.view(), 50.f); nv = onGPU_d.view().nvFinal(); - memcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float)); + memcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float)); #endif for (auto j = 0U; j < nv; ++j) @@ -294,7 +294,7 @@ int main() { gpuVertexFinder::fitVertices(onGPU_d.view(), ws_d.view(), 5000.f); gpuVertexFinder::sortByPt2(onGPU_d.view(), ws_d.view()); nv = onGPU_d.view().nvFinal(); - memcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float)); + memcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float)); #endif if (nv == 0) { @@ -303,12 +303,12 @@ int main() { } #ifdef __CUDACC__ - cudaCheck(cudaMemcpy(zv, onGPU_d.view().zv(), nv * sizeof(float), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(wv, onGPU_d.view().wv(), nv * sizeof(float), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2(), nv * sizeof(float), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(ptv2, onGPU_d.view().ptv2(), nv * sizeof(float), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(ind, onGPU_d.view().sortInd(), nv * sizeof(uint16_t), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(zv, onGPU_d.view().zv().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(wv, onGPU_d.view().wv().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(chi2, onGPU_d.view().chi2().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(ptv2, onGPU_d.view().ptv2().data(), nv * sizeof(float), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(nn, onGPU_d.view().ndof().data(), nv * sizeof(int32_t), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(ind, onGPU_d.view().sortInd().data(), nv * sizeof(uint16_t), cudaMemcpyDeviceToHost)); #endif for (auto j = 0U; j < nv; ++j) if (nn[j] > 0) @@ -353,7 +353,7 @@ int main() { std::cout << "min max rms " << *mx.first << ' ' << *mx.second << ' ' << rms << std::endl; } // loop on events - } // lopp on ave vert + } // lopp on ave vert return 0; } diff --git a/RecoVertex/PixelVertexFinding/test/alpaka/VertexFinder_t.dev.cc b/RecoVertex/PixelVertexFinding/test/alpaka/VertexFinder_t.dev.cc index 4e2e500745bff..4d475bda53a1a 100644 --- a/RecoVertex/PixelVertexFinding/test/alpaka/VertexFinder_t.dev.cc +++ b/RecoVertex/PixelVertexFinding/test/alpaka/VertexFinder_t.dev.cc @@ -204,8 +204,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { if (vertices_h.view().ndof()[j] > 0) vertices_h.view().chi2()[j] /= float(vertices_h.view().ndof()[j]); { - auto mx = - std::minmax_element(vertices_h.view().chi2(), vertices_h.view().chi2() + vertices_h.view().nvFinal()); + auto mx = std::minmax_element(vertices_h.view().chi2().data(), + vertices_h.view().chi2().data() + vertices_h.view().nvFinal()); std::cout << "after fit nv, min max chi2 " << vertices_h.view().nvFinal() << " " << *mx.first << ' ' << *mx.second << std::endl; } @@ -224,8 +224,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { if (vertices_h.view().ndof()[j] > 0) vertices_h.view().chi2()[j] /= float(vertices_h.view().ndof()[j]); { - auto mx = - std::minmax_element(vertices_h.view().chi2(), vertices_h.view().chi2() + vertices_h.view().nvFinal()); + auto mx = std::minmax_element(vertices_h.view().chi2().data(), + vertices_h.view().chi2().data() + vertices_h.view().nvFinal()); std::cout << "before splitting nv, min max chi2 " << vertices_h.view().nvFinal() << " " << *mx.first << ' ' << *mx.second << std::endl; } @@ -271,21 +271,22 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { if (vertices_h.view().ndof()[j] > 0) vertices_h.view().chi2()[j] /= float(vertices_h.view().ndof()[j]); { - auto mx = - std::minmax_element(vertices_h.view().chi2(), vertices_h.view().chi2() + vertices_h.view().nvFinal()); + auto mx = std::minmax_element(vertices_h.view().chi2().data(), + vertices_h.view().chi2().data() + vertices_h.view().nvFinal()); std::cout << "nv, min max chi2 " << vertices_h.view().nvFinal() << " " << *mx.first << ' ' << *mx.second << std::endl; } { - auto mx = std::minmax_element(vertices_h.view().wv(), vertices_h.view().wv() + vertices_h.view().nvFinal()); + auto mx = std::minmax_element(vertices_h.view().wv().data(), + vertices_h.view().wv().data() + vertices_h.view().nvFinal()); std::cout << "min max error " << 1. / std::sqrt(*mx.first) << ' ' << 1. / std::sqrt(*mx.second) << std::endl; } { - auto mx = - std::minmax_element(vertices_h.view().ptv2(), vertices_h.view().ptv2() + vertices_h.view().nvFinal()); + auto mx = std::minmax_element(vertices_h.view().ptv2().data(), + vertices_h.view().ptv2().data() + vertices_h.view().nvFinal()); std::cout << "min max ptv2 " << *mx.first << ' ' << *mx.second << std::endl; std::cout << "min max ptv2 " << vertices_h.view().ptv2()[vertices_h.view().sortInd()[0]] << ' ' << vertices_h.view().ptv2()[vertices_h.view().sortInd()[vertices_h.view().nvFinal() - 1]]