diff --git a/CUDADataFormats/Common/BuildFile.xml b/CUDADataFormats/Common/BuildFile.xml
new file mode 100644
index 0000000000000..1046b76eef0f7
--- /dev/null
+++ b/CUDADataFormats/Common/BuildFile.xml
@@ -0,0 +1,7 @@
+
+
+
+
+
+
+
diff --git a/CUDADataFormats/Common/interface/CUDAProduct.h b/CUDADataFormats/Common/interface/CUDAProduct.h
new file mode 100644
index 0000000000000..181024f068c7a
--- /dev/null
+++ b/CUDADataFormats/Common/interface/CUDAProduct.h
@@ -0,0 +1,51 @@
+#ifndef CUDADataFormats_Common_CUDAProduct_h
+#define CUDADataFormats_Common_CUDAProduct_h
+
+#include
+
+#include
+
+#include "CUDADataFormats/Common/interface/CUDAProductBase.h"
+
+namespace edm {
+ template class Wrapper;
+}
+
+/**
+ * The purpose of this class is to wrap CUDA data to edm::Event in a
+ * way which forces correct use of various utilities.
+ *
+ * The non-default construction has to be done with CUDAScopedContext
+ * (in order to properly register the CUDA event).
+ *
+ * The default constructor is needed only for the ROOT dictionary generation.
+ *
+ * The CUDA event is in practice needed only for stream-stream
+ * synchronization, but someone with long-enough lifetime has to own
+ * it. Here is a somewhat natural place. If overhead is too much, we
+ * can e.g. make CUDAService own them (creating them on demand) and
+ * use them only where synchronization between streams is needed.
+ */
+template
+class CUDAProduct: public CUDAProductBase {
+public:
+ CUDAProduct() = default; // Needed only for ROOT dictionary generation
+
+ CUDAProduct(const CUDAProduct&) = delete;
+ CUDAProduct& operator=(const CUDAProduct&) = delete;
+ CUDAProduct(CUDAProduct&&) = default;
+ CUDAProduct& operator=(CUDAProduct&&) = default;
+
+private:
+ friend class CUDAScopedContext;
+ friend class edm::Wrapper>;
+
+ explicit CUDAProduct(int device, std::shared_ptr> stream, std::shared_ptr event, T data):
+ CUDAProductBase(device, std::move(stream), std::move(event)),
+ data_(std::move(data))
+ {}
+
+ T data_; //!
+};
+
+#endif
diff --git a/CUDADataFormats/Common/interface/CUDAProductBase.h b/CUDADataFormats/Common/interface/CUDAProductBase.h
new file mode 100644
index 0000000000000..f54b1c0548ef4
--- /dev/null
+++ b/CUDADataFormats/Common/interface/CUDAProductBase.h
@@ -0,0 +1,41 @@
+#ifndef CUDADataFormats_Common_CUDAProductBase_h
+#define CUDADataFormats_Common_CUDAProductBase_h
+
+#include
+
+#include
+
+/**
+ * Base class for all instantiations of CUDA to hold the
+ * non-T-dependent members.
+ */
+class CUDAProductBase {
+public:
+ CUDAProductBase() = default; // Needed only for ROOT dictionary generation
+
+ bool isValid() const { return stream_.get() != nullptr; }
+ bool isAvailable() const;
+
+ int device() const { return device_; }
+
+ const cuda::stream_t<>& stream() const { return *stream_; }
+ cuda::stream_t<>& stream() { return *stream_; }
+ const std::shared_ptr>& streamPtr() const { return stream_; }
+
+ const cuda::event_t *event() const { return event_.get(); }
+ cuda::event_t *event() { return event_.get(); }
+
+protected:
+ explicit CUDAProductBase(int device, std::shared_ptr> stream, std::shared_ptr event);
+
+private:
+ // The cuda::stream_t is really shared among edm::Event products, so
+ // using shared_ptr also here
+ std::shared_ptr> stream_; //!
+ // shared_ptr because of caching in CUDAService
+ std::shared_ptr event_; //!
+
+ int device_ = -1; //!
+};
+
+#endif
diff --git a/CUDADataFormats/Common/src/CUDAProductBase.cc b/CUDADataFormats/Common/src/CUDAProductBase.cc
new file mode 100644
index 0000000000000..331c4514eb7f7
--- /dev/null
+++ b/CUDADataFormats/Common/src/CUDAProductBase.cc
@@ -0,0 +1,19 @@
+#include "CUDADataFormats/Common/interface/CUDAProductBase.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+
+CUDAProductBase::CUDAProductBase(int device, std::shared_ptr> stream, std::shared_ptr event):
+ stream_(std::move(stream)),
+ event_(std::move(event)),
+ device_(device)
+{}
+
+bool CUDAProductBase::isAvailable() const {
+ // In absence of event, the product was available already at the end
+ // of produce() of the producer.
+ if(not event_) {
+ return true;
+ }
+ return event_->has_occurred();
+}
diff --git a/CUDADataFormats/Common/test/BuildFile.xml b/CUDADataFormats/Common/test/BuildFile.xml
new file mode 100644
index 0000000000000..5e804fe80a736
--- /dev/null
+++ b/CUDADataFormats/Common/test/BuildFile.xml
@@ -0,0 +1,5 @@
+
+
+
+
+
diff --git a/CUDADataFormats/Common/test/test_CUDAProduct.cc b/CUDADataFormats/Common/test/test_CUDAProduct.cc
new file mode 100644
index 0000000000000..308bfe27b29db
--- /dev/null
+++ b/CUDADataFormats/Common/test/test_CUDAProduct.cc
@@ -0,0 +1,69 @@
+#include "catch.hpp"
+
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
+
+#include
+
+namespace cudatest {
+ class TestCUDAScopedContext {
+ public:
+ static
+ CUDAScopedContext make(int dev, bool createEvent) {
+ auto device = cuda::device::get(dev);
+ std::unique_ptr event;
+ if(createEvent) {
+ event = std::make_unique(device.create_event());
+ }
+ return CUDAScopedContext(dev,
+ std::make_unique>(device.create_stream(cuda::stream::implicitly_synchronizes_with_default_stream)),
+ std::move(event));
+ }
+ };
+}
+
+TEST_CASE("Use of CUDAProduct template", "[CUDACore]") {
+ SECTION("Default constructed") {
+ auto foo = CUDAProduct();
+ REQUIRE(!foo.isValid());
+
+ auto bar = std::move(foo);
+ }
+
+ exitSansCUDADevices();
+
+ constexpr int defaultDevice = 0;
+ {
+ auto ctx = cudatest::TestCUDAScopedContext::make(defaultDevice, true);
+ std::unique_ptr> dataPtr = ctx.wrap(10);
+ auto& data = *dataPtr;
+
+ SECTION("Construct from CUDAScopedContext") {
+ REQUIRE(data.isValid());
+ REQUIRE(data.device() == defaultDevice);
+ REQUIRE(data.stream().id() == ctx.stream().id());
+ REQUIRE(data.event() != nullptr);
+ }
+
+ SECTION("Move constructor") {
+ auto data2 = CUDAProduct(std::move(data));
+ REQUIRE(data2.isValid());
+ REQUIRE(!data.isValid());
+ }
+
+ SECTION("Move assignment") {
+ CUDAProduct data2;
+ data2 = std::move(data);
+ REQUIRE(data2.isValid());
+ REQUIRE(!data.isValid());
+ }
+ }
+
+ // Destroy and clean up all resources so that the next test can
+ // assume to start from a clean state.
+ cudaCheck(cudaSetDevice(defaultDevice));
+ cudaCheck(cudaDeviceSynchronize());
+ cudaDeviceReset();
+}
diff --git a/CUDADataFormats/Common/test/test_main.cc b/CUDADataFormats/Common/test/test_main.cc
new file mode 100644
index 0000000000000..0c7c351f437f5
--- /dev/null
+++ b/CUDADataFormats/Common/test/test_main.cc
@@ -0,0 +1,2 @@
+#define CATCH_CONFIG_MAIN
+#include "catch.hpp"
diff --git a/CUDADataFormats/SiPixelCluster/BuildFile.xml b/CUDADataFormats/SiPixelCluster/BuildFile.xml
new file mode 100644
index 0000000000000..d34658faa2573
--- /dev/null
+++ b/CUDADataFormats/SiPixelCluster/BuildFile.xml
@@ -0,0 +1,9 @@
+
+
+
+
+
+
+
+
+
diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
new file mode 100644
index 0000000000000..f25a8a25f0808
--- /dev/null
+++ b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
@@ -0,0 +1,76 @@
+#ifndef CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+#define CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+#include
+
+class SiPixelClustersCUDA {
+public:
+ SiPixelClustersCUDA() = default;
+ explicit SiPixelClustersCUDA(size_t maxClusters, cuda::stream_t<>& stream);
+ ~SiPixelClustersCUDA() = default;
+
+ SiPixelClustersCUDA(const SiPixelClustersCUDA&) = delete;
+ SiPixelClustersCUDA& operator=(const SiPixelClustersCUDA&) = delete;
+ SiPixelClustersCUDA(SiPixelClustersCUDA&&) = default;
+ SiPixelClustersCUDA& operator=(SiPixelClustersCUDA&&) = default;
+
+ void setNClusters(uint32_t nClusters) {
+ nClusters_h = nClusters;
+ }
+
+ uint32_t nClusters() const { return nClusters_h; }
+
+ uint32_t *moduleStart() { return moduleStart_d.get(); }
+ uint32_t *clusInModule() { return clusInModule_d.get(); }
+ uint32_t *moduleId() { return moduleId_d.get(); }
+ uint32_t *clusModuleStart() { return clusModuleStart_d.get(); }
+
+ uint32_t const *moduleStart() const { return moduleStart_d.get(); }
+ uint32_t const *clusInModule() const { return clusInModule_d.get(); }
+ 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;
+
+#ifdef __CUDACC__
+ __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); }
+#endif
+
+ friend SiPixelClustersCUDA;
+
+ private:
+ uint32_t const *moduleStart_;
+ uint32_t const *clusInModule_;
+ uint32_t const *moduleId_;
+ uint32_t const *clusModuleStart_;
+ };
+
+ DeviceConstView *view() const { return view_d.get(); }
+
+private:
+ cudautils::device::unique_ptr moduleStart_d; // index of the first pixel of each module
+ cudautils::device::unique_ptr clusInModule_d; // number of clusters found in each module
+ cudautils::device::unique_ptr moduleId_d; // module id of each module
+
+ // originally from rechits
+ cudautils::device::unique_ptr clusModuleStart_d;
+
+ cudautils::device::unique_ptr view_d; // "me" pointer
+
+ uint32_t nClusters_h;
+};
+
+#endif
diff --git a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
new file mode 100644
index 0000000000000..d88a1b0a6370b
--- /dev/null
+++ b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
@@ -0,0 +1,23 @@
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+
+SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxClusters, cuda::stream_t<>& stream) {
+ edm::Service cs;
+
+ moduleStart_d = cs->make_device_unique(maxClusters+1, stream);
+ clusInModule_d = cs->make_device_unique(maxClusters, stream);
+ moduleId_d = cs->make_device_unique(maxClusters, stream);
+ clusModuleStart_d = cs->make_device_unique(maxClusters+1, stream);
+
+ auto view = cs->make_host_unique(stream);
+ view->moduleStart_ = moduleStart_d.get();
+ view->clusInModule_ = clusInModule_d.get();
+ view->moduleId_ = moduleId_d.get();
+ view->clusModuleStart_ = clusModuleStart_d.get();
+
+ view_d = cs->make_device_unique(stream);
+ cudautils::copyAsync(view_d, view, stream);
+}
diff --git a/CUDADataFormats/SiPixelCluster/src/classes.h b/CUDADataFormats/SiPixelCluster/src/classes.h
new file mode 100644
index 0000000000000..08d46244adc7d
--- /dev/null
+++ b/CUDADataFormats/SiPixelCluster/src/classes.h
@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats_SiPixelCluster_classes_h
+#define CUDADataFormats_SiPixelCluster_classes_h
+
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif
diff --git a/CUDADataFormats/SiPixelCluster/src/classes_def.xml b/CUDADataFormats/SiPixelCluster/src/classes_def.xml
new file mode 100644
index 0000000000000..ba0706ac4b8aa
--- /dev/null
+++ b/CUDADataFormats/SiPixelCluster/src/classes_def.xml
@@ -0,0 +1,4 @@
+
+
+
+
diff --git a/CUDADataFormats/SiPixelDigi/BuildFile.xml b/CUDADataFormats/SiPixelDigi/BuildFile.xml
new file mode 100644
index 0000000000000..29ec13098819c
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/BuildFile.xml
@@ -0,0 +1,9 @@
+
+
+
+
+
+
+
+
+
diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
new file mode 100644
index 0000000000000..e9c8c0f644722
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
@@ -0,0 +1,40 @@
+#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
+#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
+
+#include "DataFormats/SiPixelDigi/interface/PixelErrors.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/GPUSimpleVector.h"
+
+#include
+
+class SiPixelDigiErrorsCUDA {
+public:
+ SiPixelDigiErrorsCUDA() = default;
+ explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cuda::stream_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; }
+
+ GPU::SimpleVector *error() { return error_d.get(); }
+ GPU::SimpleVector const *error() const { return error_d.get(); }
+ GPU::SimpleVector const *c_error() const { return error_d.get(); }
+
+ using HostDataError = std::pair, cudautils::host::unique_ptr>;
+ HostDataError dataErrorToHostAsync(cuda::stream_t<>& stream) const;
+
+ void copyErrorToHostAsync(cuda::stream_t<>& stream);
+
+private:
+ cudautils::device::unique_ptr data_d;
+ cudautils::device::unique_ptr> error_d;
+ cudautils::host::unique_ptr> error_h;
+ PixelFormatterErrors formatterErrors_h;
+};
+
+#endif
diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
new file mode 100644
index 0000000000000..6a52545483eb8
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
@@ -0,0 +1,99 @@
+#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+#include
+
+class SiPixelDigisCUDA {
+public:
+ SiPixelDigisCUDA() = default;
+ explicit SiPixelDigisCUDA(size_t maxFedWords, cuda::stream_t<>& stream);
+ ~SiPixelDigisCUDA() = default;
+
+ SiPixelDigisCUDA(const SiPixelDigisCUDA&) = delete;
+ SiPixelDigisCUDA& operator=(const SiPixelDigisCUDA&) = delete;
+ SiPixelDigisCUDA(SiPixelDigisCUDA&&) = default;
+ SiPixelDigisCUDA& operator=(SiPixelDigisCUDA&&) = default;
+
+ void setNModulesDigis(uint32_t nModules, uint32_t nDigis) {
+ nModules_h = nModules;
+ nDigis_h = nDigis;
+ }
+
+ uint32_t nModules() const { return nModules_h; }
+ uint32_t nDigis() const { return nDigis_h; }
+
+ uint16_t * xx() { return xx_d.get(); }
+ uint16_t * yy() { return yy_d.get(); }
+ uint16_t * adc() { return adc_d.get(); }
+ uint16_t * moduleInd() { return moduleInd_d.get(); }
+ int32_t * clus() { return clus_d.get(); }
+ uint32_t * pdigi() { return pdigi_d.get(); }
+ uint32_t * rawIdArr() { return rawIdArr_d.get(); }
+
+ uint16_t const *xx() const { return xx_d.get(); }
+ uint16_t const *yy() const { return yy_d.get(); }
+ uint16_t const *adc() const { return adc_d.get(); }
+ uint16_t const *moduleInd() const { return moduleInd_d.get(); }
+ int32_t const *clus() const { return clus_d.get(); }
+ 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(); }
+
+ cudautils::host::unique_ptr adcToHostAsync(cuda::stream_t<>& stream) const;
+ cudautils::host::unique_ptr< int32_t[]> clusToHostAsync(cuda::stream_t<>& stream) const;
+ cudautils::host::unique_ptr pdigiToHostAsync(cuda::stream_t<>& stream) const;
+ cudautils::host::unique_ptr rawIdArrToHostAsync(cuda::stream_t<>& stream) const;
+
+ class DeviceConstView {
+ public:
+ DeviceConstView() = default;
+
+#ifdef __CUDACC__
+ __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); }
+#endif
+
+ friend class SiPixelDigisCUDA;
+
+ private:
+ uint16_t const *xx_;
+ uint16_t const *yy_;
+ uint16_t const *adc_;
+ uint16_t const *moduleInd_;
+ int32_t const *clus_;
+ };
+
+ const DeviceConstView *view() const { return view_d.get(); }
+
+private:
+ // These are consumed by downstream device code
+ cudautils::device::unique_ptr xx_d; // local coordinates of each pixel
+ cudautils::device::unique_ptr yy_d; //
+ cudautils::device::unique_ptr adc_d; // ADC of each pixel
+ cudautils::device::unique_ptr moduleInd_d; // module id of each pixel
+ cudautils::device::unique_ptr clus_d; // cluster id of each pixel
+ cudautils::device::unique_ptr view_d; // "me" pointer
+
+ // These are for CPU output; should we (eventually) place them to a
+ // separate product?
+ cudautils::device::unique_ptr pdigi_d;
+ cudautils::device::unique_ptr rawIdArr_d;
+
+ uint32_t nModules_h = 0;
+ uint32_t nDigis_h = 0;
+};
+
+#endif
diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
new file mode 100644
index 0000000000000..92aab1ec9d578
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
@@ -0,0 +1,44 @@
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/memsetAsync.h"
+
+SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cuda::stream_t<>& stream):
+ formatterErrors_h(std::move(errors))
+{
+ edm::Service cs;
+
+ error_d = cs->make_device_unique>(stream);
+ data_d = cs->make_device_unique(maxFedWords, stream);
+
+ cudautils::memsetAsync(data_d, 0x00, maxFedWords, stream);
+
+ error_h = cs->make_host_unique>(stream);
+ GPU::make_SimpleVector(error_h.get(), maxFedWords, data_d.get());
+ assert(error_h->size() == 0);
+ assert(error_h->capacity() == static_cast(maxFedWords));
+
+ cudautils::copyAsync(error_d, error_h, stream);
+}
+
+void SiPixelDigiErrorsCUDA::copyErrorToHostAsync(cuda::stream_t<>& stream) {
+ cudautils::copyAsync(error_h, error_d, stream);
+}
+
+SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync(cuda::stream_t<>& stream) const {
+ edm::Service cs;
+ // 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 = cs->make_host_unique(error_h->capacity(), stream);
+
+ // but transfer only the required amount
+ if(error_h->size() > 0) {
+ cudautils::copyAsync(data, data_d, error_h->size(), stream);
+ }
+ auto err = *error_h;
+ err.set_data(data.get());
+ return HostDataError(std::move(err), std::move(data));
+}
diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc
new file mode 100644
index 0000000000000..ef13ed9612dbf
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc
@@ -0,0 +1,56 @@
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+
+SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cuda::stream_t<>& stream) {
+ edm::Service cs;
+
+ xx_d = cs->make_device_unique(maxFedWords, stream);
+ yy_d = cs->make_device_unique(maxFedWords, stream);
+ adc_d = cs->make_device_unique(maxFedWords, stream);
+ moduleInd_d = cs->make_device_unique(maxFedWords, stream);
+ clus_d = cs->make_device_unique< int32_t[]>(maxFedWords, stream);
+
+ pdigi_d = cs->make_device_unique(maxFedWords, stream);
+ rawIdArr_d = cs->make_device_unique(maxFedWords, stream);
+
+ auto view = cs->make_host_unique(stream);
+ view->xx_ = xx_d.get();
+ view->yy_ = yy_d.get();
+ view->adc_ = adc_d.get();
+ view->moduleInd_ = moduleInd_d.get();
+ view->clus_ = clus_d.get();
+
+ view_d = cs->make_device_unique(stream);
+ cudautils::copyAsync(view_d, view, stream);
+}
+
+cudautils::host::unique_ptr SiPixelDigisCUDA::adcToHostAsync(cuda::stream_t<>& stream) const {
+ edm::Service cs;
+ auto ret = cs->make_host_unique(nDigis(), stream);
+ cudautils::copyAsync(ret, adc_d, nDigis(), stream);
+ return ret;
+}
+
+cudautils::host::unique_ptr SiPixelDigisCUDA::clusToHostAsync(cuda::stream_t<>& stream) const {
+ edm::Service cs;
+ auto ret = cs->make_host_unique(nDigis(), stream);
+ cudautils::copyAsync(ret, clus_d, nDigis(), stream);
+ return ret;
+}
+
+cudautils::host::unique_ptr SiPixelDigisCUDA::pdigiToHostAsync(cuda::stream_t<>& stream) const {
+ edm::Service cs;
+ auto ret = cs->make_host_unique(nDigis(), stream);
+ cudautils::copyAsync(ret, pdigi_d, nDigis(), stream);
+ return ret;
+}
+
+cudautils::host::unique_ptr SiPixelDigisCUDA::rawIdArrToHostAsync(cuda::stream_t<>& stream) const {
+ edm::Service cs;
+ auto ret = cs->make_host_unique(nDigis(), stream);
+ cudautils::copyAsync(ret, rawIdArr_d, nDigis(), stream);
+ return ret;
+}
diff --git a/CUDADataFormats/SiPixelDigi/src/classes.h b/CUDADataFormats/SiPixelDigi/src/classes.h
new file mode 100644
index 0000000000000..41b135640b883
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/src/classes.h
@@ -0,0 +1,9 @@
+#ifndef CUDADataFormats_SiPixelDigi_classes_h
+#define CUDADataFormats_SiPixelDigi_classes_h
+
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif
diff --git a/CUDADataFormats/SiPixelDigi/src/classes_def.xml b/CUDADataFormats/SiPixelDigi/src/classes_def.xml
new file mode 100644
index 0000000000000..9d6816ed3b14c
--- /dev/null
+++ b/CUDADataFormats/SiPixelDigi/src/classes_def.xml
@@ -0,0 +1,7 @@
+
+
+
+
+
+
+
diff --git a/CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h b/CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h
new file mode 100644
index 0000000000000..afb682e5d451f
--- /dev/null
+++ b/CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h
@@ -0,0 +1,14 @@
+#ifndef CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+#define CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+
+#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"
+#include "FWCore/Framework/interface/DependentRecordImplementation.h"
+
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include "boost/mpl/vector.hpp"
+
+class SiPixelGainCalibrationForHLTGPURcd : public edm::eventsetup::DependentRecordImplementation > {};
+
+#endif
diff --git a/CalibTracker/Records/src/SiPixelGainCalibrationForHLTGPURcd.cc b/CalibTracker/Records/src/SiPixelGainCalibrationForHLTGPURcd.cc
new file mode 100644
index 0000000000000..e6020eca80b1f
--- /dev/null
+++ b/CalibTracker/Records/src/SiPixelGainCalibrationForHLTGPURcd.cc
@@ -0,0 +1,5 @@
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+EVENTSETUP_RECORD_REG(SiPixelGainCalibrationForHLTGPURcd);
diff --git a/CalibTracker/SiPixelESProducers/BuildFile.xml b/CalibTracker/SiPixelESProducers/BuildFile.xml
index e9d22b32f0afb..69d258da21ed1 100644
--- a/CalibTracker/SiPixelESProducers/BuildFile.xml
+++ b/CalibTracker/SiPixelESProducers/BuildFile.xml
@@ -7,7 +7,9 @@
+
+
diff --git a/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h b/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h
new file mode 100644
index 0000000000000..96989c8a2c3b2
--- /dev/null
+++ b/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h
@@ -0,0 +1,32 @@
+#ifndef CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+#define CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+
+#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+
+#include
+
+class SiPixelGainCalibrationForHLT;
+class SiPixelGainForHLTonGPU;
+struct SiPixelGainForHLTonGPU_DecodingStructure;
+class TrackerGeometry;
+
+class SiPixelGainCalibrationForHLTGPU {
+public:
+ explicit SiPixelGainCalibrationForHLTGPU(const SiPixelGainCalibrationForHLT& gains, const TrackerGeometry& geom);
+ ~SiPixelGainCalibrationForHLTGPU();
+
+ const SiPixelGainForHLTonGPU *getGPUProductAsync(cuda::stream_t<>& cudaStream) const;
+
+private:
+ const SiPixelGainCalibrationForHLT *gains_ = nullptr;
+ SiPixelGainForHLTonGPU *gainForHLTonHost_ = nullptr;
+ struct GPUData {
+ ~GPUData();
+ SiPixelGainForHLTonGPU *gainForHLTonGPU = nullptr;
+ SiPixelGainForHLTonGPU_DecodingStructure *gainDataOnGPU = nullptr;
+ };
+ CUDAESProduct gpuData_;
+};
+
+#endif
diff --git a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml
index 44db9d9ba0582..b33657e273036 100644
--- a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml
+++ b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml
@@ -6,6 +6,8 @@
+
+
diff --git a/CalibTracker/SiPixelESProducers/plugins/SiPixelGainCalibrationForHLTGPUESProducer.cc b/CalibTracker/SiPixelESProducers/plugins/SiPixelGainCalibrationForHLTGPUESProducer.cc
new file mode 100644
index 0000000000000..186bb2d72c3f3
--- /dev/null
+++ b/CalibTracker/SiPixelESProducers/plugins/SiPixelGainCalibrationForHLTGPUESProducer.cc
@@ -0,0 +1,47 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "FWCore/Framework/interface/ESProducer.h"
+#include "FWCore/Framework/interface/EventSetup.h"
+#include "FWCore/Framework/interface/ESHandle.h"
+#include "FWCore/Framework/interface/ModuleFactory.h"
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include
+
+class SiPixelGainCalibrationForHLTGPUESProducer: public edm::ESProducer {
+public:
+ explicit SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig);
+ std::unique_ptr produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord);
+
+ static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
+private:
+};
+
+SiPixelGainCalibrationForHLTGPUESProducer::SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig) {
+ setWhatProduced(this);
+}
+
+void SiPixelGainCalibrationForHLTGPUESProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
+ edm::ParameterSetDescription desc;
+ descriptions.add("siPixelGainCalibrationForHLTGPU", desc);
+}
+
+std::unique_ptr SiPixelGainCalibrationForHLTGPUESProducer::produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord) {
+ edm::ESHandle gains;
+ iRecord.getRecord().get(gains);
+
+ edm::ESHandle geom;
+ iRecord.getRecord().get(geom);
+
+ return std::make_unique(*gains, *geom);
+}
+
+#include "FWCore/Framework/interface/MakerMacros.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+
+DEFINE_FWK_EVENTSETUP_MODULE(SiPixelGainCalibrationForHLTGPUESProducer);
diff --git a/CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTGPU.cc b/CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTGPU.cc
new file mode 100644
index 0000000000000..80932fb468f71
--- /dev/null
+++ b/CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTGPU.cc
@@ -0,0 +1,4 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+TYPELOOKUP_DATA_REG(SiPixelGainCalibrationForHLTGPU);
diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc
new file mode 100644
index 0000000000000..3aef3f44c8f67
--- /dev/null
+++ b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc
@@ -0,0 +1,98 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "Geometry/CommonDetUnit/interface/GeomDetType.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+
+#include
+
+SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGainCalibrationForHLT& gains, const TrackerGeometry& geom):
+ gains_(&gains)
+{
+ // bizzarre logic (looking for fist strip-det) don't ask
+ auto const & dus = geom.detUnits();
+ unsigned m_detectors = dus.size();
+ for(unsigned int i=1;i<7;++i) {
+ if(geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) != dus.size() &&
+ dus[geom.offsetDU(GeomDetEnumerators::tkDetEnum[i])]->type().isTrackerStrip()) {
+ if(geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) < m_detectors) m_detectors = geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]);
+ }
+ }
+
+ /*
+ std::cout << "caching calibs for " << m_detectors << " pixel detectors of size " << gains.data().size() << std::endl;
+ std::cout << "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' ' << sizeof(SiPixelGainForHLTonGPU::DecodingStructure) << std::endl;
+ */
+
+ cudaCheck(cudaMallocHost((void**) & gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU)));
+ //gainForHLTonHost_->v_pedestals = gainDataOnGPU_; // how to do this?
+
+ // do not read back from the (possibly write-combined) memory buffer
+ auto minPed = gains.getPedLow();
+ auto maxPed = gains.getPedHigh();
+ auto minGain = gains.getGainLow();
+ auto maxGain = gains.getGainHigh();
+ auto nBinsToUseForEncoding = 253;
+
+ // we will simplify later (not everything is needed....)
+ gainForHLTonHost_->minPed_ = minPed;
+ gainForHLTonHost_->maxPed_ = maxPed;
+ gainForHLTonHost_->minGain_= minGain;
+ gainForHLTonHost_->maxGain_= maxGain;
+
+ gainForHLTonHost_->numberOfRowsAveragedOver_ = 80;
+ gainForHLTonHost_->nBinsToUseForEncoding_ = nBinsToUseForEncoding;
+ gainForHLTonHost_->deadFlag_ = 255;
+ gainForHLTonHost_->noisyFlag_ = 254;
+
+ gainForHLTonHost_->pedPrecision = static_cast(maxPed - minPed) / nBinsToUseForEncoding;
+ gainForHLTonHost_->gainPrecision = static_cast(maxGain - minGain) / nBinsToUseForEncoding;
+
+ /*
+ std::cout << "precisions g " << gainForHLTonHost_->pedPrecision << ' ' << gainForHLTonHost_->gainPrecision << std::endl;
+ */
+
+ // fill the index map
+ auto const & ind = gains.getIndexes();
+ /*
+ std::cout << ind.size() << " " << m_detectors << std::endl;
+ */
+
+ for (auto i=0U; igeographicalId().rawId(),SiPixelGainCalibrationForHLT::StrictWeakOrdering());
+ assert (p!=ind.end() && p->detid==dus[i]->geographicalId());
+ assert(p->iend<=gains.data().size());
+ assert(p->iend>=p->ibegin);
+ assert(0==p->ibegin%2);
+ assert(0==p->iend%2);
+ assert(p->ibegin!=p->iend);
+ assert(p->ncols>0);
+ gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(p->ibegin,p->iend), p->ncols);
+ // if (ind[i].detid!=dus[i]->geographicalId()) std::cout << ind[i].detid<<"!="<geographicalId() << std::endl;
+ // gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(ind[i].ibegin,ind[i].iend), ind[i].ncols);
+ }
+
+}
+
+SiPixelGainCalibrationForHLTGPU::~SiPixelGainCalibrationForHLTGPU() {
+ cudaCheck(cudaFreeHost(gainForHLTonHost_));
+}
+
+SiPixelGainCalibrationForHLTGPU::GPUData::~GPUData() {
+ cudaCheck(cudaFree(gainForHLTonGPU));
+ cudaCheck(cudaFree(gainDataOnGPU));
+}
+
+const SiPixelGainForHLTonGPU *SiPixelGainCalibrationForHLTGPU::getGPUProductAsync(cuda::stream_t<>& cudaStream) const {
+ const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, [this](GPUData& data, cuda::stream_t<>& stream) {
+ cudaCheck(cudaMalloc((void**) & data.gainForHLTonGPU, sizeof(SiPixelGainForHLTonGPU)));
+ cudaCheck(cudaMalloc((void**) & data.gainDataOnGPU, this->gains_->data().size())); // TODO: this could be changed to cuda::memory::device::unique_ptr<>
+ // gains.data().data() is used also for non-GPU code, we cannot allocate it on aligned and write-combined memory
+ cudaCheck(cudaMemcpyAsync(data.gainDataOnGPU, this->gains_->data().data(), this->gains_->data().size(), cudaMemcpyDefault, stream.id()));
+
+ cudaCheck(cudaMemcpyAsync(data.gainForHLTonGPU, this->gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU), cudaMemcpyDefault, stream.id()));
+ cudaCheck(cudaMemcpyAsync(&(data.gainForHLTonGPU->v_pedestals), &(data.gainDataOnGPU), sizeof(SiPixelGainForHLTonGPU_DecodingStructure*), cudaMemcpyDefault, stream.id()));
+ });
+ return data.gainForHLTonGPU;
+}
diff --git a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h
new file mode 100644
index 0000000000000..931ee7e65f295
--- /dev/null
+++ b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h
@@ -0,0 +1,73 @@
+#ifndef CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+#define CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+
+#include
+#include
+#include
+
+#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
+
+struct SiPixelGainForHLTonGPU_DecodingStructure{
+ uint8_t gain;
+ uint8_t ped;
+};
+
+
+// copy of SiPixelGainCalibrationForHLT
+class SiPixelGainForHLTonGPU {
+
+ public:
+
+ using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;
+
+ using Range = std::pair;
+
+
+ inline __host__ __device__
+ std::pair getPedAndGain(uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn ) const {
+
+
+ auto range = rangeAndCols[moduleInd].first;
+ auto nCols = rangeAndCols[moduleInd].second;
+
+ // determine what averaged data block we are in (there should be 1 or 2 of these depending on if plaquette is 1 by X or 2 by X
+ unsigned int lengthOfColumnData = (range.second-range.first)/nCols;
+ unsigned int lengthOfAveragedDataInEachColumn = 2; // we always only have two values per column averaged block
+ unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;
+
+
+ auto offset = range.first + col*lengthOfColumnData + lengthOfAveragedDataInEachColumn*numberOfDataBlocksToSkip;
+
+ assert(offset rangeAndCols[2000];
+
+ float minPed_, maxPed_, minGain_, maxGain_;
+
+ float pedPrecision, gainPrecision;
+
+ unsigned int numberOfRowsAveragedOver_; // this is 80!!!!
+ unsigned int nBinsToUseForEncoding_;
+ unsigned int deadFlag_;
+ unsigned int noisyFlag_;
+};
+
+#endif // CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
diff --git a/Configuration/Applications/python/ConfigBuilder.py b/Configuration/Applications/python/ConfigBuilder.py
index ea36a1cdeeecb..5dd2915049c3d 100644
--- a/Configuration/Applications/python/ConfigBuilder.py
+++ b/Configuration/Applications/python/ConfigBuilder.py
@@ -920,6 +920,8 @@ def define_Configs(self):
self.loadAndRemember('SimGeneral.HepPDTESSource.'+self._options.particleTable+'_cfi')
self.loadAndRemember('FWCore/MessageService/MessageLogger_cfi')
+ # Eventually replace with some more generic file to load
+ self.loadAndRemember('HeterogeneousCore/CUDAServices/CUDAService_cfi')
self.ALCADefaultCFF="Configuration/StandardSequences/AlCaRecoStreams_cff"
self.GENDefaultCFF="Configuration/StandardSequences/Generator_cff"
diff --git a/Configuration/ProcessModifiers/python/gpu_cff.py b/Configuration/ProcessModifiers/python/gpu_cff.py
new file mode 100644
index 0000000000000..993f71804fbc1
--- /dev/null
+++ b/Configuration/ProcessModifiers/python/gpu_cff.py
@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing CPU modules with GPU counterparts
+
+gpu = cms.Modifier()
diff --git a/Configuration/ProcessModifiers/python/riemannFitGPU_cff.py b/Configuration/ProcessModifiers/python/riemannFitGPU_cff.py
new file mode 100644
index 0000000000000..ef622f26b2da0
--- /dev/null
+++ b/Configuration/ProcessModifiers/python/riemannFitGPU_cff.py
@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing the default pixel track "fitting" with Riemann fit on GPU
+
+riemannFitGPU = cms.Modifier()
diff --git a/Configuration/ProcessModifiers/python/riemannFit_cff.py b/Configuration/ProcessModifiers/python/riemannFit_cff.py
new file mode 100644
index 0000000000000..f97f50df63fb6
--- /dev/null
+++ b/Configuration/ProcessModifiers/python/riemannFit_cff.py
@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing the default pixel track "fitting" with Riemann fit
+
+riemannFit = cms.Modifier()
diff --git a/Configuration/PyReleaseValidation/python/relval_2017.py b/Configuration/PyReleaseValidation/python/relval_2017.py
index 1891cac97950a..aeaeced9bb1c7 100644
--- a/Configuration/PyReleaseValidation/python/relval_2017.py
+++ b/Configuration/PyReleaseValidation/python/relval_2017.py
@@ -29,7 +29,7 @@
10024.1,10024.2,10024.3,10024.4,10024.5,
10801.0,10802.0,10803.0,10804.0,10805.0,10806.0,10807.0,10808.0,10809.0,10859.0,10871.0,
10842.0,10824.0,10825.0,10826.0,10823.0,11024.0,11025.0,11224.0,
- 10824.1,10824.5,
+ 10824.1,10824.5,10824.51,10824.52,10824.53,
10824.6,11024.6,11224.6,
10824.7,10842.7,
10824.8,
diff --git a/Configuration/PyReleaseValidation/python/relval_standard.py b/Configuration/PyReleaseValidation/python/relval_standard.py
index c2768582ae616..55e42293de8db 100644
--- a/Configuration/PyReleaseValidation/python/relval_standard.py
+++ b/Configuration/PyReleaseValidation/python/relval_standard.py
@@ -349,6 +349,9 @@
workflows[136.862] = ['',['RunEGamma2018B','HLTDR2_2018','RECODR2_2018reHLT_skimEGamma_Prompt_L1TEgDQM','HARVEST2018_L1TEgDQM']]
workflows[136.863] = ['',['RunDoubleMuon2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt','HARVEST2018']]
workflows[136.864] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_skimJetHT_Prompt','HARVEST2018']]
+workflows[136.8645] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt_pixelTrackingOnly','HARVEST2018_pixelTrackingOnly']]
+workflows[136.86452] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt_pixelTrackingOnlyGPU','HARVEST2018_pixelTrackingOnly']]
+
workflows[136.865] = ['',['RunMET2018B','HLTDR2_2018','RECODR2_2018reHLT_skimMET_Prompt','HARVEST2018']]
workflows[136.866] = ['',['RunMuonEG2018B','HLTDR2_2018','RECODR2_2018reHLT_skimMuonEG_Prompt','HARVEST2018']]
workflows[136.867] = ['',['RunSingleMu2018B','HLTDR2_2018','RECODR2_2018reHLT_skimSingleMu_Prompt_Lumi','HARVEST2018_L1TMuDQM']]
diff --git a/Configuration/PyReleaseValidation/python/relval_steps.py b/Configuration/PyReleaseValidation/python/relval_steps.py
index b74222f392fa0..54b6dc4475c74 100644
--- a/Configuration/PyReleaseValidation/python/relval_steps.py
+++ b/Configuration/PyReleaseValidation/python/relval_steps.py
@@ -2047,6 +2047,15 @@ def gen2018HiMix(fragment,howMuch):
'--datatier': 'GEN-SIM-RECO,DQMIO',
'--eventcontent': 'RECOSIM,DQM',
}
+step3_riemannFit = {
+ '--procModifiers': 'riemannFit',
+}
+step3_riemannFitGPU = {
+ '--procModifiers': 'riemannFitGPU',
+}
+step3_gpu = {
+ '--procModifiers': 'gpu',
+}
step3_trackingLowPU = {
'--era': 'Run2_2016_trackingLowPU'
}
@@ -2171,6 +2180,8 @@ def gen2018HiMix(fragment,howMuch):
steps['RECODR2_2017reHLTSiPixelCalZeroBias_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2017reHLTSiPixelCalZeroBias']])
steps['RECODR2_2018reHLT_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT']])
+steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnly']=merge([{'-s': 'RAW2DIGI:RawToDigi_pixelOnly,RECO:reconstruction_pixelTrackingOnly,DQM:@pixelTrackingOnlyDQM'},steps['RECODR2_2018reHLT_Prompt']])
+steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnlyGPU']=merge([step3_gpu, steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnly']])
steps['RECODR2_2018reHLT_skimEGamma_Prompt_L1TEgDQM']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimEGamma_L1TEgDQM']])
steps['RECODR2_2018reHLT_skimJetHT_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimJetHT']])
steps['RECODR2_2018reHLT_skimDisplacedJet_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimDisplacedJet']])
@@ -2471,6 +2482,7 @@ def gen2018HiMix(fragment,howMuch):
steps['HARVEST2018'] = merge([ {'--conditions':'auto:run2_data_relval','--era':'Run2_2018','--conditions':'auto:run2_data_promptlike',}, steps['HARVESTD'] ])
steps['HARVEST2018_L1TEgDQM'] = merge([ {'-s':'HARVESTING:@standardDQM+@ExtraHLT+@miniAODDQM+@L1TEgamma'}, steps['HARVEST2018'] ])
steps['HARVEST2018_L1TMuDQM'] = merge([ {'-s':'HARVESTING:@standardDQM+@ExtraHLT+@miniAODDQM+@L1TMuon'}, steps['HARVEST2018'] ])
+steps['HARVEST2018_pixelTrackingOnly'] = merge([ {'-s':'HARVESTING:@pixelTrackingOnlyDQM'}, steps['HARVEST2018'] ])
steps['HARVEST2018_hBStar'] = merge([ {'--era' : 'Run2_2018_highBetaStar'}, steps['HARVEST2018'] ])
steps['HARVEST2018_HEfail'] = merge([ {'--conditions':'auto:run2_data_promptlike_HEfail'}, steps['HARVEST2018'] ])
steps['HARVEST2018_BadHcalMitig'] = merge([ {'--era' : 'Run2_2018,pf_badHcalMitigation','--conditions':'auto:run2_data_promptlike_HEfail'}, steps['HARVEST2018'] ])
@@ -3124,6 +3136,21 @@ def gen2018HiMix(fragment,howMuch):
if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_pixelTrackingOnly, upgradeStepDict[step][k]])
elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+ for step in upgradeSteps['pixelTrackingOnlyRiemannFit']['steps']:
+ stepName = step + upgradeSteps['pixelTrackingOnlyRiemannFit']['suffix']
+ if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_riemannFit, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+ elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
+ for step in upgradeSteps['pixelTrackingOnlyRiemannFitGPU']['steps']:
+ stepName = step + upgradeSteps['pixelTrackingOnlyRiemannFitGPU']['suffix']
+ if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_riemannFitGPU, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+ elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
+ for step in upgradeSteps['pixelTrackingOnlyGPU']['steps']:
+ stepName = step + upgradeSteps['pixelTrackingOnlyGPU']['suffix']
+ if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_gpu, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+ elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
for step in upgradeSteps['trackingRun2']['steps']:
stepName = step + upgradeSteps['trackingRun2']['suffix']
if 'Reco' in step and upgradeStepDict[step][k]['--era']=='Run2_2017':
diff --git a/Configuration/PyReleaseValidation/python/relval_upgrade.py b/Configuration/PyReleaseValidation/python/relval_upgrade.py
index 6554343efcac3..8c182549affa3 100644
--- a/Configuration/PyReleaseValidation/python/relval_upgrade.py
+++ b/Configuration/PyReleaseValidation/python/relval_upgrade.py
@@ -5,7 +5,7 @@
# here only define the workflows as a combination of the steps defined above:
workflows = Matrix()
-# each workflow defines a name and a list of steps to be done.
+# each workflow defines a name and a list of steps to be done.
# if no explicit name/label given for the workflow (first arg),
# the name of step1 will be used
@@ -30,13 +30,13 @@ def makeStepName(key,frag,step,suffix):
for stepType in upgradeSteps.keys():
stepList[stepType] = []
hasHarvest = False
- for step in upgradeProperties[year][key]['ScenToRun']:
+ for step in upgradeProperties[year][key]['ScenToRun']:
stepMaker = makeStepName
if 'Sim' in step:
if 'HLBeamSpotFull' in step and '14TeV' in frag:
step = 'GenSimHLBeamSpotFull14'
stepMaker = makeStepNameSim
-
+
if 'HARVEST' in step: hasHarvest = True
for stepType in upgradeSteps.keys():
@@ -79,7 +79,7 @@ def makeStepName(key,frag,step,suffix):
# special workflows for tracker
if (upgradeDatasetFromFragment[frag]=="TTbar_13" or upgradeDatasetFromFragment[frag]=="TTbar_14TeV") and not 'PU' in key and hasHarvest:
# skip ALCA and Nano
- trackingVariations = ['trackingOnly','trackingRun2','trackingOnlyRun2','trackingLowPU','pixelTrackingOnly']
+ trackingVariations = ['trackingOnly','trackingRun2','trackingOnlyRun2','trackingLowPU','pixelTrackingOnly','pixelTrackingOnlyRiemannFit','pixelTrackingOnlyRiemannFitGPU','pixelTrackingOnlyGPU']
for tv in trackingVariations:
stepList[tv] = [s for s in stepList[tv] if (("ALCA" not in s) and ("Nano" not in s))]
workflows[numWF+upgradeSteps['trackingOnly']['offset']] = [ upgradeDatasetFromFragment[frag], stepList['trackingOnly']]
@@ -87,7 +87,10 @@ def makeStepName(key,frag,step,suffix):
for tv in trackingVariations[1:]:
workflows[numWF+upgradeSteps[tv]['offset']] = [ upgradeDatasetFromFragment[frag], stepList[tv]]
elif '2018' in key:
- workflows[numWF+upgradeSteps['pixelTrackingOnly']['offset']] = [ upgradeDatasetFromFragment[frag], stepList['pixelTrackingOnly']]
+ for tv in trackingVariations:
+ if not "pixelTrackingOnly" in tv:
+ continue
+ workflows[numWF+upgradeSteps[tv]['offset']] = [ upgradeDatasetFromFragment[frag], stepList[tv]]
# special workflows for HE
if upgradeDatasetFromFragment[frag]=="TTbar_13" and '2018' in key:
diff --git a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py
index 4bdbb2130e311..a5f9742bc51a4 100644
--- a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py
+++ b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py
@@ -158,6 +158,39 @@
'suffix' : '_pixelTrackingOnly',
'offset' : 0.5,
}
+upgradeSteps['pixelTrackingOnlyRiemannFit'] = {
+ 'steps' : [
+ 'RecoFull',
+ 'HARVESTFull',
+ 'RecoFullGlobal',
+ 'HARVESTFullGlobal',
+ ],
+ 'PU' : [],
+ 'suffix' : '_pixelTrackingOnlyRiemannFit',
+ 'offset' : 0.51,
+}
+upgradeSteps['pixelTrackingOnlyGPU'] = {
+ 'steps' : [
+ 'RecoFull',
+ 'HARVESTFull',
+ 'RecoFullGlobal',
+ 'HARVESTFullGlobal',
+ ],
+ 'PU' : [],
+ 'suffix' : '_pixelTrackingOnlyGPU',
+ 'offset' : 0.52,
+}
+upgradeSteps['pixelTrackingOnlyRiemannFitGPU'] = {
+ 'steps' : [
+ 'RecoFull',
+ 'HARVESTFull',
+ 'RecoFullGlobal',
+ 'HARVESTFullGlobal',
+ ],
+ 'PU' : [],
+ 'suffix' : '_pixelTrackingOnlyRiemannFitGPU',
+ 'offset' : 0.53,
+}
upgradeSteps['Timing'] = {
'steps' : upgradeSteps['baseline']['steps'],
'PU' : upgradeSteps['baseline']['PU'],
@@ -488,7 +521,7 @@
'DoubleMuPt1000Extended_pythia8_cfi',
'TenMuE_0_200_pythia8_cfi',
'SinglePiE50HCAL_pythia8_cfi',
- 'MinBias_13TeV_pythia8_TuneCUETP8M1_cfi',
+ 'MinBias_13TeV_pythia8_TuneCUETP8M1_cfi',
'TTbar_13TeV_TuneCUETP8M1_cfi',
'ZEE_13TeV_TuneCUETP8M1_cfi',
'QCD_Pt_600_800_13TeV_TuneCUETP8M1_cfi',
diff --git a/Configuration/StandardSequences/python/RawToDigi_cff.py b/Configuration/StandardSequences/python/RawToDigi_cff.py
index 35e57a29778b2..5d3393d181ed1 100644
--- a/Configuration/StandardSequences/python/RawToDigi_cff.py
+++ b/Configuration/StandardSequences/python/RawToDigi_cff.py
@@ -3,7 +3,7 @@
# This object is used to selectively make changes for different running
# scenarios. In this case it makes changes for Run 2.
-from EventFilter.SiPixelRawToDigi.SiPixelRawToDigi_cfi import *
+from EventFilter.SiPixelRawToDigi.siPixelDigis_cff import *
from EventFilter.SiStripRawToDigi.SiStripDigis_cfi import *
@@ -49,7 +49,7 @@
from EventFilter.CTPPSRawToDigi.ctppsRawToDigi_cff import *
RawToDigiTask = cms.Task(L1TRawToDigiTask,
- siPixelDigis,
+ siPixelDigisTask,
siStripDigis,
ecalDigis,
ecalPreshowerDigis,
@@ -64,14 +64,14 @@
)
RawToDigi = cms.Sequence(RawToDigiTask)
-RawToDigiTask_noTk = RawToDigiTask.copyAndExclude([siPixelDigis, siStripDigis])
+RawToDigiTask_noTk = RawToDigiTask.copyAndExclude([siPixelDigisTask, siStripDigis])
RawToDigi_noTk = cms.Sequence(RawToDigiTask_noTk)
-RawToDigiTask_pixelOnly = cms.Task(siPixelDigis)
+RawToDigiTask_pixelOnly = cms.Task(siPixelDigisTask, scalersRawToDigi)
RawToDigi_pixelOnly = cms.Sequence(RawToDigiTask_pixelOnly)
scalersRawToDigi.scalersInputTag = 'rawDataCollector'
-siPixelDigis.InputLabel = 'rawDataCollector'
+siPixelDigis.cpu.InputLabel = 'rawDataCollector'
#false by default anyways ecalDigis.DoRegional = False
ecalDigis.InputLabel = 'rawDataCollector'
ecalPreshowerDigis.sourceTag = 'rawDataCollector'
diff --git a/Configuration/StandardSequences/python/Reconstruction_cff.py b/Configuration/StandardSequences/python/Reconstruction_cff.py
index 4d4edf170e296..36ac26b2197c6 100644
--- a/Configuration/StandardSequences/python/Reconstruction_cff.py
+++ b/Configuration/StandardSequences/python/Reconstruction_cff.py
@@ -16,7 +16,7 @@
siPixelClusterShapeCachePreSplitting = siPixelClusterShapeCache.clone(
src = 'siPixelClustersPreSplitting'
- )
+)
# Global reco
from RecoEcal.Configuration.RecoEcal_cff import *
diff --git a/DQM/TrackingMonitorClient/python/pixelTrackingEffFromHitPattern_cff.py b/DQM/TrackingMonitorClient/python/pixelTrackingEffFromHitPattern_cff.py
index 15ceaf93ed20a..cff85e56d94f7 100644
--- a/DQM/TrackingMonitorClient/python/pixelTrackingEffFromHitPattern_cff.py
+++ b/DQM/TrackingMonitorClient/python/pixelTrackingEffFromHitPattern_cff.py
@@ -21,7 +21,10 @@ def _layers(suffix, quant, histoPostfix):
]
pixelTrackingEffFromHitPattern = DQMEDHarvester("DQMGenericClient",
- subDirs = cms.untracked.vstring("Tracking/PixelTrackParameters/HitEffFromHitPattern*"),
+ subDirs = cms.untracked.vstring("Tracking/PixelTrackParameters/pixelTracks/HitEffFromHitPattern*",
+ "Tracking/PixelTrackParameters/dzPV0p1/HitEffFromHitPattern*",
+ "Tracking/PixelTrackParameters/pt_0to1/HitEffFromHitPattern*",
+ "Tracking/PixelTrackParameters/pt_1/HitEffFromHitPattern*"),
efficiency = cms.vstring(
_layers("PU", "GoodNumVertices", "") +
_layers("BX", "BX", "VsBX") +
diff --git a/DQM/TrackingMonitorClient/python/pixelVertexResolutionClient_cfi.py b/DQM/TrackingMonitorClient/python/pixelVertexResolutionClient_cfi.py
new file mode 100644
index 0000000000000..2558e88d26012
--- /dev/null
+++ b/DQM/TrackingMonitorClient/python/pixelVertexResolutionClient_cfi.py
@@ -0,0 +1,7 @@
+import FWCore.ParameterSet.Config as cms
+
+from DQM.TrackingMonitorClient.primaryVertexResolutionClient_cfi import primaryVertexResolutionClient as _primaryVertexResolutionClient
+
+pixelVertexResolutionClient = _primaryVertexResolutionClient.clone(
+ subDirs = ["OfflinePixelPV/Resolution/*"]
+)
diff --git a/DQM/TrackingMonitorSource/python/pixelTracksMonitoring_cff.py b/DQM/TrackingMonitorSource/python/pixelTracksMonitoring_cff.py
index a075f671f05ce..d5deba78b46c8 100644
--- a/DQM/TrackingMonitorSource/python/pixelTracksMonitoring_cff.py
+++ b/DQM/TrackingMonitorSource/python/pixelTracksMonitoring_cff.py
@@ -1,23 +1,77 @@
import FWCore.ParameterSet.Config as cms
import DQM.TrackingMonitor.TrackerCollisionTrackingMonitor_cfi
-pixelTracksMonitoring = DQM.TrackingMonitor.TrackerCollisionTrackingMonitor_cfi.TrackerCollisionTrackMon.clone()
-pixelTracksMonitoring.FolderName = 'Tracking/PixelTrackParameters'
-pixelTracksMonitoring.TrackProducer = 'pixelTracks'
-pixelTracksMonitoring.allTrackProducer = 'pixelTracks'
-pixelTracksMonitoring.beamSpot = 'offlineBeamSpot'
-pixelTracksMonitoring.primaryVertex = 'pixelVertices'
-pixelTracksMonitoring.pvNDOF = 1
-pixelTracksMonitoring.doAllPlots = True
-pixelTracksMonitoring.doLumiAnalysis = True
-pixelTracksMonitoring.doProfilesVsLS = True
-pixelTracksMonitoring.doDCAPlots = True
-pixelTracksMonitoring.doProfilesVsLS = True
-pixelTracksMonitoring.doPlotsVsGoodPVtx = True
-pixelTracksMonitoring.doEffFromHitPatternVsPU = False
-pixelTracksMonitoring.doEffFromHitPatternVsBX = False
-pixelTracksMonitoring.doEffFromHitPatternVsLUMI = False
-pixelTracksMonitoring.doPlotsVsGoodPVtx = True
-pixelTracksMonitoring.doPlotsVsLUMI = True
-pixelTracksMonitoring.doPlotsVsBX = True
+pixelTracksMonitor = DQM.TrackingMonitor.TrackerCollisionTrackingMonitor_cfi.TrackerCollisionTrackMon.clone()
+pixelTracksMonitor.FolderName = 'Tracking/PixelTrackParameters/pixelTracks'
+pixelTracksMonitor.TrackProducer = 'pixelTracks'
+pixelTracksMonitor.allTrackProducer = 'pixelTracks'
+pixelTracksMonitor.beamSpot = 'offlineBeamSpot'
+pixelTracksMonitor.primaryVertex = 'pixelVertices'
+pixelTracksMonitor.pvNDOF = 1
+pixelTracksMonitor.doAllPlots = True
+pixelTracksMonitor.doLumiAnalysis = True
+pixelTracksMonitor.doProfilesVsLS = True
+pixelTracksMonitor.doDCAPlots = True
+pixelTracksMonitor.doProfilesVsLS = True
+pixelTracksMonitor.doPlotsVsGoodPVtx = True
+pixelTracksMonitor.doEffFromHitPatternVsPU = False
+pixelTracksMonitor.doEffFromHitPatternVsBX = False
+pixelTracksMonitor.doEffFromHitPatternVsLUMI = False
+pixelTracksMonitor.doPlotsVsGoodPVtx = True
+pixelTracksMonitor.doPlotsVsLUMI = True
+pixelTracksMonitor.doPlotsVsBX = True
+_trackSelector = cms.EDFilter('TrackSelector',
+ src = cms.InputTag('pixelTracks'),
+ cut = cms.string("")
+)
+
+pixelTracksPt0to1 = _trackSelector.clone(cut = "pt >= 0 & pt < 1 ")
+pixelTracksPt1 = _trackSelector.clone(cut = "pt >= 1 ")
+from DQM.TrackingMonitorSource.TrackCollections2monitor_cff import highPurityPV0p1 as _highPurityPV0p1
+pixelTracksPV0p1 = _highPurityPV0p1.clone(
+ src = "pixelTracks",
+ quality = "",
+ vertexTag = "goodPixelVertices"
+)
+
+pixelTracksMonitorPt0to1 = pixelTracksMonitor.clone(
+ TrackProducer = "pixelTracksPt0to1",
+ FolderName = "Tracking/PixelTrackParameters/pt_0to1"
+)
+pixelTracksMonitorPt1 = pixelTracksMonitor.clone(
+ TrackProducer = "pixelTracksPt1",
+ FolderName = "Tracking/PixelTrackParameters/pt_1"
+)
+pixelTracksMonitorPV0p1 = pixelTracksMonitor.clone(
+ TrackProducer = "pixelTracksPV0p1",
+ FolderName = "Tracking/PixelTrackParameters/dzPV0p1"
+)
+
+
+from CommonTools.ParticleFlow.goodOfflinePrimaryVertices_cfi import goodOfflinePrimaryVertices as _goodOfflinePrimaryVertices
+goodPixelVertices = _goodOfflinePrimaryVertices.clone(
+ src = "pixelVertices",
+)
+
+from DQM.TrackingMonitor.primaryVertexResolution_cfi import primaryVertexResolution as _primaryVertexResolution
+pixelVertexResolution = _primaryVertexResolution.clone(
+ vertexSrc = "goodPixelVertices",
+ rootFolder = "OfflinePixelPV/Resolution",
+)
+
+pixelTracksMonitoringTask = cms.Task(
+ goodPixelVertices,
+ pixelTracksPt0to1,
+ pixelTracksPt1,
+ pixelTracksPV0p1,
+)
+
+pixelTracksMonitoring = cms.Sequence(
+ pixelTracksMonitor +
+ pixelTracksMonitorPt0to1 +
+ pixelTracksMonitorPt1 +
+ pixelTracksMonitorPV0p1 +
+ pixelVertexResolution,
+ pixelTracksMonitoringTask
+)
diff --git a/DQMOffline/Configuration/python/DQMOffline_SecondStep_cff.py b/DQMOffline/Configuration/python/DQMOffline_SecondStep_cff.py
index 25b759bc27589..35c6082146f68 100644
--- a/DQMOffline/Configuration/python/DQMOffline_SecondStep_cff.py
+++ b/DQMOffline/Configuration/python/DQMOffline_SecondStep_cff.py
@@ -46,6 +46,7 @@
from DQMOffline.JetMET.dataCertificationJetMET_cff import *
from DQM.TrackingMonitorClient.TrackingClientConfig_Tier0_cff import *
from DQM.TrackingMonitorClient.pixelTrackingEffFromHitPattern_cff import *
+from DQM.TrackingMonitorClient.pixelVertexResolutionClient_cfi import *
from DQM.SiOuterTracker.OuterTrackerClientConfig_cff import *
DQMOffline_SecondStep_PrePOG = cms.Sequence( TrackingOfflineDQMClient *
@@ -117,7 +118,8 @@
DQMHarvestTracking = cms.Sequence( TrackingOfflineDQMClient *
dqmFastTimerServiceClient )
-DQMHarvestPixelTracking = cms.Sequence( pixelTrackingEffFromHitPattern )
+DQMHarvestPixelTracking = cms.Sequence( pixelTrackingEffFromHitPattern *
+ pixelVertexResolutionClient )
DQMHarvestOuterTracker = cms.Sequence( dqmRefHistoRootFileGetter *
dqmDcsInfoClient *
diff --git a/DQMOffline/Configuration/python/DQMOffline_cff.py b/DQMOffline/Configuration/python/DQMOffline_cff.py
index b84ef35e48b56..5cb9af6d3a960 100644
--- a/DQMOffline/Configuration/python/DQMOffline_cff.py
+++ b/DQMOffline/Configuration/python/DQMOffline_cff.py
@@ -49,6 +49,7 @@
from Validation.RecoTau.DQMSequences_cfi import *
from DQM.TrackingMonitorSource.TrackingSourceConfig_Tier0_cff import *
from DQM.TrackingMonitorSource.pixelTracksMonitoring_cff import *
+from DQMOffline.RecoB.PixelVertexMonitor_cff import *
from DQM.SiOuterTracker.OuterTrackerSourceConfig_cff import *
# miniAOD DQM validation
from Validation.RecoParticleFlow.miniAODDQM_cff import *
@@ -114,7 +115,8 @@
materialDumperAnalyzer
)
-DQMOfflinePixelTracking = cms.Sequence( pixelTracksMonitoring )
+DQMOfflinePixelTracking = cms.Sequence( pixelTracksMonitoring +
+ pixelPVMonitor )
DQMOuterTracker = cms.Sequence( dqmDcsInfo *
OuterTrackerSource *
diff --git a/DQMOffline/RecoB/python/PixelVertexMonitor_cff.py b/DQMOffline/RecoB/python/PixelVertexMonitor_cff.py
new file mode 100644
index 0000000000000..3c2e3d7d6700e
--- /dev/null
+++ b/DQMOffline/RecoB/python/PixelVertexMonitor_cff.py
@@ -0,0 +1,7 @@
+import FWCore.ParameterSet.Config as cms
+
+from DQMOffline.RecoB.PrimaryVertexMonitor_cff import pvMonitor as _pvMonitor
+pixelPVMonitor = _pvMonitor.clone(
+ TopFolderName = "OfflinePixelPV",
+ vertexLabel = "pixelVertices",
+)
diff --git a/DataFormats/CaloRecHit/test/BuildFile.xml b/DataFormats/CaloRecHit/test/BuildFile.xml
index 983e853f47698..6daf8cf086086 100644
--- a/DataFormats/CaloRecHit/test/BuildFile.xml
+++ b/DataFormats/CaloRecHit/test/BuildFile.xml
@@ -1,13 +1,14 @@
-
-
-
-
-
+
+
+
+
+
+
-
+
diff --git a/DataFormats/CaloRecHit/test/test_calo_rechit.cu b/DataFormats/CaloRecHit/test/test_calo_rechit.cu
index 595127824d61e..76475bc0aabdb 100644
--- a/DataFormats/CaloRecHit/test/test_calo_rechit.cu
+++ b/DataFormats/CaloRecHit/test/test_calo_rechit.cu
@@ -1,10 +1,11 @@
+#include
+#include
+
#include
#include
-#include
-#include
-
#include "DataFormats/CaloRecHit/interface/CaloRecHit.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
__global__ void kernel_test_calo_rechit(CaloRecHit* other) {
CaloRecHit rh{DetId(0), 10, 1, 0, 0};
@@ -42,12 +43,9 @@ void test_calo_rechit() {
}
int main(int argc, char** argv) {
- int nDevices;
- cudaGetDeviceCount(&nDevices);
- std::cout << "nDevices = " << nDevices << std::endl;
+ exitSansCUDADevices();
- if (nDevices > 0)
- test_calo_rechit();
+ test_calo_rechit();
std::cout << "all good!" << std::endl;
return 0;
diff --git a/DataFormats/DetId/test/BuildFile.xml b/DataFormats/DetId/test/BuildFile.xml
index 0cccd9fb0d26b..376a8bdc397ad 100644
--- a/DataFormats/DetId/test/BuildFile.xml
+++ b/DataFormats/DetId/test/BuildFile.xml
@@ -1,6 +1,7 @@
-
+
+
diff --git a/DataFormats/DetId/test/test_detid.cu b/DataFormats/DetId/test/test_detid.cu
index 8e67d0e0dc1bc..60d2acd3e3f8f 100644
--- a/DataFormats/DetId/test/test_detid.cu
+++ b/DataFormats/DetId/test/test_detid.cu
@@ -1,10 +1,12 @@
-#include
+#include
+#include
+
#include
+#include
-#include
-#include
#include "DataFormats/DetId/interface/DetId.h"
#include "DataFormats/HcalDetId/interface/HcalDetId.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
__global__ void test_gen_detid(DetId* id, uint32_t const rawid) {
DetId did{rawid};
@@ -27,11 +29,7 @@ void test_detid() {
}
int main(int argc, char** argv) {
- int nDevices;
- cudaGetDeviceCount(&nDevices);
- std::cout << "nDevices = " << nDevices << std::endl;
+ exitSansCUDADevices();
- // test det id functionality
- if (nDevices > 0)
- test_detid();
+ test_detid();
}
diff --git a/DataFormats/GeometrySurface/interface/SOARotation.h b/DataFormats/GeometrySurface/interface/SOARotation.h
index 1373a4091c5e5..14f3df34b993c 100644
--- a/DataFormats/GeometrySurface/interface/SOARotation.h
+++ b/DataFormats/GeometrySurface/interface/SOARotation.h
@@ -139,6 +139,48 @@ class SOAFrame {
ux+=px; uy+=py; uz+=pz;
}
+ constexpr inline
+ void toGlobal(
+ T cxx,
+ T cxy,
+ T cyy,
+ T * gl) const {
+
+ auto const & r = rot;
+ gl[0] = r.xx()*(r.xx()*cxx+r.yx()*cxy) + r.yx()*(r.xx()*cxy+r.yx()*cyy);
+ gl[1] = r.xx()*(r.xy()*cxx+r.yy()*cxy) + r.yx()*(r.xy()*cxy+r.yy()*cyy);
+ gl[2] = r.xy()*(r.xy()*cxx+r.yy()*cxy) + r.yy()*(r.xy()*cxy+r.yy()*cyy);
+ gl[3] = r.xx()*(r.xz()*cxx+r.yz()*cxy) + r.yx()*(r.xz()*cxy+r.yz()*cyy);
+ gl[4] = r.xy()*(r.xz()*cxx+r.yz()*cxy) + r.yy()*(r.xz()*cxy+r.yz()*cyy);
+ gl[5] = r.xz()*(r.xz()*cxx+r.yz()*cxy) + r.yz()*(r.xz()*cxy+r.yz()*cyy);
+ }
+
+ constexpr inline
+ void toLocal(
+ T const * ge,
+ T & lxx,
+ T & lxy,
+ T & lyy ) const {
+
+ auto const & r = rot;
+
+ T cxx = ge[0]; T cyx = ge[1]; T cyy = ge[2];
+ T czx = ge[3]; T czy = ge[4]; T czz = ge[5];
+
+ lxx
+ = r.xx()*(r.xx()*cxx + r.xy()*cyx + r.xz()*czx)
+ + r.xy()*(r.xx()*cyx + r.xy()*cyy + r.xz()*czy)
+ + r.xz()*(r.xx()*czx + r.xy()*czy + r.xz()*czz);
+ lxy
+ = r.yx()*(r.xx()*cxx + r.xy()*cyx + r.xz()*czx)
+ + r.yy()*(r.xx()*cyx + r.xy()*cyy + r.xz()*czy)
+ + r.yz()*(r.xx()*czx + r.xy()*czy + r.xz()*czz);
+ lyy
+ = r.yx()*(r.yx()*cxx + r.yy()*cyx + r.yz()*czx)
+ + r.yy()*(r.yx()*cyx + r.yy()*cyy + r.yz()*czy)
+ + r.yz()*(r.yx()*czx + r.yy()*czy + r.yz()*czz);
+ }
+
constexpr inline
T x() const { return px; }
diff --git a/DataFormats/GeometrySurface/test/BuildFile.xml b/DataFormats/GeometrySurface/test/BuildFile.xml
index ed732f292ebf1..eda35a6f42eef 100644
--- a/DataFormats/GeometrySurface/test/BuildFile.xml
+++ b/DataFormats/GeometrySurface/test/BuildFile.xml
@@ -14,3 +14,18 @@
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu b/DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu
new file mode 100644
index 0000000000000..f4c72d1ad8a02
--- /dev/null
+++ b/DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu
@@ -0,0 +1,38 @@
+#include
+#include
+#include
+
+#include "cuda/api_wrappers.h"
+
+#include "DataFormats/GeometrySurface/interface/SOARotation.h"
+
+__global__
+void toGlobal(SOAFrame const * frame,
+ float const * xl, float const * yl,
+ float * x, float * y, float * z,
+ float const * le, float * ge,
+ uint32_t n)
+{
+ int i = blockDim.x * blockIdx.x + threadIdx.x;
+ if (i >= n) return;
+
+ frame[0].toGlobal(xl[i],yl[i],x[i],y[i],z[i]);
+ frame[0].toGlobal(le[3*i],le[3*i+1],le[3*i+2],ge+6*i);
+}
+
+void toGlobalWrapper(SOAFrame const * frame,
+ float const * xl, float const * yl,
+ float * x, float * y, float * z,
+ float const * le, float * ge,
+ uint32_t n)
+{
+ int threadsPerBlock = 256;
+ int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
+ std::cout
+ << "CUDA toGlobal kernel launch with " << blocksPerGrid
+ << " blocks of " << threadsPerBlock << " threads" << std::endl;
+
+ cuda::launch(toGlobal,
+ { blocksPerGrid, threadsPerBlock },
+ frame, xl, yl, x, y, z, le, ge, n);
+}
diff --git a/DataFormats/GeometrySurface/test/gpuFrameTransformTest.cpp b/DataFormats/GeometrySurface/test/gpuFrameTransformTest.cpp
new file mode 100644
index 0000000000000..24b4f8f2441cc
--- /dev/null
+++ b/DataFormats/GeometrySurface/test/gpuFrameTransformTest.cpp
@@ -0,0 +1,113 @@
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+
+#include
+
+#include "DataFormats/GeometrySurface/interface/GloballyPositioned.h"
+#include "DataFormats/GeometrySurface/interface/SOARotation.h"
+#include "DataFormats/GeometrySurface/interface/TkRotation.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
+
+void toGlobalWrapper(SOAFrame const * frame,
+ float const * xl, float const * yl,
+ float * x, float * y, float * z,
+ float const * le, float * ge,
+ uint32_t n);
+
+int main(void)
+{
+ exitSansCUDADevices();
+
+ typedef float T;
+ typedef TkRotation Rotation;
+ typedef SOARotation SRotation;
+ typedef GloballyPositioned Frame;
+ typedef SOAFrame SFrame;
+ typedef typename Frame::PositionType Position;
+ typedef typename Frame::GlobalVector GlobalVector;
+ typedef typename Frame::GlobalPoint GlobalPoint;
+ typedef typename Frame::LocalVector LocalVector;
+ typedef typename Frame::LocalPoint LocalPoint;
+
+ constexpr uint32_t size = 10000;
+ constexpr uint32_t size32 = size*sizeof(float);
+
+ float xl[size],yl[size];
+ float x[size],y[size],z[size];
+
+ // errors
+ float le[3*size];
+ float ge[6*size];
+
+ auto current_device = cuda::device::current::get();
+ auto d_xl = cuda::memory::device::make_unique(current_device, size);
+ auto d_yl = cuda::memory::device::make_unique(current_device, size);
+
+ auto d_x = cuda::memory::device::make_unique(current_device, size);
+ auto d_y = cuda::memory::device::make_unique(current_device, size);
+ auto d_z = cuda::memory::device::make_unique(current_device, size);
+
+ auto d_le = cuda::memory::device::make_unique(current_device, 3*size);
+ auto d_ge = cuda::memory::device::make_unique(current_device, 6*size);
+
+ double a = 0.01;
+ double ca = std::cos(a);
+ double sa = std::sin(a);
+
+ Rotation r1(ca, sa, 0,
+ -sa, ca, 0,
+ 0, 0, 1);
+ Frame f1(Position(2,3,4), r1);
+ std::cout << "f1.position() " << f1.position() << std::endl;
+ std::cout << "f1.rotation() " << '\n' << f1.rotation() << std::endl;
+
+ SFrame sf1(f1.position().x(),
+ f1.position().y(),
+ f1.position().z(),
+ f1.rotation()
+ );
+
+ // auto d_sf = cuda::memory::device::make_unique(current_device, 1);
+ auto d_sf = cuda::memory::device::make_unique(current_device, sizeof(SFrame));
+ cuda::memory::copy(d_sf.get(), &sf1, sizeof(SFrame));
+
+ for (auto i=0U; isize/2) ? 1.f : 0.04f;
+ le[2*i+1]=0.;
+ }
+ std::random_shuffle(xl,xl+size);
+ std::random_shuffle(yl,yl+size);
+
+ cuda::memory::copy(d_xl.get(), xl, size32);
+ cuda::memory::copy(d_yl.get(), yl, size32);
+ cuda::memory::copy(d_le.get(), le, 3*size32);
+
+ toGlobalWrapper((SFrame const *)(d_sf.get()), d_xl.get(), d_yl.get(), d_x.get(), d_y.get(), d_z.get(),
+ d_le.get(), d_ge.get(), size
+ );
+
+ cuda::memory::copy(x,d_x.get(), size32);
+ cuda::memory::copy(y,d_y.get(), size32);
+ cuda::memory::copy(z,d_z.get(), size32);
+ cuda::memory::copy(ge,d_ge.get(), 6*size32);
+
+ float eps=0.;
+ for (auto i=0U; i
+
diff --git a/DataFormats/HcalDetId/test/test_hcal_detid.cu b/DataFormats/HcalDetId/test/test_hcal_detid.cu
index ea2099ef2971f..86ba3746f7819 100644
--- a/DataFormats/HcalDetId/test/test_hcal_detid.cu
+++ b/DataFormats/HcalDetId/test/test_hcal_detid.cu
@@ -1,10 +1,12 @@
-#include
+#include
+#include
+
#include
+#include
-#include
-#include
#include "DataFormats/DetId/interface/DetId.h"
#include "DataFormats/HcalDetId/interface/HcalDetId.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
__global__ void test_gen_detid(DetId* id) {
DetId did;
@@ -63,17 +65,13 @@ void test_hcal_detid() {
}
int main(int argc, char** argv) {
- int nDevices;
- cudaGetDeviceCount(&nDevices);
- std::cout << "nDevices = " << nDevices << std::endl;
+ exitSansCUDADevices();
// test det id functionality
- if (nDevices>0)
- test_detid();
+ test_detid();
// test hcal det ids
- if (nDevices>0)
- test_hcal_detid();
+ test_hcal_detid();
return 0;
}
diff --git a/DataFormats/HcalDigi/test/BuildFile.xml b/DataFormats/HcalDigi/test/BuildFile.xml
index d8cdba8988013..6c9921512d2c2 100644
--- a/DataFormats/HcalDigi/test/BuildFile.xml
+++ b/DataFormats/HcalDigi/test/BuildFile.xml
@@ -1,15 +1,16 @@
-
-
-
-
-
-
+
+
+
+
+
+
-
-
-
-
+
+
+
+
+
diff --git a/DataFormats/HcalDigi/test/test_hcal_digi.cu b/DataFormats/HcalDigi/test/test_hcal_digi.cu
index fafa52cbc6ee4..0f03548b0c524 100644
--- a/DataFormats/HcalDigi/test/test_hcal_digi.cu
+++ b/DataFormats/HcalDigi/test/test_hcal_digi.cu
@@ -1,15 +1,17 @@
-#include
+#include
+#include
+
#include
+#include
-#include
-#include
+#include "DataFormats/Common/interface/DataFrame.h"
#include "DataFormats/DetId/interface/DetId.h"
#include "DataFormats/HcalDetId/interface/HcalDetId.h"
#include "DataFormats/HcalDigi/interface/HBHEDataFrame.h"
+#include "DataFormats/HcalDigi/interface/HcalDigiCollections.h"
#include "DataFormats/HcalDigi/interface/QIE10DataFrame.h"
#include "DataFormats/HcalDigi/interface/QIE11DataFrame.h"
-#include "DataFormats/HcalDigi/interface/HcalDigiCollections.h"
-#include "DataFormats/Common/interface/DataFrame.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
__global__ void kernel_test_hcal_qiesample(HcalQIESample* sample, uint16_t value) {
printf("kernel: testing hcal qie sampel\n");
@@ -69,7 +71,6 @@ void test_hcal_qie1011_digis() {
constexpr int samples = 10;
constexpr int detid = 2;
HcalDataFrameContainer coll{samples, detid};
- TDF *d_dfs;
uint16_t *d_data;
uint32_t *d_out;
uint32_t h_out[size], h_test_out[size];
@@ -164,22 +165,18 @@ void test_hcal_qie8_hbhedf() {
}
int main(int argc, char** argv) {
- int nDevices;
- cudaGetDeviceCount(&nDevices);
- std::cout << "nDevices = " << nDevices << std::endl;
-
- if (nDevices > 0) {
- // qie8
- test_hcal_qiesample();
- test_hcal_qie8_hbhedf();
- test_hcal_qie8_digis();
- test_hcal_qie8_digis();
- test_hcal_qie8_digis();
-
- // qie1011
- test_hcal_qie1011_digis();
- test_hcal_qie1011_digis();
- }
+ exitSansCUDADevices();
+
+ // qie8
+ test_hcal_qiesample();
+ test_hcal_qie8_hbhedf();
+ test_hcal_qie8_digis();
+ test_hcal_qie8_digis();
+ test_hcal_qie8_digis();
+
+ // qie1011
+ test_hcal_qie1011_digis();
+ test_hcal_qie1011_digis();
return 0;
}
diff --git a/DataFormats/HcalRecHit/test/BuildFile.xml b/DataFormats/HcalRecHit/test/BuildFile.xml
index 2772404fe9b89..e0d2280530a2b 100644
--- a/DataFormats/HcalRecHit/test/BuildFile.xml
+++ b/DataFormats/HcalRecHit/test/BuildFile.xml
@@ -1,16 +1,17 @@
-
-
-
-
-
-
+
+
+
+
+
+
-
-
-
-
-
+
+
+
+
+
+
diff --git a/DataFormats/HcalRecHit/test/test_hcal_reco.cu b/DataFormats/HcalRecHit/test/test_hcal_reco.cu
index 4119d3592898c..4e95adb856d81 100644
--- a/DataFormats/HcalRecHit/test/test_hcal_reco.cu
+++ b/DataFormats/HcalRecHit/test/test_hcal_reco.cu
@@ -10,6 +10,7 @@
#include "DataFormats/HcalRecHit/interface/HORecHit.h"
#include "DataFormats/HcalRecHit/interface/HFQIE10Info.h"
#include "DataFormats/HcalRecHit/interface/HBHEChannelInfo.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
template
__global__ void kernel_test_hcal_rechits(T *other) {
@@ -114,18 +115,13 @@ void test_hcal_hbhechinfo() {
}
int main(int argc, char ** argv) {
- int nDevices;
- cudaGetDeviceCount(&nDevices);
- std::cout << "nDevices = " << nDevices << std::endl;
+ exitSansCUDADevices();
- if (nDevices > 0) {
- test_hcal_rechits();
- test_hcal_rechits();
- test_hcal_rechits();
- test_hcal_hbhechinfo();
-
- std::cout << "all good" << std::endl;
- }
+ test_hcal_rechits();
+ test_hcal_rechits();
+ test_hcal_rechits();
+ test_hcal_hbhechinfo();
+ std::cout << "all good" << std::endl;
return 0;
}
diff --git a/DataFormats/Math/test/BuildFile.xml b/DataFormats/Math/test/BuildFile.xml
index 160e9df5535f4..7d3f50abe7c91 100644
--- a/DataFormats/Math/test/BuildFile.xml
+++ b/DataFormats/Math/test/BuildFile.xml
@@ -1,27 +1,27 @@
-
-
-
-
+
+
+
+
-
-
-
-
+
+
+
+
-
+
-
+
-
+
-
+
@@ -29,54 +29,56 @@
-
+
-
+
-
+
-
+
-
+
-
+
-
+
-
+
-
+
-
+
-
+
-
+
+
-
+
+
@@ -84,7 +86,8 @@
-
+
+
@@ -92,7 +95,8 @@
-
+
+
diff --git a/DataFormats/Math/test/cudaAtan2Test.cu b/DataFormats/Math/test/cudaAtan2Test.cu
index 077c0d7238259..38b453b9c520f 100644
--- a/DataFormats/Math/test/cudaAtan2Test.cu
+++ b/DataFormats/Math/test/cudaAtan2Test.cu
@@ -29,6 +29,7 @@ end
#include "cuda/api_wrappers.h"
#include "DataFormats/Math/interface/approx_atan2.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
constexpr float xmin=-100.001; // avoid 0
constexpr float incr = 0.04;
@@ -103,16 +104,7 @@ void go() {
}
int main() {
- int count = 0;
- auto status = cudaGetDeviceCount(& count);
- if (status != cudaSuccess) {
- std::cerr << "Failed to initialise the CUDA runtime, the test will be skipped." << "\n";
- exit(EXIT_SUCCESS);
- }
- if (count == 0) {
- std::cerr << "No CUDA devices on this system, the test will be skipped." << "\n";
- exit(EXIT_SUCCESS);
- }
+ exitSansCUDADevices();
try {
go<3>();
diff --git a/DataFormats/Math/test/cudaMathTest.cu b/DataFormats/Math/test/cudaMathTest.cu
index 8088e9c6e6fc3..dc97cda39c6cf 100644
--- a/DataFormats/Math/test/cudaMathTest.cu
+++ b/DataFormats/Math/test/cudaMathTest.cu
@@ -28,10 +28,6 @@ end
#include "cuda/api_wrappers.h"
-#include
-#include
-#include
-
#ifdef __CUDACC__
#define inline __host__ __device__ inline
#include
@@ -40,6 +36,11 @@ end
#include
#endif
+#include "DataFormats/Math/interface/approx_log.h"
+#include "DataFormats/Math/interface/approx_exp.h"
+#include "DataFormats/Math/interface/approx_atan2.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"
+
std::mt19937 eng;
std::mt19937 eng2;
std::uniform_real_distribution rgen(0.,1.);
@@ -205,22 +206,12 @@ void go()
}
int main() {
- int count = 0;
- auto status = cudaGetDeviceCount(& count);
- if (status != cudaSuccess) {
- std::cerr << "Failed to initialise the CUDA runtime, the test will be skipped." << "\n";
- exit(EXIT_SUCCESS);
- }
- if (count == 0) {
- std::cerr << "No CUDA devices on this system, the test will be skipped." << "\n";
- exit(EXIT_SUCCESS);
- }
+ exitSansCUDADevices();
try {
go();
go();
go();
-
go();
} catch(cuda::runtime_error &ex) {
std::cerr << "CUDA error: " << ex.what() << std::endl;
diff --git a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h
index 22f9cb1020814..ba75447e945bb 100644
--- a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h
+++ b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h
@@ -21,6 +21,7 @@
#include
#include
#include
+#include
class PixelDigi;
@@ -196,7 +197,10 @@ class SiPixelCluster {
float getSplitClusterErrorX() const { return err_x; }
float getSplitClusterErrorY() const { return err_y; }
-
+ // the original id (they get sorted)
+ auto originalId() const { return theOriginalClusterId;}
+ void setOriginalId(uint16_t id) { theOriginalClusterId=id;}
+
private:
std::vector thePixelOffset;
@@ -207,6 +211,8 @@ class SiPixelCluster {
uint16_t theMinPixelCol=MAXPOS; // Minimum pixel index in the y direction (left edge).
uint8_t thePixelRowSpan=0; // Span pixel index in the x direction (low edge).
uint8_t thePixelColSpan=0; // Span pixel index in the y direction (left edge).
+
+ uint16_t theOriginalClusterId=std::numeric_limits::max();
float err_x=-99999.9f;
float err_y=-99999.9f;
diff --git a/DataFormats/SiPixelCluster/src/classes_def.xml b/DataFormats/SiPixelCluster/src/classes_def.xml
index 55c9fd8538417..d43f062877eb0 100644
--- a/DataFormats/SiPixelCluster/src/classes_def.xml
+++ b/DataFormats/SiPixelCluster/src/classes_def.xml
@@ -4,6 +4,7 @@
+
diff --git a/DataFormats/SiPixelDigi/interface/PixelErrors.h b/DataFormats/SiPixelDigi/interface/PixelErrors.h
new file mode 100644
index 0000000000000..5231b7d1f372a
--- /dev/null
+++ b/DataFormats/SiPixelDigi/interface/PixelErrors.h
@@ -0,0 +1,21 @@
+#ifndef DataFormats_SiPixelDigi_PixelErrors_h
+#define DataFormats_SiPixelDigi_PixelErrors_h
+
+#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h"
+#include "FWCore/Utilities/interface/typedefs.h"
+
+#include