Skip to content

Commit

Permalink
Address pixel local reco PR review comments (#575)
Browse files Browse the repository at this point in the history
Remove SiPixelDigiHeterogeneousConverter as obsolete, should have been removed as part of #100.

Address review comments for SiPixelClustersCUDA:
  - remove commented out default constructor and private: from DeviceConstView;
    this is perhaps the best compromise between non-default constructors not
    being preferred for device allocations, and the use case in
    SiPixelRecHitSoAFromLegacy (for the expected life time of this class)
  - remove const getters with c_ prefix
  - improve constructor parameter name
  - use more initializer list
  - initialize nClusters_h

Address review comments for SiPixelDigiErrorsCUDA:
  - use type alias
  - remove const getters with c_ prefix and other unnecessary methods
  - use more initializer list

Address review comments for SiPixelDigisCUDA:
  - remove const getters with c_ prefix and other unnecessary methods
  - remove commented out default constructor and private: from DeviceConstView
  - add comments for remaining SiPixelDigisCUDA member arrays

Move PixelErrorsCompact and SiPixelDigiErrorsSoa to DataFormats/SiPixelRawData, rename classes

Address review comments for SiPixelErrorsSoA
  - remove redundant assert
  - move constructor inline

Address review comments for SiPixelDigisSoA
  - remove redundant assert
  - add comments

Enable if constexpr also for CUDA in TrackingRecHit2DHeterogeneous

Move dictionary of HostProduct<unsigned int[]> to CUDADataFormats/Common
  • Loading branch information
makortel authored and fwyzard committed Dec 29, 2020
1 parent 0d9b5d3 commit 6a08ba1
Show file tree
Hide file tree
Showing 24 changed files with 146 additions and 105 deletions.
11 changes: 6 additions & 5 deletions CUDADataFormats/Common/BuildFile.xml
Original file line number Diff line number Diff line change
@@ -1,7 +1,8 @@
<iftool name="cuda-gcc-support">
<use name="HeterogeneousCore/CUDAUtilities"/>

<export>
<iftool name="cuda">
<use name="DataFormats/Common"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="rootcore"/>
<export>
<lib name="1"/>
</export>
</export>
</iftool>
7 changes: 7 additions & 0 deletions CUDADataFormats/Common/src/classes.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
#ifndef CUDADataFormats_Common_src_classes_h
#define CUDADataFormats_Common_src_classes_h

#include "CUDADataFormats/Common/interface/HostProduct.h"
#include "DataFormats/Common/interface/Wrapper.h"

#endif // CUDADataFormats_Common_src_classes_h
4 changes: 4 additions & 0 deletions CUDADataFormats/Common/src/classes_def.xml
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
<lcgdict>
<class name="HostProduct<unsigned int[]>" persistent="false"/>
<class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
</lcgdict>
14 changes: 2 additions & 12 deletions CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
class SiPixelClustersCUDA {
public:
SiPixelClustersCUDA() = default;
explicit SiPixelClustersCUDA(size_t maxClusters, cudaStream_t stream);
explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
~SiPixelClustersCUDA() = default;

SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
Expand All @@ -32,23 +32,13 @@ class SiPixelClustersCUDA {
uint32_t const *moduleId() const { return moduleId_d.get(); }
uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); }

uint32_t const *c_moduleStart() const { return moduleStart_d.get(); }
uint32_t const *c_clusInModule() const { return clusInModule_d.get(); }
uint32_t const *c_moduleId() const { return moduleId_d.get(); }
uint32_t const *c_clusModuleStart() const { return clusModuleStart_d.get(); }

class DeviceConstView {
public:
// DeviceConstView() = default;

__device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_ + i); }
__device__ __forceinline__ uint32_t clusInModule(int i) const { return __ldg(clusInModule_ + i); }
__device__ __forceinline__ uint32_t moduleId(int i) const { return __ldg(moduleId_ + i); }
__device__ __forceinline__ uint32_t clusModuleStart(int i) const { return __ldg(clusModuleStart_ + i); }

friend SiPixelClustersCUDA;

// private:
uint32_t const *moduleStart_;
uint32_t const *clusInModule_;
uint32_t const *moduleId_;
Expand All @@ -67,7 +57,7 @@ class SiPixelClustersCUDA {

cms::cuda::device::unique_ptr<DeviceConstView> view_d; // "me" pointer

uint32_t nClusters_h;
uint32_t nClusters_h = 0;
};

#endif
11 changes: 5 additions & 6 deletions CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
Original file line number Diff line number Diff line change
Expand Up @@ -4,12 +4,11 @@
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"

SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxClusters, cudaStream_t stream) {
moduleStart_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
clusInModule_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
moduleId_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
clusModuleStart_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters + 1, stream);

SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
: moduleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)),
clusInModule_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
moduleId_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
clusModuleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)) {
auto view = cms::cuda::make_host_unique<DeviceConstView>(stream);
view->moduleStart_ = moduleStart_d.get();
view->clusInModule_ = clusInModule_d.get();
Expand Down
25 changes: 13 additions & 12 deletions CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
Original file line number Diff line number Diff line change
@@ -1,7 +1,8 @@
#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h

#include "DataFormats/SiPixelDigi/interface/PixelErrors.h"
#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h"
#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h"
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h"
Expand All @@ -10,32 +11,32 @@

class SiPixelDigiErrorsCUDA {
public:
using SiPixelErrorCompactVector = cms::cuda::SimpleVector<SiPixelErrorCompact>;

SiPixelDigiErrorsCUDA() = default;
explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream);
explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream);
~SiPixelDigiErrorsCUDA() = default;

SiPixelDigiErrorsCUDA(const SiPixelDigiErrorsCUDA&) = delete;
SiPixelDigiErrorsCUDA& operator=(const SiPixelDigiErrorsCUDA&) = delete;
SiPixelDigiErrorsCUDA(SiPixelDigiErrorsCUDA&&) = default;
SiPixelDigiErrorsCUDA& operator=(SiPixelDigiErrorsCUDA&&) = default;

const PixelFormatterErrors& formatterErrors() const { return formatterErrors_h; }
const SiPixelFormatterErrors& formatterErrors() const { return formatterErrors_h; }

cms::cuda::SimpleVector<PixelErrorCompact>* error() { return error_d.get(); }
cms::cuda::SimpleVector<PixelErrorCompact> const* error() const { return error_d.get(); }
cms::cuda::SimpleVector<PixelErrorCompact> const* c_error() const { return error_d.get(); }
SiPixelErrorCompactVector* error() { return error_d.get(); }
SiPixelErrorCompactVector const* error() const { return error_d.get(); }

using HostDataError =
std::pair<cms::cuda::SimpleVector<PixelErrorCompact>, cms::cuda::host::unique_ptr<PixelErrorCompact[]>>;
using HostDataError = std::pair<SiPixelErrorCompactVector, cms::cuda::host::unique_ptr<SiPixelErrorCompact[]>>;
HostDataError dataErrorToHostAsync(cudaStream_t stream) const;

void copyErrorToHostAsync(cudaStream_t stream);

private:
cms::cuda::device::unique_ptr<PixelErrorCompact[]> data_d;
cms::cuda::device::unique_ptr<cms::cuda::SimpleVector<PixelErrorCompact>> error_d;
cms::cuda::host::unique_ptr<cms::cuda::SimpleVector<PixelErrorCompact>> error_h;
PixelFormatterErrors formatterErrors_h;
cms::cuda::device::unique_ptr<SiPixelErrorCompact[]> data_d;
cms::cuda::device::unique_ptr<SiPixelErrorCompactVector> error_d;
cms::cuda::host::unique_ptr<SiPixelErrorCompactVector> error_h;
SiPixelFormatterErrors formatterErrors_h;
};

#endif
17 changes: 2 additions & 15 deletions CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
Original file line number Diff line number Diff line change
Expand Up @@ -42,32 +42,19 @@ class SiPixelDigisCUDA {
uint32_t const *pdigi() const { return pdigi_d.get(); }
uint32_t const *rawIdArr() const { return rawIdArr_d.get(); }

uint16_t const *c_xx() const { return xx_d.get(); }
uint16_t const *c_yy() const { return yy_d.get(); }
uint16_t const *c_adc() const { return adc_d.get(); }
uint16_t const *c_moduleInd() const { return moduleInd_d.get(); }
int32_t const *c_clus() const { return clus_d.get(); }
uint32_t const *c_pdigi() const { return pdigi_d.get(); }
uint32_t const *c_rawIdArr() const { return rawIdArr_d.get(); }

cms::cuda::host::unique_ptr<uint16_t[]> adcToHostAsync(cudaStream_t stream) const;
cms::cuda::host::unique_ptr<int32_t[]> clusToHostAsync(cudaStream_t stream) const;
cms::cuda::host::unique_ptr<uint32_t[]> pdigiToHostAsync(cudaStream_t stream) const;
cms::cuda::host::unique_ptr<uint32_t[]> rawIdArrToHostAsync(cudaStream_t stream) const;

class DeviceConstView {
public:
// DeviceConstView() = default;

__device__ __forceinline__ uint16_t xx(int i) const { return __ldg(xx_ + i); }
__device__ __forceinline__ uint16_t yy(int i) const { return __ldg(yy_ + i); }
__device__ __forceinline__ uint16_t adc(int i) const { return __ldg(adc_ + i); }
__device__ __forceinline__ uint16_t moduleInd(int i) const { return __ldg(moduleInd_ + i); }
__device__ __forceinline__ int32_t clus(int i) const { return __ldg(clus_ + i); }

friend class SiPixelDigisCUDA;

// private:
uint16_t const *xx_;
uint16_t const *yy_;
uint16_t const *adc_;
Expand All @@ -88,8 +75,8 @@ class SiPixelDigisCUDA {

// These are for CPU output; should we (eventually) place them to a
// separate product?
cms::cuda::device::unique_ptr<uint32_t[]> pdigi_d;
cms::cuda::device::unique_ptr<uint32_t[]> rawIdArr_d;
cms::cuda::device::unique_ptr<uint32_t[]> pdigi_d; // packed digi (row, col, adc) of each pixel
cms::cuda::device::unique_ptr<uint32_t[]> rawIdArr_d; // DetId of each pixel

uint32_t nModules_h = 0;
uint32_t nDigis_h = 0;
Expand Down
13 changes: 6 additions & 7 deletions CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
Original file line number Diff line number Diff line change
Expand Up @@ -7,14 +7,13 @@

#include <cassert>

SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream)
: formatterErrors_h(std::move(errors)) {
error_d = cms::cuda::make_device_unique<cms::cuda::SimpleVector<PixelErrorCompact>>(stream);
data_d = cms::cuda::make_device_unique<PixelErrorCompact[]>(maxFedWords, stream);

SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream)
: data_d(cms::cuda::make_device_unique<SiPixelErrorCompact[]>(maxFedWords, stream)),
error_d(cms::cuda::make_device_unique<SiPixelErrorCompactVector>(stream)),
error_h(cms::cuda::make_host_unique<SiPixelErrorCompactVector>(stream)),
formatterErrors_h(std::move(errors)) {
cms::cuda::memsetAsync(data_d, 0x00, maxFedWords, stream);

error_h = cms::cuda::make_host_unique<cms::cuda::SimpleVector<PixelErrorCompact>>(stream);
cms::cuda::make_SimpleVector(error_h.get(), maxFedWords, data_d.get());
assert(error_h->empty());
assert(error_h->capacity() == static_cast<int>(maxFedWords));
Expand All @@ -30,7 +29,7 @@ SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync
// On one hand size() could be sufficient. On the other hand, if
// someone copies the SimpleVector<>, (s)he might expect the data
// buffer to actually have space for capacity() elements.
auto data = cms::cuda::make_host_unique<PixelErrorCompact[]>(error_h->capacity(), stream);
auto data = cms::cuda::make_host_unique<SiPixelErrorCompact[]>(error_h->capacity(), stream);

// but transfer only the required amount
if (not error_h->empty()) {
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -84,11 +84,7 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH

// if empy do not bother
if (0 == nHits) {
if
#ifndef __CUDACC__
constexpr
#endif
(std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
cms::cuda::copyAsync(m_view, view, stream);
} else {
m_view.reset(view.release()); // NOLINT: std::move() breaks CUDA version
Expand Down
1 change: 0 additions & 1 deletion CUDADataFormats/TrackingRecHit/src/classes.h
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,6 @@
#define CUDADataFormats_SiPixelCluster_src_classes_h

#include "CUDADataFormats/Common/interface/Product.h"
#include "CUDADataFormats/Common/interface/HostProduct.h"
#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
#include "DataFormats/Common/interface/Wrapper.h"

Expand Down
2 changes: 0 additions & 2 deletions CUDADataFormats/TrackingRecHit/src/classes_def.xml
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,4 @@
<class name="edm::Wrapper<TrackingRecHit2DCPU>" persistent="false"/>
<class name="edm::Wrapper<TrackingRecHit2DHost>" persistent="false"/>
<class name="edm::Wrapper<cms::cuda::Product<TrackingRecHit2DGPU>>" persistent="false"/>
<class name="HostProduct<unsigned int[]>" persistent="false"/>
<class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
</lcgdict>
14 changes: 10 additions & 4 deletions DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,12 @@
#include <cstdint>
#include <vector>

// The main purpose of this class is to deliver digi and cluster data
// from an EDProducer that transfers the data from GPU to host to an
// EDProducer that converts the SoA to legacy data products. The class
// is independent of any GPU technology, and in prunciple could be
// produced by host code, and be used for other purposes than
// conversion-to-legacy as well.
class SiPixelDigisSoA {
public:
SiPixelDigisSoA() = default;
Expand All @@ -24,10 +30,10 @@ class SiPixelDigisSoA {
const std::vector<int32_t>& clusVector() const { return clus_; }

private:
std::vector<uint32_t> pdigi_;
std::vector<uint32_t> rawIdArr_;
std::vector<uint16_t> adc_;
std::vector<int32_t> clus_;
std::vector<uint32_t> pdigi_; // packed digi (row, col, adc) of each pixel
std::vector<uint32_t> rawIdArr_; // DetId of each pixel
std::vector<uint16_t> adc_; // ADC of each pixel
std::vector<int32_t> clus_; // cluster id of each pixel
};

#endif
4 changes: 1 addition & 3 deletions DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,4 @@ SiPixelDigisSoA::SiPixelDigisSoA(
: pdigi_(pdigi, pdigi + nDigis),
rawIdArr_(rawIdArr, rawIdArr + nDigis),
adc_(adc, adc + nDigis),
clus_(clus, clus + nDigis) {
assert(pdigi_.size() == nDigis);
}
clus_(clus, clus + nDigis) {}
1 change: 0 additions & 1 deletion DataFormats/SiPixelDigi/src/classes.h
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,6 @@
#include "DataFormats/SiPixelDigi/interface/SiPixelCalibDigi.h"
#include "DataFormats/SiPixelDigi/interface/SiPixelCalibDigiError.h"
#include "DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h"
#include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h"
#include "DataFormats/Common/interface/Wrapper.h"
#include "DataFormats/Common/interface/DetSetVector.h"
#include "DataFormats/Common/interface/DetSetVectorNew.h"
Expand Down
3 changes: 0 additions & 3 deletions DataFormats/SiPixelDigi/src/classes_def.xml
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,4 @@

<class name="SiPixelDigisSoA" persistent="false"/>
<class name="edm::Wrapper<SiPixelDigisSoA>" persistent="false"/>

<class name="SiPixelDigiErrorsSoA" persistent="false"/>
<class name="edm::Wrapper<SiPixelDigiErrorsSoA>" persistent="false"/>
</lcgdict>
13 changes: 13 additions & 0 deletions DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
#ifndef DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h
#define DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h

#include <cstdint>

struct SiPixelErrorCompact {
uint32_t rawId;
uint32_t word;
uint8_t errorType;
uint8_t fedId;
};

#endif // DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h
30 changes: 30 additions & 0 deletions DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
#ifndef DataFormats_SiPixelDigi_interface_SiPixelErrorsSoA_h
#define DataFormats_SiPixelDigi_interface_SiPixelErrorsSoA_h

#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h"
#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h"

#include <cstdint>
#include <vector>

class SiPixelErrorsSoA {
public:
SiPixelErrorsSoA() = default;
explicit SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err)
: error_(error, error + nErrors), formatterErrors_(err) {}
~SiPixelErrorsSoA() = default;

auto size() const { return error_.size(); }

const SiPixelFormatterErrors *formatterErrors() const { return formatterErrors_; }

const SiPixelErrorCompact &error(size_t i) const { return error_[i]; }

const std::vector<SiPixelErrorCompact> &errorVector() const { return error_; }

private:
std::vector<SiPixelErrorCompact> error_;
const SiPixelFormatterErrors *formatterErrors_ = nullptr;
};

#endif
12 changes: 12 additions & 0 deletions DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
#ifndef DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h
#define DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h

#include <map>
#include <vector>

#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h"
#include "FWCore/Utilities/interface/typedefs.h"

using SiPixelFormatterErrors = std::map<cms_uint32_t, std::vector<SiPixelRawDataError>>;

#endif // DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h
1 change: 1 addition & 0 deletions DataFormats/SiPixelRawData/src/classes.h
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,7 @@
#define SIPIXELRAWDATA_CLASSES_H

#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h"
#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h"
#include "DataFormats/Common/interface/Wrapper.h"
#include "DataFormats/Common/interface/DetSetVector.h"
#include <vector>
Expand Down
3 changes: 3 additions & 0 deletions DataFormats/SiPixelRawData/src/classes_def.xml
Original file line number Diff line number Diff line change
Expand Up @@ -14,4 +14,7 @@
<class name="edm::Wrapper< edm::DetSet<SiPixelRawDataError> >"/>
<class name="edm::Wrapper< std::vector<edm::DetSet<SiPixelRawDataError> > >"/>
<class name="edm::Wrapper< edm::DetSetVector<SiPixelRawDataError> >"/>

<class name="SiPixelErrorsSoA" persistent="false"/>
<class name="edm::Wrapper<SiPixelErrorsSoA>" persistent="false"/>
</lcgdict>
Loading

0 comments on commit 6a08ba1

Please sign in to comment.