From a839a446701d7b406f5d6c939d48120e991d207f Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Mon, 17 Jan 2022 16:07:19 +0100 Subject: [PATCH] [cudadev] Ran clang-format to format code. --- .../CUDADataFormats/SiPixelClustersCUDA.cc | 5 +- .../CUDADataFormats/SiPixelClustersCUDA.h | 67 +-- .../CUDADataFormats/SiPixelDigisCUDA.cc | 44 +- .../CUDADataFormats/SiPixelDigisCUDA.h | 127 ++-- .../TrackingRecHit2DHeterogeneous.h | 40 +- .../TrackingRecHit2DHostSOAStore.cc | 15 +- .../TrackingRecHit2DHostSOAStore.h | 13 +- .../CUDADataFormats/TrackingRecHit2DSOAView.h | 170 +++--- .../CondFormats/SiPixelROCsStatusAndMapping.h | 40 +- .../SiPixelROCsStatusAndMappingWrapper.cc | 20 +- .../SiPixelROCsStatusAndMappingWrapper.h | 27 +- src/cudadev/DataFormats/SoACommon.h | 476 +++++++-------- src/cudadev/DataFormats/SoALayout.h | 89 +-- src/cudadev/DataFormats/SoAView.h | 553 +++++++++--------- .../SiPixelRawToClusterGPUKernel.cu | 3 +- .../SiPixelRawToClusterGPUKernel.h | 2 +- src/cudadev/test/SoALayoutAndView_t.cu | 150 +++-- 17 files changed, 919 insertions(+), 922 deletions(-) diff --git a/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.cc b/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.cc index 90361b048..2c71cdabf 100644 --- a/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.cc +++ b/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.cc @@ -3,10 +3,9 @@ #include "CUDACore/host_unique_ptr.h" #include "CUDADataFormats/SiPixelClustersCUDA.h" -SiPixelClustersCUDA::SiPixelClustersCUDA(): data_d(), deviceLayout_(data_d.get(), 0), deviceView_(deviceLayout_) {} +SiPixelClustersCUDA::SiPixelClustersCUDA() : data_d(), deviceLayout_(data_d.get(), 0), deviceView_(deviceLayout_) {} SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream) : data_d(cms::cuda::make_device_unique(DeviceLayout::computeDataSize(maxModules), stream)), deviceLayout_(data_d.get(), maxModules), - deviceView_(deviceLayout_) -{} + deviceView_(deviceLayout_) {} diff --git a/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.h b/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.h index 714654098..9f7451239 100644 --- a/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.h +++ b/src/cudadev/CUDADataFormats/SiPixelClustersCUDA.h @@ -12,45 +12,40 @@ class SiPixelClustersCUDA { public: GENERATE_SOA_LAYOUT(DeviceLayoutTemplate, - SOA_COLUMN(uint32_t, moduleStart), // index of the first pixel of each module - SOA_COLUMN(uint32_t, clusInModule), // number of clusters found in each module - SOA_COLUMN(uint32_t, moduleId), // module id of each module - - // originally from rechits - SOA_COLUMN(uint32_t, clusModuleStart) // index of the first cluster of each module - ) - + SOA_COLUMN(uint32_t, moduleStart), // index of the first pixel of each module + SOA_COLUMN(uint32_t, clusInModule), // number of clusters found in each module + SOA_COLUMN(uint32_t, moduleId), // module id of each module + + // originally from rechits + SOA_COLUMN(uint32_t, clusModuleStart)) // index of the first cluster of each module + // We use all defaults for the template parameters. using DeviceLayout = DeviceLayoutTemplate<>; - GENERATE_SOA_VIEW(DeviceViewTemplate, - SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceLayout, deviceLayout)), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(deviceLayout, moduleStart), // index of the first pixel of each module - SOA_VIEW_VALUE(deviceLayout, clusInModule), // number of clusters found in each module - SOA_VIEW_VALUE(deviceLayout, moduleId), // module id of each module - - // originally from rechits - SOA_VIEW_VALUE(deviceLayout, clusModuleStart) // index of the first cluster of each module - ) - ) - + GENERATE_SOA_VIEW( + DeviceViewTemplate, + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceLayout, deviceLayout)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(deviceLayout, moduleStart), // index of the first pixel of each module + SOA_VIEW_VALUE(deviceLayout, clusInModule), // number of clusters found in each module + SOA_VIEW_VALUE(deviceLayout, moduleId), // module id of each module + + // originally from rechits + SOA_VIEW_VALUE(deviceLayout, clusModuleStart))) // index of the first cluster of each module + using DeviceView = DeviceViewTemplate<>; - - GENERATE_SOA_CONST_VIEW(DeviceConstViewTemplate, - SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceView, deviceView)), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(deviceView, moduleStart), // index of the first pixel of each module - SOA_VIEW_VALUE(deviceView, clusInModule), // number of clusters found in each module - SOA_VIEW_VALUE(deviceView, moduleId), // module id of each module - - // originally from rechits - SOA_VIEW_VALUE(deviceView, clusModuleStart) // index of the first cluster of each module - ) - ) - + + GENERATE_SOA_CONST_VIEW( + DeviceConstViewTemplate, + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceView, deviceView)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(deviceView, moduleStart), // index of the first pixel of each module + SOA_VIEW_VALUE(deviceView, clusInModule), // number of clusters found in each module + SOA_VIEW_VALUE(deviceView, moduleId), // module id of each module + + // originally from rechits + SOA_VIEW_VALUE(deviceView, clusModuleStart))) // index of the first cluster of each module + using DeviceConstView = DeviceConstViewTemplate<>; - + explicit SiPixelClustersCUDA(); explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream); ~SiPixelClustersCUDA() = default; @@ -77,10 +72,10 @@ class SiPixelClustersCUDA { DeviceConstView view() const { return DeviceConstView(deviceView_); } private: - cms::cuda::device::unique_ptr data_d; // Single SoA storage + cms::cuda::device::unique_ptr data_d; // Single SoA storage DeviceLayout deviceLayout_; DeviceView deviceView_; - + uint32_t nClusters_h = 0; }; diff --git a/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.cc b/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.cc index a816806f8..55837fa92 100644 --- a/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.cc +++ b/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.cc @@ -6,28 +6,22 @@ SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream) : data_d(cms::cuda::make_device_unique( - DeviceOnlyLayout::computeDataSize(maxFedWords) + - HostDeviceLayout::computeDataSize(maxFedWords), - stream)), + DeviceOnlyLayout::computeDataSize(maxFedWords) + HostDeviceLayout::computeDataSize(maxFedWords), stream)), deviceOnlyLayout_d(data_d.get(), maxFedWords), hostDeviceLayout_d(deviceOnlyLayout_d.soaMetadata().nextByte(), maxFedWords), deviceFullView_(deviceOnlyLayout_d, hostDeviceLayout_d), - devicePixelConstView_(deviceFullView_) -{} + devicePixelConstView_(deviceFullView_) {} SiPixelDigisCUDA::SiPixelDigisCUDA() - : data_d(),deviceOnlyLayout_d(), hostDeviceLayout_d(), deviceFullView_(), devicePixelConstView_() -{} + : data_d(), deviceOnlyLayout_d(), hostDeviceLayout_d(), deviceFullView_(), devicePixelConstView_() {} -SiPixelDigisCUDA::HostStore::HostStore() - : data_h(), hostLayout_(nullptr, 0), hostView_(hostLayout_) -{} +SiPixelDigisCUDA::HostStore::HostStore() : data_h(), hostLayout_(nullptr, 0), hostView_(hostLayout_) {} SiPixelDigisCUDA::HostStore::HostStore(size_t maxFedWords, cudaStream_t stream) - : data_h(cms::cuda::make_host_unique(SiPixelDigisCUDA::HostDeviceLayout::computeDataSize(maxFedWords), stream)), - hostLayout_(data_h.get(), maxFedWords), - hostView_(hostLayout_) -{} + : data_h(cms::cuda::make_host_unique(SiPixelDigisCUDA::HostDeviceLayout::computeDataSize(maxFedWords), + stream)), + hostLayout_(data_h.get(), maxFedWords), + hostView_(hostLayout_) {} void SiPixelDigisCUDA::HostStore::reset() { hostLayout_ = HostDeviceLayout(); @@ -38,7 +32,8 @@ void SiPixelDigisCUDA::HostStore::reset() { cms::cuda::host::unique_ptr SiPixelDigisCUDA::adcToHostAsync(cudaStream_t stream) const { auto ret = cms::cuda::make_host_unique(nDigis(), stream); // TODO: this is downgraded from cms::cuda::copyAsync as we copy data from within a block but not the full block. - cudaCheck(cudaMemcpyAsync(ret.get(), deviceFullView_.adc(), nDigis() * sizeof(decltype(ret[0])), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync( + ret.get(), deviceFullView_.adc(), nDigis() * sizeof(decltype(ret[0])), cudaMemcpyDeviceToHost, stream)); return ret; } @@ -49,13 +44,20 @@ SiPixelDigisCUDA::HostStore SiPixelDigisCUDA::dataToHostAsync(cudaStream_t strea HostStore ret(nDigis(), stream); auto rhlsm = ret.hostLayout_.soaMetadata(); auto hdlsm_d = hostDeviceLayout_d.soaMetadata(); - cudaCheck(cudaMemcpyAsync(rhlsm.addressOf_adc(), hdlsm_d.addressOf_adc(), nDigis_h * sizeof(*rhlsm.addressOf_adc()), - cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(rhlsm.addressOf_adc(), + hdlsm_d.addressOf_adc(), + nDigis_h * sizeof(*rhlsm.addressOf_adc()), + cudaMemcpyDeviceToHost, + stream)); // Copy the other columns, realigning the data in shorter arrays. clus is the first but all 3 columns (clus, pdigis, rawIdArr) have // the same geometry. - cudaCheck(cudaMemcpy2DAsync(rhlsm.addressOf_clus(), rhlsm.clusPitch(), - hdlsm_d.addressOf_clus(), hdlsm_d.clusPitch(), - 3 /* rows */, - nDigis() * sizeof(decltype (*ret.hostView_.clus())), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpy2DAsync(rhlsm.addressOf_clus(), + rhlsm.clusPitch(), + hdlsm_d.addressOf_clus(), + hdlsm_d.clusPitch(), + 3 /* rows */, + nDigis() * sizeof(decltype(*ret.hostView_.clus())), + cudaMemcpyDeviceToHost, + stream)); return ret; } \ No newline at end of file diff --git a/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.h b/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.h index e334c102f..734b3631b 100644 --- a/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.h +++ b/src/cudadev/CUDADataFormats/SiPixelDigisCUDA.h @@ -11,75 +11,71 @@ class SiPixelDigisCUDA { public: - GENERATE_SOA_LAYOUT(DeviceOnlyLayoutTemplate, - /* These are consumed by downstream device code */ - SOA_COLUMN(uint16_t, xx), /* local coordinates of each pixel */ - SOA_COLUMN(uint16_t, yy), /* */ - SOA_COLUMN(uint16_t, moduleInd) /* module id of each pixel */ + GENERATE_SOA_LAYOUT( + DeviceOnlyLayoutTemplate, + /* These are consumed by downstream device code */ + SOA_COLUMN(uint16_t, xx), /* local coordinates of each pixel */ + SOA_COLUMN(uint16_t, yy), /* */ + SOA_COLUMN(uint16_t, moduleInd) /* module id of each pixel */ ) - + using DeviceOnlyLayout = DeviceOnlyLayoutTemplate<>; - - GENERATE_SOA_LAYOUT(HostDeviceLayoutTemplate, - /* These are also transferred to host (see HostDataView) */ - SOA_COLUMN(uint16_t, adc), /* ADC of each pixel */ - SOA_COLUMN(int32_t, clus), /* cluster id of each pixel */ - /* These are for CPU output; should we (eventually) place them to a */ - /* separate product? */ - SOA_COLUMN(uint32_t, pdigi), /* packed digi (row, col, adc) of each pixel */ - SOA_COLUMN(uint32_t, rawIdArr) /* DetId of each pixel */ + + GENERATE_SOA_LAYOUT( + HostDeviceLayoutTemplate, + /* These are also transferred to host (see HostDataView) */ + SOA_COLUMN(uint16_t, adc), /* ADC of each pixel */ + SOA_COLUMN(int32_t, clus), /* cluster id of each pixel */ + /* These are for CPU output; should we (eventually) place them to a */ + /* separate product? */ + SOA_COLUMN(uint32_t, pdigi), /* packed digi (row, col, adc) of each pixel */ + SOA_COLUMN(uint32_t, rawIdArr) /* DetId of each pixel */ ) - + using HostDeviceLayout = HostDeviceLayoutTemplate<>; - + GENERATE_SOA_VIEW(HostDeviceViewTemplate, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */ - SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */ - SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */ - SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */ - ) - ) - + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)), + SOA_VIEW_VALUE_LIST( + SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */ + SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */ + SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */ + SOA_VIEW_VALUE(hostDevice, + rawIdArr) /* DetId of each pixel */ + )) + using HostDeviceView = HostDeviceViewTemplate<>; - - GENERATE_SOA_VIEW(DeviceFullViewTemplate, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(DeviceOnlyLayout, deviceOnly), - SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(deviceOnly, xx), /* local coordinates of each pixel */ - SOA_VIEW_VALUE(deviceOnly, yy), /* */ - SOA_VIEW_VALUE(deviceOnly, moduleInd),/* module id of each pixel */ - SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */ - SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */ - SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */ - SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */ - ) - ) - + + GENERATE_SOA_VIEW( + DeviceFullViewTemplate, + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceOnlyLayout, deviceOnly), + SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)), + SOA_VIEW_VALUE_LIST( + SOA_VIEW_VALUE(deviceOnly, xx), /* local coordinates of each pixel */ + SOA_VIEW_VALUE(deviceOnly, yy), /* */ + SOA_VIEW_VALUE(deviceOnly, moduleInd), /* module id of each pixel */ + SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */ + SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */ + SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */ + SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */ + )) + using DeviceFullView = DeviceFullViewTemplate<>; /* Device pixel view: this is a second generation view (view from view) */ - GENERATE_SOA_CONST_VIEW(DevicePixelConstViewTemplate, - /* We get out data from the DeviceFullView */ - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(DeviceFullView, deviceFullView) - ), - /* These are consumed by downstream device code */ - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(deviceFullView, xx), /* local coordinates of each pixel */ - SOA_VIEW_VALUE(deviceFullView, yy), /* */ - SOA_VIEW_VALUE(deviceFullView, moduleInd), /* module id of each pixel */ - SOA_VIEW_VALUE(deviceFullView, adc), /* ADC of each pixel */ - SOA_VIEW_VALUE(deviceFullView, clus) /* cluster id of each pixel */ - ) - ) - + GENERATE_SOA_CONST_VIEW( + DevicePixelConstViewTemplate, + /* We get out data from the DeviceFullView */ + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceFullView, deviceFullView)), + /* These are consumed by downstream device code */ + SOA_VIEW_VALUE_LIST( + SOA_VIEW_VALUE(deviceFullView, xx), /* local coordinates of each pixel */ + SOA_VIEW_VALUE(deviceFullView, yy), /* */ + SOA_VIEW_VALUE(deviceFullView, moduleInd), /* module id of each pixel */ + SOA_VIEW_VALUE(deviceFullView, adc), /* ADC of each pixel */ + SOA_VIEW_VALUE(deviceFullView, clus) /* cluster id of each pixel */ + )) + using DevicePixelConstView = DevicePixelConstViewTemplate<>; explicit SiPixelDigisCUDA(); @@ -117,27 +113,28 @@ class SiPixelDigisCUDA { class HostStore { friend SiPixelDigisCUDA; + public: HostStore(); const SiPixelDigisCUDA::HostDeviceView view() { return hostView_; } void reset(); + private: HostStore(size_t maxFedWords, cudaStream_t stream); cms::cuda::host::unique_ptr data_h; HostDeviceLayout hostLayout_; HostDeviceView hostView_; - }; HostStore dataToHostAsync(cudaStream_t stream) const; - // Special copy for validation - cms::cuda::host::unique_ptr adcToHostAsync(cudaStream_t stream) const; + // Special copy for validation + cms::cuda::host::unique_ptr adcToHostAsync(cudaStream_t stream) const; - const DevicePixelConstView& pixelConstView() const { return devicePixelConstView_; } + const DevicePixelConstView &pixelConstView() const { return devicePixelConstView_; } private: // These are consumed by downstream device code - cms::cuda::device::unique_ptr data_d; // Single SoA storage + cms::cuda::device::unique_ptr data_d; // Single SoA storage DeviceOnlyLayout deviceOnlyLayout_d; HostDeviceLayout hostDeviceLayout_d; DeviceFullView deviceFullView_; diff --git a/src/cudadev/CUDADataFormats/TrackingRecHit2DHeterogeneous.h b/src/cudadev/CUDADataFormats/TrackingRecHit2DHeterogeneous.h index 26a4ca75a..5294328a8 100644 --- a/src/cudadev/CUDADataFormats/TrackingRecHit2DHeterogeneous.h +++ b/src/cudadev/CUDADataFormats/TrackingRecHit2DHeterogeneous.h @@ -40,21 +40,21 @@ class TrackingRecHit2DHeterogeneous { // Transfer the local and global coordinates, charge and size TrackingRecHit2DHostSOAStore hitsToHostAsync(cudaStream_t stream) const; - + // apparently unused //cms::cuda::host::unique_ptr hitsModuleStartToHostAsync(cudaStream_t stream) const; private: static_assert(sizeof(uint32_t) == sizeof(float)); // just stating the obvious - + unique_ptr m_PhiBinnerStore; //! unique_ptr m_AverageGeometryStore; //! unique_ptr m_store; //! uint32_t m_nHits; - - unique_ptr m_hitsSupportLayerStartStore; //! + + unique_ptr m_hitsSupportLayerStartStore; //! uint32_t const* m_hitsModuleStart; // needed for legacy, this is on GPU! @@ -101,41 +101,41 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous(uint32_t nH //m_store16 = Traits::template make_device_unique(nHits * n16, stream); //m_store32 = // Traits::template make_device_unique(nHits * n32 + phase1PixelTopology::numberOfLayers + 1, stream); - // We need to store all SoA rows for TrackingRecHit2DSOAView::HitsView(nHits) + + // We need to store all SoA rows for TrackingRecHit2DSOAView::HitsView(nHits) + // (phase1PixelTopology::numberOfLayers + 1) TrackingRecHit2DSOAView::PhiBinner::index_type. - // As mentioned above, alignment is not important, yet we want to have 32 bits + // As mentioned above, alignment is not important, yet we want to have 32 bits // (TrackingRecHit2DSOAView::PhiBinner::index_type exactly) alignement for the second part. // In order to simplify code, we align all to the minimum necessary size (sizeof(TrackingRecHit2DSOAStore::PhiBinner::index_type)). { // Simplify a bit following computations const size_t phiBinnerByteSize = - (phase1PixelTopology::numberOfLayers + 1) * sizeof (TrackingRecHit2DSOAStore::PhiBinner::index_type); + (phase1PixelTopology::numberOfLayers + 1) * sizeof(TrackingRecHit2DSOAStore::PhiBinner::index_type); // Allocate the buffer - m_hitsSupportLayerStartStore = Traits::template make_device_unique ( - TrackingRecHit2DSOAStore::HitsLayout::computeDataSize(m_nHits) + - TrackingRecHit2DSOAStore::SupportObjectsLayout::computeDataSize(m_nHits) + - phiBinnerByteSize, - stream); + m_hitsSupportLayerStartStore = Traits::template make_device_unique( + TrackingRecHit2DSOAStore::HitsLayout::computeDataSize(m_nHits) + + TrackingRecHit2DSOAStore::SupportObjectsLayout::computeDataSize(m_nHits) + phiBinnerByteSize, + stream); // Split the buffer in stores and array store->m_hitsLayout = TrackingRecHit2DSOAStore::HitsLayout(m_hitsSupportLayerStartStore.get(), nHits); - store->m_supportObjectsLayout = TrackingRecHit2DSOAStore::SupportObjectsLayout(store->m_hitsLayout.soaMetadata().nextByte(), nHits); - m_hitsLayerStart = store->m_hitsLayerStart = reinterpret_cast (store->m_supportObjectsLayout.soaMetadata().nextByte()); + store->m_supportObjectsLayout = + TrackingRecHit2DSOAStore::SupportObjectsLayout(store->m_hitsLayout.soaMetadata().nextByte(), nHits); + m_hitsLayerStart = store->m_hitsLayerStart = + reinterpret_cast(store->m_supportObjectsLayout.soaMetadata().nextByte()); // Record additional references - store->m_hitsAndSupportView = TrackingRecHit2DSOAStore::HitsAndSupportView( - store->m_hitsLayout, - store->m_supportObjectsLayout - ); + store->m_hitsAndSupportView = + TrackingRecHit2DSOAStore::HitsAndSupportView(store->m_hitsLayout, store->m_supportObjectsLayout); m_phiBinnerStorage = store->m_phiBinnerStorage = store->m_hitsAndSupportView.phiBinnerStorage(); m_iphi = store->m_hitsAndSupportView.iphi(); } m_PhiBinnerStore = Traits::template make_device_unique(stream); static_assert(sizeof(TrackingRecHit2DSOAStore::hindex_type) == sizeof(float)); - static_assert(sizeof(TrackingRecHit2DSOAStore::hindex_type) == sizeof(TrackingRecHit2DSOAStore::PhiBinner::index_type)); + static_assert(sizeof(TrackingRecHit2DSOAStore::hindex_type) == + sizeof(TrackingRecHit2DSOAStore::PhiBinner::index_type)); // copy all the pointers m_phiBinner = store->m_phiBinner = m_PhiBinnerStore.get(); - + // transfer view if constexpr (std::is_same::value) { cms::cuda::copyAsync(m_store, store, stream); diff --git a/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.cc b/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.cc index 530daa600..6e2338f6c 100644 --- a/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.cc +++ b/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.cc @@ -1,16 +1,15 @@ #include "CUDADataFormats/TrackingRecHit2DHostSOAStore.h" -TrackingRecHit2DHostSOAStore::TrackingRecHit2DHostSOAStore(): - hitsLayout_(hits_h.get(), 0 /* size */, 1 /* byte alignement */) -{} +TrackingRecHit2DHostSOAStore::TrackingRecHit2DHostSOAStore() + : hitsLayout_(hits_h.get(), 0 /* size */, 1 /* byte alignement */) {} void TrackingRecHit2DHostSOAStore::reset() { hits_h.reset(); hitsLayout_ = TrackingRecHit2DSOAStore::HitsLayout(); } -TrackingRecHit2DHostSOAStore::TrackingRecHit2DHostSOAStore(size_t size, cudaStream_t stream): - hits_h(cms::cuda::make_host_unique(TrackingRecHit2DSOAStore::HitsLayout::computeDataSize(size), stream)), - hitsLayout_(hits_h.get(), size), - hitsView_(hitsLayout_) -{} +TrackingRecHit2DHostSOAStore::TrackingRecHit2DHostSOAStore(size_t size, cudaStream_t stream) + : hits_h(cms::cuda::make_host_unique(TrackingRecHit2DSOAStore::HitsLayout::computeDataSize(size), + stream)), + hitsLayout_(hits_h.get(), size), + hitsView_(hitsLayout_) {} diff --git a/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.h b/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.h index f2b34e9bc..e587932d4 100644 --- a/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.h +++ b/src/cudadev/CUDADataFormats/TrackingRecHit2DHostSOAStore.h @@ -6,16 +6,20 @@ #include "CUDACore/host_unique_ptr.h" template -class TrackingRecHit2DHeterogeneous; +class TrackingRecHit2DHeterogeneous; struct TrackingRecHit2DHostSOAStore { template friend class TrackingRecHit2DHeterogeneous; + public: TrackingRecHit2DHostSOAStore(); void reset(); - __device__ __forceinline__ const auto operator[](size_t i) const { return hitsView_[i]; } - __device__ __forceinline__ size_t size() { return /* TODO: move to view when view will embed size */hitsLayout_.soaMetadata().size(); } + __device__ __forceinline__ const auto operator[](size_t i) const { return hitsView_[i]; } + __device__ __forceinline__ size_t size() { + return /* TODO: move to view when view will embed size */ hitsLayout_.soaMetadata().size(); + } + private: TrackingRecHit2DHostSOAStore(size_t size, cudaStream_t stream); cms::cuda::host::unique_ptr hits_h; @@ -23,5 +27,4 @@ struct TrackingRecHit2DHostSOAStore { TrackingRecHit2DSOAStore::HitsView hitsView_; }; - -#endif // ndef CUDADataFormats_TrackingRecHit_interface_TrackingRecHit2DSOAHostStore_h \ No newline at end of file +#endif // ndef CUDADataFormats_TrackingRecHit_interface_TrackingRecHit2DSOAHostStore_h \ No newline at end of file diff --git a/src/cudadev/CUDADataFormats/TrackingRecHit2DSOAView.h b/src/cudadev/CUDADataFormats/TrackingRecHit2DSOAView.h index c983dc6e5..7ee03c1ea 100644 --- a/src/cudadev/CUDADataFormats/TrackingRecHit2DSOAView.h +++ b/src/cudadev/CUDADataFormats/TrackingRecHit2DSOAView.h @@ -24,115 +24,103 @@ class TrackingRecHit2DSOAStore { template friend class TrackingRecHit2DHeterogeneous; - + __device__ __forceinline__ uint32_t nHits() const { return m_nHits; } // Our arrays do not require specific alignment as access will not be coalesced in the current implementation // Sill, we need the 32 bits integers to be aligned, so we simply declare the SoA with the 32 bits fields first - // and the 16 bits behind (as they have a looser alignment requirement. Then the SoA can be create with a byte + // and the 16 bits behind (as they have a looser alignment requirement. Then the SoA can be create with a byte // alignment of 1) GENERATE_SOA_LAYOUT(HitsLayoutTemplate, - // 32 bits section - // local coord - SOA_COLUMN(float, xLocal), - SOA_COLUMN(float, yLocal), - SOA_COLUMN(float, xerrLocal), - SOA_COLUMN(float, yerrLocal), - - // global coord - SOA_COLUMN(float, xGlobal), - SOA_COLUMN(float, yGlobal), - SOA_COLUMN(float, zGlobal), - SOA_COLUMN(float, rGlobal), - // global coordinates continue in the 16 bits section - - // cluster properties - SOA_COLUMN(int32_t, charge), - - // 16 bits section (and cluster properties immediately continued) - SOA_COLUMN(int16_t, clusterSizeX), - SOA_COLUMN(int16_t, clusterSizeY) - ) - + // 32 bits section + // local coord + SOA_COLUMN(float, xLocal), + SOA_COLUMN(float, yLocal), + SOA_COLUMN(float, xerrLocal), + SOA_COLUMN(float, yerrLocal), + + // global coord + SOA_COLUMN(float, xGlobal), + SOA_COLUMN(float, yGlobal), + SOA_COLUMN(float, zGlobal), + SOA_COLUMN(float, rGlobal), + // global coordinates continue in the 16 bits section + + // cluster properties + SOA_COLUMN(int32_t, charge), + + // 16 bits section (and cluster properties immediately continued) + SOA_COLUMN(int16_t, clusterSizeX), + SOA_COLUMN(int16_t, clusterSizeY)) + // The hits layout does not use default alignment but a more relaxed one. using HitsLayout = HitsLayoutTemplate; - + GENERATE_SOA_VIEW(HitsViewTemplate, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(HitsLayout, hitsLayout) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(hitsLayout, xLocal), - SOA_VIEW_VALUE(hitsLayout, yLocal), - SOA_VIEW_VALUE(hitsLayout, xerrLocal), - SOA_VIEW_VALUE(hitsLayout, yerrLocal), - - SOA_VIEW_VALUE(hitsLayout, xGlobal), - SOA_VIEW_VALUE(hitsLayout, yGlobal), - SOA_VIEW_VALUE(hitsLayout, zGlobal), - SOA_VIEW_VALUE(hitsLayout, rGlobal), - - SOA_VIEW_VALUE(hitsLayout, charge), - SOA_VIEW_VALUE(hitsLayout, clusterSizeX), - SOA_VIEW_VALUE(hitsLayout, clusterSizeY) - ) - ) - + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(HitsLayout, hitsLayout)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(hitsLayout, xLocal), + SOA_VIEW_VALUE(hitsLayout, yLocal), + SOA_VIEW_VALUE(hitsLayout, xerrLocal), + SOA_VIEW_VALUE(hitsLayout, yerrLocal), + + SOA_VIEW_VALUE(hitsLayout, xGlobal), + SOA_VIEW_VALUE(hitsLayout, yGlobal), + SOA_VIEW_VALUE(hitsLayout, zGlobal), + SOA_VIEW_VALUE(hitsLayout, rGlobal), + + SOA_VIEW_VALUE(hitsLayout, charge), + SOA_VIEW_VALUE(hitsLayout, clusterSizeX), + SOA_VIEW_VALUE(hitsLayout, clusterSizeY))) + using HitsView = HitsViewTemplate<>; - + GENERATE_SOA_LAYOUT(SupportObjectsLayoutTemplate, - // This is the end of the data which is transferred to host. The following columns are supporting - // objects, not transmitted - - // Supporting data (32 bits aligned) - SOA_COLUMN(TrackingRecHit2DSOAStore::PhiBinner::index_type, phiBinnerStorage), - - // global coordinates (not transmitted) - SOA_COLUMN(int16_t, iphi), - - // cluster properties (not transmitted) - SOA_COLUMN(uint16_t, detectorIndex) - ); - + // This is the end of the data which is transferred to host. The following columns are supporting + // objects, not transmitted + + // Supporting data (32 bits aligned) + SOA_COLUMN(TrackingRecHit2DSOAStore::PhiBinner::index_type, phiBinnerStorage), + + // global coordinates (not transmitted) + SOA_COLUMN(int16_t, iphi), + + // cluster properties (not transmitted) + SOA_COLUMN(uint16_t, detectorIndex)) + // The support objects layouts also not use default alignment but a more relaxed one. using SupportObjectsLayout = SupportObjectsLayoutTemplate; - + GENERATE_SOA_VIEW(HitsAndSupportViewTemplate, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(HitsLayout, hitsLayout), - SOA_VIEW_LAYOUT(SupportObjectsLayout, supportObjectsLayout) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(hitsLayout, xLocal), - SOA_VIEW_VALUE(hitsLayout, yLocal), - SOA_VIEW_VALUE(hitsLayout, xerrLocal), - SOA_VIEW_VALUE(hitsLayout, yerrLocal), - - SOA_VIEW_VALUE(hitsLayout, xGlobal), - SOA_VIEW_VALUE(hitsLayout, yGlobal), - SOA_VIEW_VALUE(hitsLayout, zGlobal), - SOA_VIEW_VALUE(hitsLayout, rGlobal), - - SOA_VIEW_VALUE(hitsLayout, charge), - SOA_VIEW_VALUE(hitsLayout, clusterSizeX), - SOA_VIEW_VALUE(hitsLayout, clusterSizeY), - - SOA_VIEW_VALUE(supportObjectsLayout, phiBinnerStorage), - SOA_VIEW_VALUE(supportObjectsLayout, iphi), - SOA_VIEW_VALUE(supportObjectsLayout, detectorIndex) - ) - ); - + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(HitsLayout, hitsLayout), + SOA_VIEW_LAYOUT(SupportObjectsLayout, supportObjectsLayout)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(hitsLayout, xLocal), + SOA_VIEW_VALUE(hitsLayout, yLocal), + SOA_VIEW_VALUE(hitsLayout, xerrLocal), + SOA_VIEW_VALUE(hitsLayout, yerrLocal), + + SOA_VIEW_VALUE(hitsLayout, xGlobal), + SOA_VIEW_VALUE(hitsLayout, yGlobal), + SOA_VIEW_VALUE(hitsLayout, zGlobal), + SOA_VIEW_VALUE(hitsLayout, rGlobal), + + SOA_VIEW_VALUE(hitsLayout, charge), + SOA_VIEW_VALUE(hitsLayout, clusterSizeX), + SOA_VIEW_VALUE(hitsLayout, clusterSizeY), + + SOA_VIEW_VALUE(supportObjectsLayout, phiBinnerStorage), + SOA_VIEW_VALUE(supportObjectsLayout, iphi), + SOA_VIEW_VALUE(supportObjectsLayout, detectorIndex))) + using HitsAndSupportView = HitsAndSupportViewTemplate; - + // Shortcut operator saving the explicit calls to view in usage. - __device__ __forceinline__ HitsAndSupportView::element operator[] (size_t index) { - return m_hitsAndSupportView[index]; + __device__ __forceinline__ HitsAndSupportView::element operator[](size_t index) { + return m_hitsAndSupportView[index]; } - __device__ __forceinline__ HitsAndSupportView::const_element operator[] (size_t index) const { + __device__ __forceinline__ HitsAndSupportView::const_element operator[](size_t index) const { return m_hitsAndSupportView[index]; } - + __device__ __forceinline__ pixelCPEforGPU::ParamsOnGPU const& cpeParams() const { return *m_cpeParams; } __device__ __forceinline__ uint32_t hitsModuleStart(int i) const { return __ldg(m_hitsModuleStart + i); } @@ -153,7 +141,7 @@ class TrackingRecHit2DSOAStore { SupportObjectsLayout m_supportObjectsLayout; // Global view simplifying usage HitsAndSupportView m_hitsAndSupportView; - + // individually defined supporting objects // m_averageGeometry is corrected for beam spot, not sure where to host it otherwise AverageGeometry* m_averageGeometry; // owned by TrackingRecHit2DHeterogeneous diff --git a/src/cudadev/CondFormats/SiPixelROCsStatusAndMapping.h b/src/cudadev/CondFormats/SiPixelROCsStatusAndMapping.h index 711675a5d..af60caf9d 100644 --- a/src/cudadev/CondFormats/SiPixelROCsStatusAndMapping.h +++ b/src/cudadev/CondFormats/SiPixelROCsStatusAndMapping.h @@ -26,34 +26,30 @@ struct SiPixelROCsStatusAndMapping { }; GENERATE_SOA_LAYOUT(SiPixelROCsStatusAndMappingLayoutTemplate, - SOA_COLUMN(unsigned int, fed), - SOA_COLUMN(unsigned int, link), - SOA_COLUMN(unsigned int, roc), - SOA_COLUMN(unsigned int, rawId), - SOA_COLUMN(unsigned int, rocInDet), - SOA_COLUMN(unsigned int, moduleId), - SOA_COLUMN(unsigned char, badRocs), - SOA_SCALAR(unsigned int, size) -) + SOA_COLUMN(unsigned int, fed), + SOA_COLUMN(unsigned int, link), + SOA_COLUMN(unsigned int, roc), + SOA_COLUMN(unsigned int, rawId), + SOA_COLUMN(unsigned int, rocInDet), + SOA_COLUMN(unsigned int, moduleId), + SOA_COLUMN(unsigned char, badRocs), + SOA_SCALAR(unsigned int, size)) using SiPixelROCsStatusAndMappingLayout = SiPixelROCsStatusAndMappingLayoutTemplate<>; GENERATE_SOA_CONST_VIEW(SiPixelROCsStatusAndMappingConstViewTemplate, - SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(SiPixelROCsStatusAndMappingLayout, mappingLayout)), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(mappingLayout, fed), - SOA_VIEW_VALUE(mappingLayout, link), - SOA_VIEW_VALUE(mappingLayout, roc), - SOA_VIEW_VALUE(mappingLayout, rawId), - SOA_VIEW_VALUE(mappingLayout, rocInDet), - SOA_VIEW_VALUE(mappingLayout, moduleId), - SOA_VIEW_VALUE(mappingLayout, badRocs), - SOA_VIEW_VALUE(mappingLayout, size) - ) -) + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(SiPixelROCsStatusAndMappingLayout, mappingLayout)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(mappingLayout, fed), + SOA_VIEW_VALUE(mappingLayout, link), + SOA_VIEW_VALUE(mappingLayout, roc), + SOA_VIEW_VALUE(mappingLayout, rawId), + SOA_VIEW_VALUE(mappingLayout, rocInDet), + SOA_VIEW_VALUE(mappingLayout, moduleId), + SOA_VIEW_VALUE(mappingLayout, badRocs), + SOA_VIEW_VALUE(mappingLayout, size))) // Slightly more complex than using, but allows forward declarations. -struct SiPixelROCsStatusAndMappingConstView: public SiPixelROCsStatusAndMappingConstViewTemplate<> { +struct SiPixelROCsStatusAndMappingConstView : public SiPixelROCsStatusAndMappingConstViewTemplate<> { using SiPixelROCsStatusAndMappingConstViewTemplate<>::SiPixelROCsStatusAndMappingConstViewTemplate; }; diff --git a/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.cc b/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.cc index 2b7cc79cd..c09ed1852 100644 --- a/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.cc +++ b/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.cc @@ -24,22 +24,22 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelRO std::copy(modToUnp.begin(), modToUnp.end(), modToUnpDefault.begin()); } -SiPixelROCsStatusAndMappingConstView SiPixelROCsStatusAndMappingWrapper::getGPUProductAsync(cudaStream_t cudaStream) const { - const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, - [this](GPUData& data, cudaStream_t stream) { - // allocate - data.allocate(stream); - // transfer - cms::cuda::copyAsync(data.cablingMapDevice, this->cablingMapHost, stream); - } - ); +SiPixelROCsStatusAndMappingConstView SiPixelROCsStatusAndMappingWrapper::getGPUProductAsync( + cudaStream_t cudaStream) const { + const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, [this](GPUData& data, cudaStream_t stream) { + // allocate + data.allocate(stream); + // transfer + cms::cuda::copyAsync(data.cablingMapDevice, this->cablingMapHost, stream); + }); return data.cablingMapDeviceView; } const unsigned char* SiPixelROCsStatusAndMappingWrapper::getModToUnpAllAsync(cudaStream_t cudaStream) const { const auto& data = modToUnp_.dataForCurrentDeviceAsync(cudaStream, [this](ModulesToUnpack& data, cudaStream_t stream) { - data.modToUnpDefault = cms::cuda::make_device_unique(pixelgpudetails::MAX_SIZE_BYTE_BOOL, stream); + data.modToUnpDefault = + cms::cuda::make_device_unique(pixelgpudetails::MAX_SIZE_BYTE_BOOL, stream); cudaCheck(cudaMemcpyAsync(data.modToUnpDefault.get(), this->modToUnpDefault.data(), this->modToUnpDefault.size() * sizeof(unsigned char), diff --git a/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.h b/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.h index 61f0f5b3c..120ce4d29 100644 --- a/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.h +++ b/src/cudadev/CondFormats/SiPixelROCsStatusAndMappingWrapper.h @@ -15,7 +15,7 @@ class SiPixelROCsStatusAndMappingWrapper { public: /* This is using a layout as the size is needed. TODO: use views when views start embedding size. */ explicit SiPixelROCsStatusAndMappingWrapper(SiPixelROCsStatusAndMapping const &cablingMap, - std::vector modToUnp); + std::vector modToUnp); bool hasQuality() const { return hasQuality_; } @@ -35,27 +35,26 @@ class SiPixelROCsStatusAndMappingWrapper { void allocate(cudaStream_t stream) { cablingMapDevice = cms::cuda::make_device_unique(stream); // Populate the view with individual column pointers - auto & cmd = *cablingMapDevice; + auto &cmd = *cablingMapDevice; cablingMapDeviceView = SiPixelROCsStatusAndMappingConstView( - pixelgpudetails::MAX_SIZE, - cmd.fed, // Those are array pointers (in device, but we won't dereference them here). - cmd.link, - cmd.roc, - cmd.rawId, - cmd.rocInDet, - cmd.moduleId, - cmd.badRocs, - &cmd.size // This is a scalar, we need the address-of operator + pixelgpudetails::MAX_SIZE, + cmd.fed, // Those are array pointers (in device, but we won't dereference them here). + cmd.link, + cmd.roc, + cmd.rawId, + cmd.rocInDet, + cmd.moduleId, + cmd.badRocs, + &cmd.size // This is a scalar, we need the address-of operator ); } cms::cuda::device::unique_ptr cablingMapDevice; - SiPixelROCsStatusAndMappingConstView cablingMapDeviceView; // map struct in GPU - + SiPixelROCsStatusAndMappingConstView cablingMapDeviceView; // map struct in GPU }; cms::cuda::ESProduct gpuData_; struct ModulesToUnpack { - cms::cuda::device::unique_ptr modToUnpDefault; // pointer to GPU + cms::cuda::device::unique_ptr modToUnpDefault; // pointer to GPU }; cms::cuda::ESProduct modToUnp_; }; diff --git a/src/cudadev/DataFormats/SoACommon.h b/src/cudadev/DataFormats/SoACommon.h index cd3ac76cb..28727319e 100644 --- a/src/cudadev/DataFormats/SoACommon.h +++ b/src/cudadev/DataFormats/SoACommon.h @@ -25,181 +25,184 @@ // Exception throwing (or willful crash in kernels) #if defined(__CUDACC__) && defined(__CUDA_ARCH__) -#define SOA_THROW_OUT_OF_RANGE(A) { printf (A); *((char *)nullptr) = 0; } +#define SOA_THROW_OUT_OF_RANGE(A) \ + { \ + printf(A); \ + *((char*)nullptr) = 0; \ + } #else -#define SOA_THROW_OUT_OF_RANGE(A) { throw std::out_of_range(A); } +#define SOA_THROW_OUT_OF_RANGE(A) \ + { throw std::out_of_range(A); } #endif // compile-time sized SoA namespace cms::soa { -enum class RestrictQualify : bool { Enabled, Disabled, Default = Disabled }; - -enum class RangeChecking: bool { Enabled, Disabled, Default = Disabled }; - -template -struct add_restrict {}; - -template -struct add_restrict { - typedef T Value; - typedef T * __restrict__ Pointer; - typedef T & __restrict__ Reference; - typedef const T ConstValue; - typedef const T * __restrict__ PointerToConst; - typedef const T & __restrict__ ReferenceToConst; -}; - -template -struct add_restrict { - typedef T Value; - typedef T * Pointer; - typedef T & Reference; - typedef const T ConstValue; - typedef const T * PointerToConst; - typedef const T & ReferenceToConst; -}; - -// Helper template managing the value within it column -// The optional compile time alignment parameter enables informing the -// compiler of alignment (enforced by caller). -template -class SoAValue { -public: - typedef add_restrict Restr; - typedef typename Restr::Value Val; - typedef typename Restr::Pointer Ptr; - typedef typename Restr::Reference Ref; - typedef typename Restr::PointerToConst PtrToConst; - typedef typename Restr::ReferenceToConst RefToConst; - SOA_HOST_DEVICE_INLINE SoAValue(size_t i, T* col) : idx_(i), col_(col) {} - /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ - SOA_HOST_DEVICE_INLINE Ref operator()() { - // Ptr type will add the restrict qualifyer if needed - Ptr col = alignedCol(); - return col[idx_]; - } - SOA_HOST_DEVICE_INLINE RefToConst operator()() const { - // PtrToConst type will add the restrict qualifyer if needed - PtrToConst col = alignedCol(); - return col[idx_]; - } - SOA_HOST_DEVICE_INLINE Ptr operator&() { return &alignedCol()[idx_]; } - SOA_HOST_DEVICE_INLINE PtrToConst operator&() const { return &alignedCol()[idx_]; } - template - SOA_HOST_DEVICE_INLINE Ref operator=(const T2& v) { - return alignedCol()[idx_] = v; - } - typedef Val valueType; - static constexpr auto valueSize = sizeof(T); + enum class RestrictQualify : bool { Enabled, Disabled, Default = Disabled }; + + enum class RangeChecking : bool { Enabled, Disabled, Default = Disabled }; + + template + struct add_restrict {}; + + template + struct add_restrict { + typedef T Value; + typedef T* __restrict__ Pointer; + typedef T& __restrict__ Reference; + typedef const T ConstValue; + typedef const T* __restrict__ PointerToConst; + typedef const T& __restrict__ ReferenceToConst; + }; + + template + struct add_restrict { + typedef T Value; + typedef T* Pointer; + typedef T& Reference; + typedef const T ConstValue; + typedef const T* PointerToConst; + typedef const T& ReferenceToConst; + }; -private: - SOA_HOST_DEVICE_INLINE Ptr alignedCol() const { - if constexpr (ALIGNMENT) { - return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + // Helper template managing the value within it column + // The optional compile time alignment parameter enables informing the + // compiler of alignment (enforced by caller). + template + class SoAValue { + public: + typedef add_restrict Restr; + typedef typename Restr::Value Val; + typedef typename Restr::Pointer Ptr; + typedef typename Restr::Reference Ref; + typedef typename Restr::PointerToConst PtrToConst; + typedef typename Restr::ReferenceToConst RefToConst; + SOA_HOST_DEVICE_INLINE SoAValue(size_t i, T* col) : idx_(i), col_(col) {} + /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ + SOA_HOST_DEVICE_INLINE Ref operator()() { + // Ptr type will add the restrict qualifyer if needed + Ptr col = alignedCol(); + return col[idx_]; } - return reinterpret_cast(col_); - } - size_t idx_; - T* col_; -}; - -// Helper template managing the value within it column -template -class SoAConstValue { -public: - typedef add_restrict Restr; - typedef typename Restr::Value Val; - typedef typename Restr::Pointer Ptr; - typedef typename Restr::Reference Ref; - typedef typename Restr::PointerToConst PtrToConst; - typedef typename Restr::ReferenceToConst RefToConst; - SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, const T* col) : idx_(i), col_(col) {} - /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ - SOA_HOST_DEVICE_INLINE RefToConst operator()() const { - // Ptr type will add the restrict qualifyer if needed - PtrToConst col = alignedCol(); - return col[idx_]; - } - SOA_HOST_DEVICE_INLINE const T* operator&() const { return &alignedCol()[idx_]; } - typedef T valueType; - static constexpr auto valueSize = sizeof(T); - -private: - SOA_HOST_DEVICE_INLINE PtrToConst alignedCol() const { - if constexpr (ALIGNMENT) { - return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + SOA_HOST_DEVICE_INLINE RefToConst operator()() const { + // PtrToConst type will add the restrict qualifyer if needed + PtrToConst col = alignedCol(); + return col[idx_]; } - return reinterpret_cast(col_) ; - } - size_t idx_; - const T* col_; -}; - -// Helper template managing the value within it column -// TODO Create a const variant to avoid leaking mutable access. -template -class SoAEigenValue { -public: - typedef C Type; - typedef Eigen::Map> MapType; - typedef Eigen::Map> CMapType; - SOA_HOST_DEVICE_INLINE SoAEigenValue(size_t i, typename C::Scalar* col, size_t stride) - : val_(col + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), - crCol_(col), - cVal_(crCol_ + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), - stride_(stride) {} - SOA_HOST_DEVICE_INLINE MapType& operator()() { return val_; } - SOA_HOST_DEVICE_INLINE const CMapType& operator()() const { return cVal_; } - SOA_HOST_DEVICE_INLINE operator C() { return val_; } - SOA_HOST_DEVICE_INLINE operator const C() const { return cVal_; } - SOA_HOST_DEVICE_INLINE C* operator&() { return &val_; } - SOA_HOST_DEVICE_INLINE const C* operator&() const { return &cVal_; } - template - SOA_HOST_DEVICE_INLINE MapType& operator=(const C2& v) { - return val_ = v; - } - typedef typename C::Scalar ValueType; - static constexpr auto valueSize = sizeof(C::Scalar); - SOA_HOST_DEVICE_INLINE size_t stride() { return stride_; } - -private: - MapType val_; - const typename C::Scalar* __restrict__ crCol_; - CMapType cVal_; - size_t stride_; -}; - -// Helper template to avoid commas in macro -template -struct EigenConstMapMaker { - typedef Eigen::Map> Type; - class DataHolder { + SOA_HOST_DEVICE_INLINE Ptr operator&() { return &alignedCol()[idx_]; } + SOA_HOST_DEVICE_INLINE PtrToConst operator&() const { return &alignedCol()[idx_]; } + template + SOA_HOST_DEVICE_INLINE Ref operator=(const T2& v) { + return alignedCol()[idx_] = v; + } + typedef Val valueType; + static constexpr auto valueSize = sizeof(T); + + private: + SOA_HOST_DEVICE_INLINE Ptr alignedCol() const { + if constexpr (ALIGNMENT) { + return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + } + return reinterpret_cast(col_); + } + size_t idx_; + T* col_; + }; + + // Helper template managing the value within it column + template + class SoAConstValue { public: - DataHolder(const typename C::Scalar* data) : data_(data) {} - EigenConstMapMaker::Type withStride(size_t stride) { - return EigenConstMapMaker::Type( - data_, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)); + typedef add_restrict Restr; + typedef typename Restr::Value Val; + typedef typename Restr::Pointer Ptr; + typedef typename Restr::Reference Ref; + typedef typename Restr::PointerToConst PtrToConst; + typedef typename Restr::ReferenceToConst RefToConst; + SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, const T* col) : idx_(i), col_(col) {} + /* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */ + SOA_HOST_DEVICE_INLINE RefToConst operator()() const { + // Ptr type will add the restrict qualifyer if needed + PtrToConst col = alignedCol(); + return col[idx_]; } + SOA_HOST_DEVICE_INLINE const T* operator&() const { return &alignedCol()[idx_]; } + typedef T valueType; + static constexpr auto valueSize = sizeof(T); private: - const typename C::Scalar* const data_; + SOA_HOST_DEVICE_INLINE PtrToConst alignedCol() const { + if constexpr (ALIGNMENT) { + return reinterpret_cast(__builtin_assume_aligned(col_, ALIGNMENT)); + } + return reinterpret_cast(col_); + } + size_t idx_; + const T* col_; }; - static DataHolder withData(const typename C::Scalar* data) { return DataHolder(data); } -}; -// Helper function to compute aligned size -inline size_t alignSize(size_t size, size_t alignment = 128) { - if (size) - return ((size - 1) / alignment + 1) * alignment; - else - return 0; -} + // Helper template managing the value within it column + // TODO Create a const variant to avoid leaking mutable access. + template + class SoAEigenValue { + public: + typedef C Type; + typedef Eigen::Map> MapType; + typedef Eigen::Map> CMapType; + SOA_HOST_DEVICE_INLINE SoAEigenValue(size_t i, typename C::Scalar* col, size_t stride) + : val_(col + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), + crCol_(col), + cVal_(crCol_ + i, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)), + stride_(stride) {} + SOA_HOST_DEVICE_INLINE MapType& operator()() { return val_; } + SOA_HOST_DEVICE_INLINE const CMapType& operator()() const { return cVal_; } + SOA_HOST_DEVICE_INLINE operator C() { return val_; } + SOA_HOST_DEVICE_INLINE operator const C() const { return cVal_; } + SOA_HOST_DEVICE_INLINE C* operator&() { return &val_; } + SOA_HOST_DEVICE_INLINE const C* operator&() const { return &cVal_; } + template + SOA_HOST_DEVICE_INLINE MapType& operator=(const C2& v) { + return val_ = v; + } + typedef typename C::Scalar ValueType; + static constexpr auto valueSize = sizeof(C::Scalar); + SOA_HOST_DEVICE_INLINE size_t stride() { return stride_; } -} // namespace cms::soa + private: + MapType val_; + const typename C::Scalar* __restrict__ crCol_; + CMapType cVal_; + size_t stride_; + }; + + // Helper template to avoid commas in macro + template + struct EigenConstMapMaker { + typedef Eigen::Map> Type; + class DataHolder { + public: + DataHolder(const typename C::Scalar* data) : data_(data) {} + EigenConstMapMaker::Type withStride(size_t stride) { + return EigenConstMapMaker::Type( + data_, C::RowsAtCompileTime, C::ColsAtCompileTime, Eigen::InnerStride(stride)); + } + + private: + const typename C::Scalar* const data_; + }; + static DataHolder withData(const typename C::Scalar* data) { return DataHolder(data); } + }; + + // Helper function to compute aligned size + inline size_t alignSize(size_t size, size_t alignment = 128) { + if (size) + return ((size - 1) / alignment + 1) * alignment; + else + return 0; + } + +} // namespace cms::soa /* declare "scalars" (one value shared across the whole SoA) and "columns" (one value per element) */ #define _VALUE_TYPE_SCALAR 0 @@ -208,9 +211,13 @@ inline size_t alignSize(size_t size, size_t alignment = 128) { namespace cms::soa { -enum class SoAColumnType { scalar = _VALUE_TYPE_SCALAR, column = _VALUE_TYPE_COLUMN, eigen = _VALUE_TYPE_EIGEN_COLUMN }; + enum class SoAColumnType { + scalar = _VALUE_TYPE_SCALAR, + column = _VALUE_TYPE_COLUMN, + eigen = _VALUE_TYPE_EIGEN_COLUMN + }; -} // namespace cms::soa +} // namespace cms::soa #define SOA_SCALAR(TYPE, NAME) (_VALUE_TYPE_SCALAR, TYPE, NAME) #define SOA_COLUMN(TYPE, NAME) (_VALUE_TYPE_COLUMN, TYPE, NAME) @@ -234,83 +241,90 @@ enum class SoAColumnType { scalar = _VALUE_TYPE_SCALAR, column = _VALUE_TYPE_COL namespace cms::soa { -/* Column accessors: templates implementing the global accesors (soa::x() and soa::x(index) */ -enum class SoAAccessType: bool { mutableAccess, constAccess }; - -template -struct SoAColumnAccessorsImpl {}; - - -// Todo: add alignment support. -// Sfinae based const/non const variants. -// Column -template -struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T * baseAddress): baseAddress_(baseAddress) {} - SOA_HOST_DEVICE_INLINE T * operator()() { return baseAddress_; } - SOA_HOST_DEVICE_INLINE T & operator()(size_t index) { return baseAddress_[index]; } -private: - T * baseAddress_; -}; - -// Const column -template -struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const T * baseAddress): baseAddress_(baseAddress) {} - SOA_HOST_DEVICE_INLINE const T * operator()() const { return baseAddress_; } - SOA_HOST_DEVICE_INLINE T operator()(size_t index) const { return baseAddress_[index]; } -private: - const T * baseAddress_; -}; - -// Scalar -template -struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T * baseAddress): baseAddress_(baseAddress) {} - SOA_HOST_DEVICE_INLINE T & operator() () { return *baseAddress_; } - SOA_HOST_DEVICE_INLINE void operator() (size_t index) const { assert (false && "Indexed access impossible for SoA scalars."); } -private: - T * baseAddress_; -}; - -// Const scalar -template -struct SoAColumnAccessorsImpl { - SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const T * baseAddress): baseAddress_(baseAddress) {} - SOA_HOST_DEVICE_INLINE T operator() () const { return *baseAddress_; } - SOA_HOST_DEVICE_INLINE void operator() (size_t index) const { assert (false && "Indexed access impossible for SoA scalars."); } -private: - const T * baseAddress_; -}; - -/* A helper template stager avoiding comma in macros */ -template -struct SoAAccessors{ - using myInt = int; - template - struct ColumnType { + /* Column accessors: templates implementing the global accesors (soa::x() and soa::x(index) */ + enum class SoAAccessType : bool { mutableAccess, constAccess }; + + template + struct SoAColumnAccessorsImpl {}; + + // Todo: add alignment support. + // Sfinae based const/non const variants. + // Column + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE T* operator()() { return baseAddress_; } + SOA_HOST_DEVICE_INLINE T& operator()(size_t index) { return baseAddress_[index]; } + + private: + T* baseAddress_; + }; + + // Const column + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE const T* operator()() const { return baseAddress_; } + SOA_HOST_DEVICE_INLINE T operator()(size_t index) const { return baseAddress_[index]; } + + private: + const T* baseAddress_; + }; + + // Scalar + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE T& operator()() { return *baseAddress_; } + SOA_HOST_DEVICE_INLINE void operator()(size_t index) const { + assert(false && "Indexed access impossible for SoA scalars."); + } + + private: + T* baseAddress_; + }; + + // Const scalar + template + struct SoAColumnAccessorsImpl { + SOA_HOST_DEVICE_INLINE SoAColumnAccessorsImpl(const T* baseAddress) : baseAddress_(baseAddress) {} + SOA_HOST_DEVICE_INLINE T operator()() const { return *baseAddress_; } + SOA_HOST_DEVICE_INLINE void operator()(size_t index) const { + assert(false && "Indexed access impossible for SoA scalars."); + } + + private: + const T* baseAddress_; + }; + + /* A helper template stager avoiding comma in macros */ + template + struct SoAAccessors { using myInt = int; - template - struct AccessType: public SoAColumnAccessorsImpl { + template + struct ColumnType { using myInt = int; - using SoAColumnAccessorsImpl::SoAColumnAccessorsImpl; + template + struct AccessType : public SoAColumnAccessorsImpl { + using myInt = int; + using SoAColumnAccessorsImpl::SoAColumnAccessorsImpl; + }; }; }; -}; -/* Enum parameters allowing templated control of layout/view behaviors */ -/* Alignement enforcement verifies every column is aligned, and + /* Enum parameters allowing templated control of layout/view behaviors */ + /* Alignement enforcement verifies every column is aligned, and * hints the compiler that it can expect column pointers to be aligned */ -enum class AlignmentEnforcement : bool { Relaxed, Enforced }; - -struct CacheLineSize { - static constexpr size_t NvidiaGPU = 128; - static constexpr size_t IntelCPU = 64; - static constexpr size_t AMDCPU = 64; - static constexpr size_t ARMCPU = 64; - static constexpr size_t defaultSize = NvidiaGPU; -}; + enum class AlignmentEnforcement : bool { Relaxed, Enforced }; + + struct CacheLineSize { + static constexpr size_t NvidiaGPU = 128; + static constexpr size_t IntelCPU = 64; + static constexpr size_t AMDCPU = 64; + static constexpr size_t ARMCPU = 64; + static constexpr size_t defaultSize = NvidiaGPU; + }; -} // namespace cms::soa +} // namespace cms::soa #endif // ndef DataStructures_SoACommon_h diff --git a/src/cudadev/DataFormats/SoALayout.h b/src/cudadev/DataFormats/SoALayout.h index 058d7d91a..0f17adde7 100644 --- a/src/cudadev/DataFormats/SoALayout.h +++ b/src/cudadev/DataFormats/SoALayout.h @@ -92,37 +92,33 @@ /** * SoAMetadata member computing column pitch */ -#define _DEFINE_METADATA_MEMBERS_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ - _SWITCH_ON_TYPE( \ - VALUE_TYPE, \ - /* Scalar */ \ - size_t BOOST_PP_CAT(NAME, Pitch()) const { \ - return (((sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * ParentClass::byteAlignment; \ - } typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ - constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::scalar; \ - CPP_TYPE const * BOOST_PP_CAT(addressOf_, NAME)() const { return parent_.BOOST_PP_CAT(NAME, _); } \ - CPP_TYPE * BOOST_PP_CAT(addressOf_, NAME)() { return parent_.BOOST_PP_CAT(NAME, _); } \ - , \ - /* Column */ \ - CPP_TYPE const * BOOST_PP_CAT(addressOf_, NAME)() const { return parent_.BOOST_PP_CAT(NAME, _); } \ - CPP_TYPE * BOOST_PP_CAT(addressOf_, NAME)() { return parent_.BOOST_PP_CAT(NAME, _); } \ - size_t BOOST_PP_CAT(NAME, Pitch()) const { \ - return (((parent_.nElements_ * sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * \ - ParentClass::byteAlignment; \ - } \ - typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ - constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::column; \ - , \ - /* Eigen column */ \ - size_t BOOST_PP_CAT(NAME, Pitch()) const { \ - return (((parent_.nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / ParentClass::byteAlignment) + 1) * \ - ParentClass::byteAlignment * CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime; \ - } \ - typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ - constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::eigen; \ - CPP_TYPE::Scalar const * BOOST_PP_CAT(addressOf_, NAME)() const { return parent_.BOOST_PP_CAT(NAME, _); } \ - CPP_TYPE::Scalar * BOOST_PP_CAT(addressOf_, NAME)() { return parent_.BOOST_PP_CAT(NAME, _); } \ - ) +#define _DEFINE_METADATA_MEMBERS_IMPL(VALUE_TYPE, CPP_TYPE, NAME) \ + _SWITCH_ON_TYPE( \ + VALUE_TYPE, /* Scalar */ \ + size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * ParentClass::byteAlignment; \ + } typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::scalar; \ + CPP_TYPE const* BOOST_PP_CAT(addressOf_, NAME)() const { \ + return parent_.BOOST_PP_CAT(NAME, _); \ + } CPP_TYPE* BOOST_PP_CAT(addressOf_, NAME)() { return parent_.BOOST_PP_CAT(NAME, _); }, /* Column */ \ + CPP_TYPE const* BOOST_PP_CAT(addressOf_, NAME)() \ + const { return parent_.BOOST_PP_CAT(NAME, _); } CPP_TYPE* BOOST_PP_CAT(addressOf_, NAME)() { \ + return parent_.BOOST_PP_CAT(NAME, _); \ + } size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((parent_.nElements_ * sizeof(CPP_TYPE) - 1) / ParentClass::byteAlignment) + 1) * \ + ParentClass::byteAlignment; \ + } typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::column; \ + , /* Eigen column */ \ + size_t BOOST_PP_CAT(NAME, Pitch()) const { \ + return (((parent_.nElements_ * sizeof(CPP_TYPE::Scalar) - 1) / ParentClass::byteAlignment) + 1) * \ + ParentClass::byteAlignment * CPP_TYPE::RowsAtCompileTime * CPP_TYPE::ColsAtCompileTime; \ + } typedef CPP_TYPE BOOST_PP_CAT(TypeOf_, NAME); \ + constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, NAME) = cms::soa::SoAColumnType::eigen; \ + CPP_TYPE::Scalar const* BOOST_PP_CAT(addressOf_, NAME)() const { \ + return parent_.BOOST_PP_CAT(NAME, _); \ + } CPP_TYPE::Scalar* BOOST_PP_CAT(addressOf_, NAME)() { return parent_.BOOST_PP_CAT(NAME, _); }) #define _DEFINE_METADATA_MEMBERS(R, DATA, TYPE_NAME) _DEFINE_METADATA_MEMBERS_IMPL TYPE_NAME @@ -227,18 +223,19 @@ /* * A macro defining a SoA layout (collection of scalars and columns of equal lengths) */ -#define GENERATE_SOA_LAYOUT(CLASS, ...) \ +// clang-format off +#define GENERATE_SOA_LAYOUT(CLASS, ...) \ template \ + cms::soa::AlignmentEnforcement ALIGNMENT_ENFORCEMENT = cms::soa::AlignmentEnforcement::Relaxed> \ struct CLASS { \ /* these could be moved to an external type trait to free up the symbol names */ \ using self_type = CLASS; \ typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ \ - /* For CUDA applications, we align to the 128 bytes of the cache lines. \ - * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ - * up to compute capability 8.X. \ - */ \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ constexpr static size_t defaultAlignment = 128; \ constexpr static size_t byteAlignment = ALIGNMENT; \ constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ @@ -271,22 +268,25 @@ return ret; \ } \ \ - /** \ - * Helper/friend class allowing SoA introspection. \ - */ \ + /** \ + * Helper/friend class allowing SoA introspection. \ + */ \ struct SoAMetadata { \ friend CLASS; \ SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ SOA_HOST_DEVICE_INLINE size_t byteSize() const { return parent_.byteSize_; } \ SOA_HOST_DEVICE_INLINE size_t byteAlignment() const { return CLASS::byteAlignment; } \ - SOA_HOST_DEVICE_INLINE std::byte* data() { return parent_.mem_; } \ + SOA_HOST_DEVICE_INLINE std::byte* data() { return parent_.mem_; } \ SOA_HOST_DEVICE_INLINE const std::byte* data() const { return parent_.mem_; } \ SOA_HOST_DEVICE_INLINE std::byte* nextByte() const { return parent_.mem_ + parent_.byteSize_; } \ - SOA_HOST_DEVICE_INLINE CLASS cloneToNewAddress(std::byte* addr) const { return CLASS(addr, parent_.nElements_); } \ + SOA_HOST_DEVICE_INLINE CLASS cloneToNewAddress(std::byte* addr) const { \ + return CLASS(addr, parent_.nElements_); \ + } \ _ITERATE_ON_ALL(_DEFINE_METADATA_MEMBERS, ~, __VA_ARGS__) \ \ - SoAMetadata & operator=(const SoAMetadata &) = delete; \ - SoAMetadata(const SoAMetadata &) = delete; \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ + \ private: \ SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ const CLASS& parent_; \ @@ -344,5 +344,6 @@ size_t byteSize_; \ _ITERATE_ON_ALL(_DECLARE_SOA_DATA_MEMBER, ~, __VA_ARGS__) \ }; +// clang-format on #endif // ndef DataStructures_SoALayout_h diff --git a/src/cudadev/DataFormats/SoAView.h b/src/cudadev/DataFormats/SoAView.h index 6aee9d1f6..4a34cf759 100644 --- a/src/cudadev/DataFormats/SoAView.h +++ b/src/cudadev/DataFormats/SoAView.h @@ -40,31 +40,33 @@ namespace cms::soa { -/* Traits for the different column type scenarios */ -/* Value traits passes the class as is in the case of column type and return + /* Traits for the different column type scenarios */ + /* Value traits passes the class as is in the case of column type and return * an empty class with functions returning non-scalar as accessors. */ -template -struct ConstValueTraits {}; + template + struct ConstValueTraits {}; -template -struct ConstValueTraits : public C { using C::C; }; + template + struct ConstValueTraits : public C { + using C::C; + }; -template -struct ConstValueTraits { - // Just take to SoAValue type to generate the right constructor. - SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::valueType *) {} - // Any attempt to do anything with the "scalar" value a const element will fail. -}; + template + struct ConstValueTraits { + // Just take to SoAValue type to generate the right constructor. + SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::valueType*) {} + // Any attempt to do anything with the "scalar" value a const element will fail. + }; -template -struct ConstValueTraits { - // Just take to SoAValue type to generate the right constructor. - SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::valueType *) {} - // TODO: implement - // Any attempt to do anything with the eigen value a const element will fail. -}; + template + struct ConstValueTraits { + // Just take to SoAValue type to generate the right constructor. + SOA_HOST_DEVICE_INLINE ConstValueTraits(size_t, const typename C::valueType*) {} + // TODO: implement + // Any attempt to do anything with the eigen value a const element will fail. + }; -} // namespace cms::soa; +} // namespace cms::soa #include /* @@ -84,14 +86,14 @@ struct ConstValueTraits { */ #define _DECLARE_VIEW_MEMBER_TYPE_ALIAS_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME, DATA) \ typedef typename BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::SoAMetadata::BOOST_PP_CAT(TypeOf_, LAYOUT_MEMBER) \ - BOOST_PP_CAT(TypeOf_, LOCAL_NAME); \ + BOOST_PP_CAT(TypeOf_, LOCAL_NAME); \ constexpr static cms::soa::SoAColumnType BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) = \ BOOST_PP_CAT(TypeOf_, LAYOUT_NAME)::SoAMetadata::BOOST_PP_CAT(ColumnTypeOf_, LAYOUT_MEMBER); \ - SOA_HOST_DEVICE_INLINE \ - DATA BOOST_PP_CAT(TypeOf_, LOCAL_NAME) * BOOST_PP_CAT(addressOf_, LOCAL_NAME)() const { \ - return parent_.BOOST_PP_CAT(LOCAL_NAME, _); \ - }; \ - static_assert(BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) != cms::soa::SoAColumnType::eigen, \ + SOA_HOST_DEVICE_INLINE \ + DATA BOOST_PP_CAT(TypeOf_, LOCAL_NAME) * BOOST_PP_CAT(addressOf_, LOCAL_NAME)() const { \ + return parent_.BOOST_PP_CAT(LOCAL_NAME, _); \ + }; \ + static_assert(BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) != cms::soa::SoAColumnType::eigen, \ "Eigen columns not supported in views."); #define _DECLARE_VIEW_MEMBER_TYPE_ALIAS(R, DATA, LAYOUT_MEMBER_NAME) \ @@ -121,21 +123,22 @@ struct ConstValueTraits { (DATA typename BOOST_PP_CAT(SoAMetadata::TypeOf_, LOCAL_NAME) * LOCAL_NAME) #define _DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS(R, DATA, LAYOUT_MEMBER_NAME) \ - BOOST_PP_EXPAND(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) + BOOST_PP_EXPAND( \ + _DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) /** * Generator of member initialization from constructor. * We use a lambda with auto return type to handle multiple possible return types. */ -#define _DECLARE_VIEW_MEMBER_INITIALIZERS_IMPL(LAYOUT, MEMBER, NAME) \ - (BOOST_PP_CAT(NAME, _)([&]() -> auto { \ - static_assert(BOOST_PP_CAT(SoAMetadata::ColumnTypeOf_, NAME) != cms::soa::SoAColumnType::eigen, \ - "Eigen values not supported in views"); \ - auto addr = LAYOUT.soaMetadata().BOOST_PP_CAT(addressOf_, MEMBER)(); \ - if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ - if (reinterpret_cast(addr) % byteAlignment) \ - throw std::out_of_range("In constructor by layout: misaligned column: " #NAME); \ - return addr; \ +#define _DECLARE_VIEW_MEMBER_INITIALIZERS_IMPL(LAYOUT, MEMBER, NAME) \ + (BOOST_PP_CAT(NAME, _)([&]() -> auto { \ + static_assert(BOOST_PP_CAT(SoAMetadata::ColumnTypeOf_, NAME) != cms::soa::SoAColumnType::eigen, \ + "Eigen values not supported in views"); \ + auto addr = LAYOUT.soaMetadata().BOOST_PP_CAT(addressOf_, MEMBER)(); \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (reinterpret_cast(addr) % byteAlignment) \ + throw std::out_of_range("In constructor by layout: misaligned column: " #NAME); \ + return addr; \ }())) #define _DECLARE_VIEW_MEMBER_INITIALIZERS(R, DATA, LAYOUT_MEMBER_NAME) \ @@ -145,28 +148,27 @@ struct ConstValueTraits { * Generator of size computation for constructor. * This is the per-layout part of the lambda checking they all have the same size. */ -#define _UPDATE_SIZE_OF_VIEW_IMPL(LAYOUT_TYPE, LAYOUT_NAME) \ - if (set) { \ - if (ret != LAYOUT_NAME.soaMetadata().size()) \ - throw std::out_of_range("In constructor by layout: different sizes from layouts."); \ - } else { \ - ret = LAYOUT_NAME.soaMetadata().size(); \ - set = true; \ +#define _UPDATE_SIZE_OF_VIEW_IMPL(LAYOUT_TYPE, LAYOUT_NAME) \ + if (set) { \ + if (ret != LAYOUT_NAME.soaMetadata().size()) \ + throw std::out_of_range("In constructor by layout: different sizes from layouts."); \ + } else { \ + ret = LAYOUT_NAME.soaMetadata().size(); \ + set = true; \ } -#define _UPDATE_SIZE_OF_VIEW(R, DATA, TYPE_NAME) \ - BOOST_PP_EXPAND(_UPDATE_SIZE_OF_VIEW_IMPL TYPE_NAME) +#define _UPDATE_SIZE_OF_VIEW(R, DATA, TYPE_NAME) BOOST_PP_EXPAND(_UPDATE_SIZE_OF_VIEW_IMPL TYPE_NAME) /** * Generator of member initialization from constructor. * We use a lambda with auto return type to handle multiple possible return types. */ -#define _DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN_IMPL(LAYOUT, MEMBER, NAME) \ - (BOOST_PP_CAT(NAME, _)([&]() -> auto { \ - if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ - if (reinterpret_cast(NAME) % byteAlignment) \ +#define _DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN_IMPL(LAYOUT, MEMBER, NAME) \ + (BOOST_PP_CAT(NAME, _)([&]() -> auto { \ + if constexpr (alignmentEnforcement == AlignmentEnforcement::Enforced) \ + if (reinterpret_cast(NAME) % byteAlignment) \ throw std::out_of_range("In constructor by column: misaligned column: " #NAME); \ - return NAME; \ + return NAME; \ }())) #define _DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN(R, DATA, LAYOUT_MEMBER_NAME) \ @@ -208,9 +210,11 @@ struct ConstValueTraits { /** * Declaration of the members accessors of the const element subclass */ -#define _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ - SOA_HOST_DEVICE_INLINE typename SoAConstValueWithConf::RefToConst LOCAL_NAME() const { \ - return BOOST_PP_CAT(LOCAL_NAME, _)(); \ +#define _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + SOA_HOST_DEVICE_INLINE \ + typename SoAConstValueWithConf::RefToConst LOCAL_NAME() \ + const { \ + return BOOST_PP_CAT(LOCAL_NAME, _)(); \ } #define _DECLARE_VIEW_CONST_ELEMENT_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ @@ -219,11 +223,10 @@ struct ConstValueTraits { /** * Declaration of the private members of the const element subclass */ -#define _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ - const cms::soa::ConstValueTraits< \ - SoAConstValueWithConf, \ - BOOST_PP_CAT(SoAMetadata::ColumnTypeOf_, LOCAL_NAME) \ - > BOOST_PP_CAT(LOCAL_NAME, _); +#define _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + const cms::soa::ConstValueTraits, \ + BOOST_PP_CAT(SoAMetadata::ColumnTypeOf_, LOCAL_NAME)> \ + BOOST_PP_CAT(LOCAL_NAME, _); #define _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER(R, DATA, LAYOUT_MEMBER_NAME) \ _DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER_IMPL LAYOUT_MEMBER_NAME @@ -231,8 +234,8 @@ struct ConstValueTraits { /** * Generator of the member-by-member copy operator of the element subclass. */ -#define _DECLARE_VIEW_ELEMENT_VALUE_COPY_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ - if constexpr (SoAMetadata:: BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) != cms::soa::SoAColumnType::scalar) \ +#define _DECLARE_VIEW_ELEMENT_VALUE_COPY_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + if constexpr (SoAMetadata::BOOST_PP_CAT(ColumnTypeOf_, LOCAL_NAME) != cms::soa::SoAColumnType::scalar) \ LOCAL_NAME() = other.LOCAL_NAME(); #define _DECLARE_VIEW_ELEMENT_VALUE_COPY(R, DATA, LAYOUT_MEMBER_NAME) \ @@ -258,18 +261,18 @@ struct ConstValueTraits { /** * Direct access to column pointer and indexed access */ -#define _DECLARE_VIEW_SOA_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ - /* Column or scalar */ \ - SOA_HOST_DEVICE_INLINE auto LOCAL_NAME() { \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType:: \ - template AccessType(BOOST_PP_CAT(LOCAL_NAME, _))(); \ - } \ - SOA_HOST_DEVICE_INLINE auto LOCAL_NAME(size_t index) { \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType:: \ - template AccessType(BOOST_PP_CAT(LOCAL_NAME, _))(index); \ - } +#define _DECLARE_VIEW_SOA_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + /* Column or scalar */ \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME() { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>(BOOST_PP_CAT(LOCAL_NAME, _))(); \ + } \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME(size_t index) { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::mutableAccess>(BOOST_PP_CAT(LOCAL_NAME, _))(index); \ + } #define _DECLARE_VIEW_SOA_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_ACCESSOR_IMPL LAYOUT_MEMBER_NAME) @@ -277,18 +280,18 @@ struct ConstValueTraits { /** * Direct access to column pointer (const) and indexed access. */ -#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ - /* Column or scalar */ \ - SOA_HOST_DEVICE_INLINE auto LOCAL_NAME() const { \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType:: \ - template AccessType(BOOST_PP_CAT(LOCAL_NAME, _))(); \ - } \ - SOA_HOST_DEVICE_INLINE auto LOCAL_NAME(size_t index) const { \ - return typename cms::soa::SoAAccessors:: \ - template ColumnType:: \ - template AccessType(BOOST_PP_CAT(LOCAL_NAME, _))(index); \ - } +#define _DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL(LAYOUT_NAME, LAYOUT_MEMBER, LOCAL_NAME) \ + /* Column or scalar */ \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME() const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, _))(); \ + } \ + SOA_HOST_DEVICE_INLINE auto LOCAL_NAME(size_t index) const { \ + return typename cms::soa::SoAAccessors:: \ + template ColumnType::template AccessType< \ + cms::soa::SoAAccessType::constAccess>(BOOST_PP_CAT(LOCAL_NAME, _))(index); \ + } #define _DECLARE_VIEW_SOA_CONST_ACCESSOR(R, DATA, LAYOUT_MEMBER_NAME) \ BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_CONST_ACCESSOR_IMPL LAYOUT_MEMBER_NAME) @@ -303,21 +306,155 @@ struct ConstValueTraits { BOOST_PP_EXPAND(_DECLARE_VIEW_SOA_MEMBER_IMPL BOOST_PP_TUPLE_PUSH_BACK(LAYOUT_MEMBER_NAME, DATA)) /* ---- MUTABLE VIEW -------------------------------------------------------------------------------------------------------------------- */ +// clang-format off +#define GENERATE_SOA_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ + template \ + struct CLASS { \ + /* these could be moved to an external type trait to free up the symbol names */ \ + using self_type = CLASS; \ + typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ + \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ + constexpr static size_t defaultAlignment = cms::soa::CacheLineSize::defaultSize; \ + constexpr static size_t byteAlignment = ALIGNMENT; \ + constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ + constexpr static size_t conditionalAlignment = \ + alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ + constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \ + constexpr static cms::soa::RangeChecking rangeChecking = RANGE_CHECKING; \ + /* Those typedefs avoid having commas in macros (which is problematic) */ \ + template \ + using SoAValueWithConf = cms::soa::SoAValue; \ + \ + template \ + using SoAConstValueWithConf = cms::soa::SoAConstValue; \ + \ + template \ + using SoAEigenValueWithConf = cms::soa::SoAEigenValue; \ + /** \ + * Helper/friend class allowing SoA introspection. \ + */ \ + struct SoAMetadata { \ + friend CLASS; \ + SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ + /* Alias layout or view types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS, ~, LAYOUTS_LIST) \ + \ + /* Alias member types to name-derived identifyer to allow simpler definitions */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, BOOST_PP_EMPTY(), VALUE_LIST) \ + \ + /* Forbid copying to avoid const correctness evasion */ \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ + \ + private: \ + SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ + const CLASS& parent_; \ + }; \ + friend SoAMetadata; \ + SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ + SOA_HOST_DEVICE_INLINE SoAMetadata soaMetadata() { return SoAMetadata(*this); } \ + \ + /* Trivial constuctor */ \ + CLASS() : nElements_(0), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_TRIVIAL_CONSTRUCTION, ~, VALUE_LIST) {} \ + \ + /* Constructor relying on user provided layouts or views */ \ + SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, BOOST_PP_EMPTY(), LAYOUTS_LIST)) \ + : nElements_([&]() -> size_t { \ + bool set = false; \ + size_t ret = 0; \ + _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ + return ret; \ + }()), \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS, ~, VALUE_LIST) {} \ + \ + /* Constructor relying on individually provided column addresses */ \ + SOA_HOST_ONLY CLASS(size_t nElements, \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, \ + BOOST_PP_EMPTY(), \ + VALUE_LIST)) \ + : nElements_(nElements), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN, ~, VALUE_LIST) {} \ + \ + struct const_element { \ + SOA_HOST_DEVICE_INLINE \ + const_element(size_t index, /* Declare parameters */ \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, const, VALUE_LIST)) \ + : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONST_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_ACCESSOR, ~, VALUE_LIST) \ + \ + private: \ + _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ + }; \ + \ + struct element { \ + SOA_HOST_DEVICE_INLINE \ + element(size_t index, /* Declare parameters */ \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, BOOST_PP_EMPTY(), VALUE_LIST)) \ + : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ + SOA_HOST_DEVICE_INLINE \ + element& operator=(const element& other) { \ + _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_COPY, ~, VALUE_LIST) \ + return *this; \ + } \ + _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ + }; \ + \ + /* AoS-like accessor (non-const) */ \ + SOA_HOST_DEVICE_INLINE \ + element operator[](size_t index) { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + } \ + return element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ + } \ + \ + /* AoS-like accessor (const) */ \ + SOA_HOST_DEVICE_INLINE \ + const_element operator[](size_t index) const { \ + if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + } \ + return const_element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ + } \ + \ + /* accessors */ \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_ACCESSOR, ~, VALUE_LIST) \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_CONST_ACCESSOR, ~, VALUE_LIST) \ + \ + /* dump the SoA internal structure */ \ + template \ + SOA_HOST_ONLY friend void dump(); \ + \ + private: \ + size_t nElements_; \ + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_MEMBER, BOOST_PP_EMPTY(), VALUE_LIST) \ + }; +// clang-format on -#define GENERATE_SOA_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ +/* ---- CONST VIEW --------------------------------------------------------------------------------------------------------------------- */ +// clang-format off +#define GENERATE_SOA_CONST_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ template \ + cms::soa::AlignmentEnforcement ALIGNMENT_ENFORCEMENT = cms::soa::AlignmentEnforcement::Relaxed, \ + cms::soa::RestrictQualify RESTRICT_QUALIFY = cms::soa::RestrictQualify::Enabled, \ + cms::soa::RangeChecking RANGE_CHECKING = cms::soa::RangeChecking::Disabled> \ struct CLASS { \ /* these could be moved to an external type trait to free up the symbol names */ \ using self_type = CLASS; \ typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ - \ - /* For CUDA applications, we align to the 128 bytes of the cache lines. \ - * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ - * up to compute capability 8.X. \ - */ \ + \ + /* For CUDA applications, we align to the 128 bytes of the cache lines. \ + * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ + * up to compute capability 8.X. \ + */ \ constexpr static size_t defaultAlignment = cms::soa::CacheLineSize::defaultSize; \ constexpr static size_t byteAlignment = ALIGNMENT; \ constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ @@ -325,7 +462,7 @@ struct ConstValueTraits { alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \ constexpr static cms::soa::RangeChecking rangeChecking = RANGE_CHECKING; \ -/* Those typedefs avoid having commas in macros (which is problematic) */ \ + /* Those typedefs avoid having commas in macros (which is problematic) */ \ template \ using SoAValueWithConf = cms::soa::SoAValue; \ \ @@ -334,47 +471,45 @@ struct ConstValueTraits { \ template \ using SoAEigenValueWithConf = cms::soa::SoAEigenValue; \ + \ /** \ * Helper/friend class allowing SoA introspection. \ - */ \ + */ \ struct SoAMetadata { \ friend CLASS; \ - SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ - /* Alias layout or view types to name-derived identifyer to allow simpler definitions */ \ + SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ + /* Alias layout/view types to name-derived identifyer to allow simpler definitions */ \ _ITERATE_ON_ALL(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS, ~, LAYOUTS_LIST) \ \ /* Alias member types to name-derived identifyer to allow simpler definitions */ \ - _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, BOOST_PP_EMPTY(), VALUE_LIST) \ + _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, const, VALUE_LIST) \ + \ + SoAMetadata& operator=(const SoAMetadata&) = delete; \ + SoAMetadata(const SoAMetadata&) = delete; \ \ - /* Forbid copying to avoid const correctness evasion */ \ - SoAMetadata & operator=(const SoAMetadata &) = delete; \ - SoAMetadata(const SoAMetadata &) = delete; \ private: \ SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ const CLASS& parent_; \ }; \ friend SoAMetadata; \ SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ - SOA_HOST_DEVICE_INLINE SoAMetadata soaMetadata() { return SoAMetadata(*this); } \ \ /* Trivial constuctor */ \ CLASS() : nElements_(0), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_TRIVIAL_CONSTRUCTION, ~, VALUE_LIST) {} \ \ /* Constructor relying on user provided layouts or views */ \ - SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, BOOST_PP_EMPTY(), LAYOUTS_LIST)) \ - : nElements_( \ - [&]() -> size_t { \ - bool set = false; \ - size_t ret = 0; \ - _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ - return ret; \ - }() \ - ), \ + SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, const, LAYOUTS_LIST)) \ + : nElements_([&]() -> size_t { \ + bool set = false; \ + size_t ret = 0; \ + _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ + return ret; \ + }()), \ _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS, ~, VALUE_LIST) {} \ \ /* Constructor relying on individually provided column addresses */ \ SOA_HOST_ONLY CLASS(size_t nElements, \ - _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, BOOST_PP_EMPTY(), VALUE_LIST)) \ + _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, const, VALUE_LIST)) \ : nElements_(nElements), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN, ~, VALUE_LIST) {} \ \ struct const_element { \ @@ -388,39 +523,17 @@ struct ConstValueTraits { _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ }; \ \ - struct element { \ - SOA_HOST_DEVICE_INLINE \ - element(size_t index, /* Declare parameters */ \ - _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, BOOST_PP_EMPTY(), VALUE_LIST)) \ - : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ - SOA_HOST_DEVICE_INLINE \ - element& operator=(const element& other) { \ - _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_COPY, ~, VALUE_LIST) \ - return *this; \ - } \ - _ITERATE_ON_ALL(_DECLARE_VIEW_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ - }; \ - \ - /* AoS-like accessor (non-const) */ \ - SOA_HOST_DEVICE_INLINE \ - element operator[](size_t index) { \ - if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ - if (index >= nElements_) SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ - } \ - return element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ - } \ - \ /* AoS-like accessor (const) */ \ SOA_HOST_DEVICE_INLINE \ const_element operator[](size_t index) const { \ if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ - if (index >= nElements_) SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ + if (index >= nElements_) \ + SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ } \ return const_element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ } \ \ /* accessors */ \ - _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_ACCESSOR, ~, VALUE_LIST) \ _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_CONST_ACCESSOR, ~, VALUE_LIST) \ \ /* dump the SoA internal structure */ \ @@ -429,150 +542,50 @@ struct ConstValueTraits { \ private: \ size_t nElements_; \ - _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_MEMBER, BOOST_PP_EMPTY(), VALUE_LIST) \ - }; - -/* ---- CONST VIEW --------------------------------------------------------------------------------------------------------------------- */ - -#define GENERATE_SOA_CONST_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \ - template \ - struct CLASS { \ - /* these could be moved to an external type trait to free up the symbol names */ \ - using self_type = CLASS; \ - typedef cms::soa::AlignmentEnforcement AlignmentEnforcement; \ - \ - /* For CUDA applications, we align to the 128 bytes of the cache lines. \ - * See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-memory-3-0 this is still valid \ - * up to compute capability 8.X. \ - */ \ - constexpr static size_t defaultAlignment = cms::soa::CacheLineSize::defaultSize; \ - constexpr static size_t byteAlignment = ALIGNMENT; \ - constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \ - constexpr static size_t conditionalAlignment = \ - alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \ - constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \ - constexpr static cms::soa::RangeChecking rangeChecking = RANGE_CHECKING; \ - /* Those typedefs avoid having commas in macros (which is problematic) */ \ - template \ - using SoAValueWithConf = cms::soa::SoAValue; \ - \ - template \ - using SoAConstValueWithConf = cms::soa::SoAConstValue; \ - \ - template \ - using SoAEigenValueWithConf = cms::soa::SoAEigenValue; \ - \ - /** \ - * Helper/friend class allowing SoA introspection. \ - */ \ - struct SoAMetadata { \ - friend CLASS; \ - SOA_HOST_DEVICE_INLINE size_t size() const { return parent_.nElements_; } \ - /* Alias layout/view types to name-derived identifyer to allow simpler definitions */ \ - _ITERATE_ON_ALL(_DECLARE_VIEW_LAYOUT_TYPE_ALIAS, ~, LAYOUTS_LIST) \ - \ - /* Alias member types to name-derived identifyer to allow simpler definitions */ \ - _ITERATE_ON_ALL(_DECLARE_VIEW_MEMBER_TYPE_ALIAS, const, VALUE_LIST) \ - \ - SoAMetadata & operator=(const SoAMetadata &) = delete; \ - SoAMetadata(const SoAMetadata &) = delete; \ - private: \ - SOA_HOST_DEVICE_INLINE SoAMetadata(const CLASS& parent) : parent_(parent) {} \ - const CLASS& parent_; \ - }; \ - friend SoAMetadata; \ - SOA_HOST_DEVICE_INLINE const SoAMetadata soaMetadata() const { return SoAMetadata(*this); } \ - \ - /* Trivial constuctor */ \ - CLASS() : nElements_(0), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_TRIVIAL_CONSTRUCTION, ~, VALUE_LIST) {} \ - \ - /* Constructor relying on user provided layouts or views */ \ - SOA_HOST_ONLY CLASS(_ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_PARAMETERS, const, LAYOUTS_LIST)) \ - : nElements_( \ - [&]() -> size_t { \ - bool set = false; \ - size_t ret = 0; \ - _ITERATE_ON_ALL(_UPDATE_SIZE_OF_VIEW, BOOST_PP_EMPTY(), LAYOUTS_LIST) \ - return ret; \ - }() \ - ), \ - _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS, ~, VALUE_LIST) {} \ - \ - /* Constructor relying on individually provided column addresses */ \ - SOA_HOST_ONLY CLASS(size_t nElements, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONSTRUCTION_BYCOLUMN_PARAMETERS, const, VALUE_LIST)) \ - : nElements_(nElements), _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_MEMBER_INITIALIZERS_BYCOLUMN, ~, VALUE_LIST) {} \ - \ - struct const_element { \ - SOA_HOST_DEVICE_INLINE \ - const_element(size_t index, /* Declare parameters */ \ - _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_VALUE_ARG, const, VALUE_LIST)) \ - : _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_CONST_ELEM_MEMBER_INIT, index, VALUE_LIST) {} \ - _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_ACCESSOR, ~, VALUE_LIST) \ - \ - private: \ - _ITERATE_ON_ALL(_DECLARE_VIEW_CONST_ELEMENT_VALUE_MEMBER, ~, VALUE_LIST) \ - }; \ - \ - /* AoS-like accessor (const) */ \ - SOA_HOST_DEVICE_INLINE \ - const_element operator[](size_t index) const { \ - if constexpr (rangeChecking == cms::soa::RangeChecking::Enabled) { \ - if (index >= nElements_) SOA_THROW_OUT_OF_RANGE("Out of range index in " #CLASS "::operator[]") \ - } \ - return const_element(index, _ITERATE_ON_ALL_COMMA(_DECLARE_VIEW_ELEMENT_CONSTR_CALL, ~, VALUE_LIST)); \ - } \ - \ - /* accessors */ \ - _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_CONST_ACCESSOR, ~, VALUE_LIST) \ - \ - /* dump the SoA internal structure */ \ - template \ - SOA_HOST_ONLY friend void dump(); \ - \ - private: \ - size_t nElements_; \ - _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_MEMBER, const, VALUE_LIST) \ - }; + _ITERATE_ON_ALL(_DECLARE_VIEW_SOA_MEMBER, const, VALUE_LIST) \ +}; +// clang-format on /** * Helper macro turning layout field declaration into view field declaration. */ -#define _VIEW_FIELD_FROM_LAYOUT_IMPL(VALUE_TYPE, CPP_TYPE, NAME, DATA) \ - (DATA, NAME, NAME) +#define _VIEW_FIELD_FROM_LAYOUT_IMPL(VALUE_TYPE, CPP_TYPE, NAME, DATA) (DATA, NAME, NAME) #define _VIEW_FIELD_FROM_LAYOUT(R, DATA, VALUE_TYPE_NAME) \ - BOOST_PP_EXPAND ((_VIEW_FIELD_FROM_LAYOUT_IMPL BOOST_PP_TUPLE_PUSH_BACK(VALUE_TYPE_NAME, DATA))) + BOOST_PP_EXPAND((_VIEW_FIELD_FROM_LAYOUT_IMPL BOOST_PP_TUPLE_PUSH_BACK(VALUE_TYPE_NAME, DATA))) /** * A macro defining both layout and view(s) in one go. */ -#define GENERATE_SOA_LAYOUT_VIEW_AND_CONST_VIEW(LAYOUT_NAME, VIEW_NAME, CONST_VIEW_NAME, ... ) \ -GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ -using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME <>; \ -GENERATE_SOA_VIEW(VIEW_NAME, \ - SOA_VIEW_LAYOUT_LIST( (BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME)) ), \ - SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); \ -GENERATE_SOA_CONST_VIEW(CONST_VIEW_NAME, \ - SOA_VIEW_LAYOUT_LIST( (BOOST_PP_CAT(LAYOUT_NAME,_default), BOOST_PP_CAT(instance_, LAYOUT_NAME)) ), \ - SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); - -#define GENERATE_SOA_LAYOUT_AND_VIEW(LAYOUT_NAME, VIEW_NAME, ... ) \ -GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ -using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME <>; \ -GENERATE_SOA_VIEW(VIEW_NAME, \ - SOA_VIEW_LAYOUT_LIST( (BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME)) ), \ - SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); - - -#define GENERATE_SOA_LAYOUT_AND_CONST_VIEW(LAYOUT_NAME, CONST_VIEW_NAME, ... ) \ -GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ -using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME <>; \ -GENERATE_SOA_CONST_VIEW(CONST_VIEW_NAME, \ - SOA_VIEW_LAYOUT_LIST( (BOOST_PP_CAT(LAYOUT_NAME,_default), BOOST_PP_CAT(instance_, LAYOUT_NAME)) ), \ - SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); +#define GENERATE_SOA_LAYOUT_VIEW_AND_CONST_VIEW(LAYOUT_NAME, VIEW_NAME, CONST_VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_VIEW(VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA( \ + _VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); \ + GENERATE_SOA_CONST_VIEW( \ + CONST_VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST( \ + _ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); + +#define GENERATE_SOA_LAYOUT_AND_VIEW(LAYOUT_NAME, VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_VIEW(VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST(_ITERATE_ON_ALL_COMMA( \ + _VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); + +#define GENERATE_SOA_LAYOUT_AND_CONST_VIEW(LAYOUT_NAME, CONST_VIEW_NAME, ...) \ + GENERATE_SOA_LAYOUT(LAYOUT_NAME, __VA_ARGS__); \ + using BOOST_PP_CAT(LAYOUT_NAME, _default) = LAYOUT_NAME<>; \ + GENERATE_SOA_CONST_VIEW( \ + CONST_VIEW_NAME, \ + SOA_VIEW_LAYOUT_LIST((BOOST_PP_CAT(LAYOUT_NAME, _default), BOOST_PP_CAT(instance_, LAYOUT_NAME))), \ + SOA_VIEW_VALUE_LIST( \ + _ITERATE_ON_ALL_COMMA(_VIEW_FIELD_FROM_LAYOUT, BOOST_PP_CAT(instance_, LAYOUT_NAME), __VA_ARGS__))); #endif // ndef DataStructures_SoAView_h diff --git a/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.cu b/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.cu index 8711740a3..331874b4f 100644 --- a/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.cu +++ b/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.cu @@ -62,7 +62,6 @@ namespace pixelgpudetails { return (PixelSubdetector::PixelBarrel == ((rawId >> DetId::kSubdetOffset) & DetId::kSubdetMask)); } - //reference http://cmsdoxygen.web.cern.ch/cmsdoxygen/CMSSW_9_2_0/doc/html/dd/d31/FrameConversion_8cc_source.html //http://cmslxr.fnal.gov/source/CondFormats/SiPixelObjects/src/PixelROC.cc?v=CMSSW_9_2_0#0071 // Convert local pixel to pixelgpudetails::global pixel @@ -494,7 +493,7 @@ namespace pixelgpudetails { // Interface to outside void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2, const SiPixelClusterThresholds clusterThresholds, - SiPixelROCsStatusAndMappingConstView & cablingMap, + SiPixelROCsStatusAndMappingConstView &cablingMap, const unsigned char *modToUnp, const SiPixelGainForHLTonGPU *gains, const WordFedAppender &wordFed, diff --git a/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.h b/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.h index 2297c296e..c3ff57103 100644 --- a/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.h +++ b/src/cudadev/plugin-SiPixelClusterizer/SiPixelRawToClusterGPUKernel.h @@ -135,7 +135,7 @@ namespace pixelgpudetails { void makeClustersAsync(bool isRun2, const SiPixelClusterThresholds clusterThresholds, - SiPixelROCsStatusAndMappingConstView & cablingMap, + SiPixelROCsStatusAndMappingConstView& cablingMap, const unsigned char* modToUnp, const SiPixelGainForHLTonGPU* gains, const WordFedAppender& wordFed, diff --git a/src/cudadev/test/SoALayoutAndView_t.cu b/src/cudadev/test/SoALayoutAndView_t.cu index f5fa1794e..71c401063 100644 --- a/src/cudadev/test/SoALayoutAndView_t.cu +++ b/src/cudadev/test/SoALayoutAndView_t.cu @@ -9,118 +9,106 @@ // Multiple stores in a buffer // Scalars, Columns of scalars and of Eigen vectors // View to each of them, from one and multiple stores. - -GENERATE_SOA_LAYOUT_AND_VIEW(SoA1LayoutTemplate, SoA1ViewTemplate, - // predefined static scalars - // size_t size; - // size_t alignment; - // columns: one value per element - SOA_COLUMN(double, x), - SOA_COLUMN(double, y), - SOA_COLUMN(double, z), - SOA_COLUMN(double, sum), - SOA_COLUMN(double, prod), - /* Leave Eigen definitions out until support is complete. +GENERATE_SOA_LAYOUT_AND_VIEW(SoA1LayoutTemplate, + SoA1ViewTemplate, + // predefined static scalars + // size_t size; + // size_t alignment; + + // columns: one value per element + SOA_COLUMN(double, x), + SOA_COLUMN(double, y), + SOA_COLUMN(double, z), + SOA_COLUMN(double, sum), + SOA_COLUMN(double, prod), + /* Leave Eigen definitions out until support is complete. SOA_EIGEN_COLUMN(Eigen::Vector3d, a), SOA_EIGEN_COLUMN(Eigen::Vector3d, b), SOA_EIGEN_COLUMN(Eigen::Vector3d, r),*/ - SOA_COLUMN(uint16_t, color), - SOA_COLUMN(int32_t, value), - SOA_COLUMN(double *, py), - SOA_COLUMN(uint32_t, count), - SOA_COLUMN(uint32_t, anotherCount), + SOA_COLUMN(uint16_t, color), + SOA_COLUMN(int32_t, value), + SOA_COLUMN(double *, py), + SOA_COLUMN(uint32_t, count), + SOA_COLUMN(uint32_t, anotherCount), - // scalars: one value for the whole structure - SOA_SCALAR(const char *, description), - SOA_SCALAR(uint32_t, someNumber) -) + // scalars: one value for the whole structure + SOA_SCALAR(const char *, description), + SOA_SCALAR(uint32_t, someNumber)) using SoA1Layout = SoA1LayoutTemplate<>; using SoA1View = SoA1ViewTemplate<>; // A partial view (artificial mix of store and view) GENERATE_SOA_VIEW(SoA1View2GTemplate, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(SoA1Layout, soa1), - SOA_VIEW_LAYOUT(SoA1View, soa1v) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(soa1, x), - SOA_VIEW_VALUE(soa1v, y), - SOA_VIEW_VALUE(soa1, color), - SOA_VIEW_VALUE(soa1v, value), - SOA_VIEW_VALUE(soa1v, count), - SOA_VIEW_VALUE(soa1, anotherCount), - SOA_VIEW_VALUE(soa1v, description), - SOA_VIEW_VALUE(soa1, someNumber) - ) -) + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(SoA1Layout, soa1), SOA_VIEW_LAYOUT(SoA1View, soa1v)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(soa1, x), + SOA_VIEW_VALUE(soa1v, y), + SOA_VIEW_VALUE(soa1, color), + SOA_VIEW_VALUE(soa1v, value), + SOA_VIEW_VALUE(soa1v, count), + SOA_VIEW_VALUE(soa1, anotherCount), + SOA_VIEW_VALUE(soa1v, description), + SOA_VIEW_VALUE(soa1, someNumber))) using SoA1View2G = SoA1View2GTemplate<>; - - // Same partial view, yet const. GENERATE_SOA_CONST_VIEW(SoA1View2Gconst, - SOA_VIEW_LAYOUT_LIST( - SOA_VIEW_LAYOUT(SoA1Layout, soa1), - SOA_VIEW_LAYOUT(SoA1View, soa1v) - ), - SOA_VIEW_VALUE_LIST( - SOA_VIEW_VALUE(soa1, x), - SOA_VIEW_VALUE(soa1v, y), -/* Eigen columns are not supported in views. + SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(SoA1Layout, soa1), SOA_VIEW_LAYOUT(SoA1View, soa1v)), + SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(soa1, x), + SOA_VIEW_VALUE(soa1v, y), + /* Eigen columns are not supported in views. SoA_view_value(soa1, a, a), SoA_view_value(soa1, b, b), SoA_view_value(soa1, r, r), */ - SOA_VIEW_VALUE(soa1, color), - SOA_VIEW_VALUE(soa1v, value), - SOA_VIEW_VALUE(soa1v, count), - SOA_VIEW_VALUE(soa1, anotherCount), - SOA_VIEW_VALUE(soa1v, description), - SOA_VIEW_VALUE(soa1, someNumber) - ) -) + SOA_VIEW_VALUE(soa1, color), + SOA_VIEW_VALUE(soa1v, value), + SOA_VIEW_VALUE(soa1v, count), + SOA_VIEW_VALUE(soa1, anotherCount), + SOA_VIEW_VALUE(soa1v, description), + SOA_VIEW_VALUE(soa1, someNumber))) // Parameter reusing kernels. The disassembly will indicate whether the compiler uses the wanted cache hits and uses // `restrict` hints avoid multiple reduce loads. // The PTX can be obtained using -ptx insterad of -c when compiling. template -__device__ void addAndMulTemplate ( - T soa, size_t size) { - auto idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx >= size) return; - auto si = soa[idx]; - si.sum() = si.x() + si.y(); - si.prod() = si.x() * si.y(); - } +__device__ void addAndMulTemplate(T soa, size_t size) { + auto idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= size) + return; + auto si = soa[idx]; + si.sum() = si.x() + si.y(); + si.prod() = si.x() * si.y(); +} __global__ void aAMDef(SoA1ViewTemplate soa, size_t size) { + cms::soa::AlignmentEnforcement::Relaxed, + cms::soa::RestrictQualify::Disabled> soa, + size_t size) { addAndMulTemplate(soa, size); } __global__ void aAMRestrict(SoA1ViewTemplate soa, size_t size) { + cms::soa::AlignmentEnforcement::Relaxed, + cms::soa::RestrictQualify::Enabled> soa, + size_t size) { addAndMulTemplate(soa, size); } -const size_t size=10000; +const size_t size = 10000; int main() { // Allocate buffer std::unique_ptr buffer( - static_cast(std::aligned_alloc(SoA1Layout::defaultAlignment, SoA1Layout::computeDataSize(size))), - std::free); + static_cast(std::aligned_alloc(SoA1Layout::defaultAlignment, SoA1Layout::computeDataSize(size))), + std::free); SoA1Layout soa1(buffer.get(), size); - SoA1View soa1view (soa1); - SoA1View2G soa1v2g (soa1, soa1view); - SoA1View2Gconst soa1v2gconst (soa1, soa1view); + SoA1View soa1view(soa1); + SoA1View2G soa1v2g(soa1, soa1view); + SoA1View2Gconst soa1v2gconst(soa1, soa1view); // Write to view - for (size_t i=0; i < size; i++) { + for (size_t i = 0; i < size; i++) { auto s = soa1view[i]; s.x = 1.0 * i; s.y = 2.0 * i; @@ -136,7 +124,7 @@ int main() { s.r() = s.a().cross(s.b());*/ } // Check direct read back - for (size_t i=0; i < size; i++) { + for (size_t i = 0; i < size; i++) { auto s = soa1view[i]; assert(s.x() == 1.0 * i); assert(s.y() == 2.0 * i); @@ -152,7 +140,7 @@ int main() { assert(s.r() == s.a().cross(s.b()));*/ } // Check readback through other views - for (size_t i=0; i < size; i++) { + for (size_t i = 0; i < size; i++) { auto sv = soa1view[i]; auto sv2g = soa1v2g[i]; auto sv2gc = soa1v2gconst[i]; @@ -167,14 +155,18 @@ int main() { assert(sv2gc.y() == 2.0 * i); assert(sv2gc.color() == i); } - + // Validation of range checking try { // Get a view like the default, except for range checking - SoA1ViewTemplate soa1viewRangeChecking(soa1); + SoA1ViewTemplate + soa1viewRangeChecking(soa1); // This should throw an exception [[maybe_unused]] auto si = soa1viewRangeChecking[soa1viewRangeChecking.soaMetadata().size()]; assert(false); - } catch (const std::out_of_range &) {} + } catch (const std::out_of_range &) { + } } \ No newline at end of file