diff --git a/CUDADataFormats/Track/BuildFile.xml b/CUDADataFormats/Track/BuildFile.xml
deleted file mode 100644
index cf07e3b540f24..0000000000000
--- a/CUDADataFormats/Track/BuildFile.xml
+++ /dev/null
@@ -1,10 +0,0 @@
-
-
-
-
-
-
-
-
-
-
diff --git a/CUDADataFormats/Track/README.md b/CUDADataFormats/Track/README.md
deleted file mode 100644
index 8f66d9e4c4467..0000000000000
--- a/CUDADataFormats/Track/README.md
+++ /dev/null
@@ -1,50 +0,0 @@
-# Track CUDA Data Formats
-
-`CUDADataFormat`s meant to be used on Host (CPU) or Device (CUDA GPU) for
-storing information about `Track`s created during the Pixel-local Reconstruction
-chain. It stores data in an SoA manner. It combines the data contained in the
-deprecated `TrackSoAHeterogeneousT` and `TrajectoryStateSoAT` classes.
-
-The host format is inheriting from `CUDADataFormats/Common/interface/PortableHostCollection.h`,
-while the device format is inheriting from `CUDADataFormats/Common/interface/PortableDeviceCollection.h`
-
-Both formats use the same SoA Layout (`TrackSoAHeterogeneousLayout`) which is generated
-via the `GENERATE_SOA_LAYOUT` macro in the `PixelTrackUtilities.h` file.
-
-## Notes
-
--`hitIndices` and `detIndices`, instances of `HitContainer`, have been added into the
-layout as `SOA_SCALAR`s, meaning that they manage their own data independently from the SoA
-`Layout`. This could be improved in the future, if `HitContainer` (aka a `OneToManyAssoc` of fixed size)
-is replaced, but there don't seem to be any conflicts in including it in the `Layout` like this.
-- Host and Device classes should **not** be created via inheritance, as they're done here,
-but via composition. See [this discussion](https://github.com/cms-sw/cmssw/pull/40465#discussion_r1066039309).
-
-## TrackSoAHeterogeneousHost
-
-The version of the data format to be used for storing `Track` data on the CPU.
-Instances of this class are to be used for:
-
-- Having a place to copy data to host from device, via `cudaMemcpy`, or
-- Running host-side algorithms using data stored in an SoA manner.
-
-## TrackSoAHeterogeneousDevice
-
-The version of the data format to be used for storing `Track` data on the GPU.
-
-Instances of `TrackSoAHeterogeneousDevice` are to be created on host and be
-used on device only. To do so, the instance's `view()` method is to be called
-to pass a `View` to any kernel launched. Accessing data from the `view()` is not
-possible on the host side.
-
-## Utilities
-
-`PixelTrackUtilities.h` contains a collection of methods which were originally
-defined as class methods inside either `TrackSoAHeterogeneousT` and `TrajectoryStateSoAT`
-which have been adapted to operate on `View` instances, so that they are callable
-from within `__global__` kernels, on both CPU and CPU.
-
-## Use case
-
-See `test/TrackSoAHeterogeneous_test.cpp` for a simple example of instantiation,
-processing and copying from device to host.
diff --git a/CUDADataFormats/Track/interface/PixelTrackUtilities.h b/CUDADataFormats/Track/interface/PixelTrackUtilities.h
deleted file mode 100644
index 6d7ea258be8d2..0000000000000
--- a/CUDADataFormats/Track/interface/PixelTrackUtilities.h
+++ /dev/null
@@ -1,243 +0,0 @@
-#ifndef CUDADataFormats_Track_PixelTrackUtilities_h
-#define CUDADataFormats_Track_PixelTrackUtilities_h
-
-#include
-#include
-#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h"
-#include "DataFormats/SoATemplate/interface/SoALayout.h"
-
-namespace pixelTrack {
-
- enum class Quality : uint8_t { bad = 0, edup, dup, loose, strict, tight, highPurity, notQuality };
- constexpr uint32_t qualitySize{uint8_t(Quality::notQuality)};
- const std::string qualityName[qualitySize]{"bad", "edup", "dup", "loose", "strict", "tight", "highPurity"};
- inline Quality qualityByName(std::string const &name) {
- auto qp = std::find(qualityName, qualityName + qualitySize, name) - qualityName;
- return static_cast(qp);
- }
-
-} // namespace pixelTrack
-
-template
-struct TrackSoA {
- static constexpr int32_t S = TrackerTraits::maxNumberOfTuples;
- static constexpr int32_t H = TrackerTraits::avgHitsPerTrack;
- // Aliases in order to not confuse the GENERATE_SOA_LAYOUT
- // macro with weird colons and angled brackets.
- using Vector5f = Eigen::Matrix;
- using Vector15f = Eigen::Matrix;
- using Quality = pixelTrack::Quality;
-
- using hindex_type = uint32_t;
-
- using HitContainer = cms::cuda::OneToManyAssoc;
-
- GENERATE_SOA_LAYOUT(TrackSoALayout,
- SOA_COLUMN(Quality, quality),
- SOA_COLUMN(float, chi2),
- SOA_COLUMN(int8_t, nLayers),
- SOA_COLUMN(float, eta),
- SOA_COLUMN(float, pt),
- SOA_EIGEN_COLUMN(Vector5f, state),
- SOA_EIGEN_COLUMN(Vector15f, covariance),
- SOA_SCALAR(int, nTracks),
- SOA_SCALAR(HitContainer, hitIndices),
- SOA_SCALAR(HitContainer, detIndices))
-};
-
-// Methods that operate on View and ConstView of the TrackSoA, and cannot be class methods.
-
-template
-struct TracksUtilities {
- using TrackSoAView = typename TrackSoA::template TrackSoALayout<>::View;
- using TrackSoAConstView = typename TrackSoA::template TrackSoALayout<>::ConstView;
- using hindex_type = typename TrackSoA::hindex_type;
-
- // State at the Beam spot
- // phi,tip,1/pt,cotan(theta),zip
- static __host__ __device__ inline float charge(const TrackSoAConstView &tracks, int32_t i) {
- return std::copysign(1.f, tracks[i].state()(2));
- }
-
- static constexpr __host__ __device__ inline float phi(const TrackSoAConstView &tracks, int32_t i) {
- return tracks[i].state()(0);
- }
-
- static constexpr __host__ __device__ inline float tip(const TrackSoAConstView &tracks, int32_t i) {
- return tracks[i].state()(1);
- }
-
- static constexpr __host__ __device__ inline float zip(const TrackSoAConstView &tracks, int32_t i) {
- return tracks[i].state()(4);
- }
-
- static constexpr __host__ __device__ inline bool isTriplet(const TrackSoAConstView &tracks, int i) {
- return tracks[i].nLayers() == 3;
- }
-
- template
- static constexpr __host__ __device__ inline void copyFromCircle(
- TrackSoAView &tracks, V3 const &cp, M3 const &ccov, V2 const &lp, M2 const &lcov, float b, int32_t i) {
- tracks[i].state() << cp.template cast(), lp.template cast();
-
- tracks[i].state()(2) = tracks[i].state()(2) * b;
- auto cov = tracks[i].covariance();
- cov(0) = ccov(0, 0);
- cov(1) = ccov(0, 1);
- cov(2) = b * float(ccov(0, 2));
- cov(4) = cov(3) = 0;
- cov(5) = ccov(1, 1);
- cov(6) = b * float(ccov(1, 2));
- cov(8) = cov(7) = 0;
- cov(9) = b * b * float(ccov(2, 2));
- cov(11) = cov(10) = 0;
- cov(12) = lcov(0, 0);
- cov(13) = lcov(0, 1);
- cov(14) = lcov(1, 1);
- }
-
- template
- static constexpr __host__ __device__ inline void copyFromDense(TrackSoAView &tracks,
- V5 const &v,
- M5 const &cov,
- int32_t i) {
- tracks[i].state() = v.template cast();
- for (int j = 0, ind = 0; j < 5; ++j)
- for (auto k = j; k < 5; ++k)
- tracks[i].covariance()(ind++) = cov(j, k);
- }
-
- template
- static constexpr __host__ __device__ inline void copyToDense(const TrackSoAConstView &tracks,
- V5 &v,
- M5 &cov,
- int32_t i) {
- v = tracks[i].state().template cast();
- for (int j = 0, ind = 0; j < 5; ++j) {
- cov(j, j) = tracks[i].covariance()(ind++);
- for (auto k = j + 1; k < 5; ++k)
- cov(k, j) = cov(j, k) = tracks[i].covariance()(ind++);
- }
- }
-
- static constexpr __host__ __device__ inline int computeNumberOfLayers(const TrackSoAConstView &tracks, int32_t i) {
- auto pdet = tracks.detIndices().begin(i);
- int nl = 1;
- auto ol = pixelTopology::getLayer(*pdet);
- for (; pdet < tracks.detIndices().end(i); ++pdet) {
- auto il = pixelTopology::getLayer(*pdet);
- if (il != ol)
- ++nl;
- ol = il;
- }
- return nl;
- }
-
- static constexpr __host__ __device__ inline int nHits(const TrackSoAConstView &tracks, int i) {
- return tracks.detIndices().size(i);
- }
-};
-
-namespace pixelTrack {
-
- template
- struct QualityCutsT {};
-
- template
- struct QualityCutsT> {
- using TrackSoAView = typename TrackSoA::template TrackSoALayout<>::View;
- using TrackSoAConstView = typename TrackSoA::template TrackSoALayout<>::ConstView;
- using tracksHelper = TracksUtilities;
- // chi2 cut = chi2Scale * (chi2Coeff[0] + pT/GeV * (chi2Coeff[1] + pT/GeV * (chi2Coeff[2] + pT/GeV * chi2Coeff[3])))
- float chi2Coeff[4];
- float chi2MaxPt; // GeV
- float chi2Scale;
-
- struct Region {
- float maxTip; // cm
- float minPt; // GeV
- float maxZip; // cm
- };
-
- Region triplet;
- Region quadruplet;
-
- __device__ __forceinline__ bool isHP(const TrackSoAConstView &tracks, int nHits, int it) const {
- // impose "region cuts" based on the fit results (phi, Tip, pt, cotan(theta)), Zip)
- // default cuts:
- // - for triplets: |Tip| < 0.3 cm, pT > 0.5 GeV, |Zip| < 12.0 cm
- // - for quadruplets: |Tip| < 0.5 cm, pT > 0.3 GeV, |Zip| < 12.0 cm
- // (see CAHitNtupletGeneratorGPU.cc)
- auto const ®ion = (nHits > 3) ? quadruplet : triplet;
- return (std::abs(tracksHelper::tip(tracks, it)) < region.maxTip) and (tracks.pt(it) > region.minPt) and
- (std::abs(tracksHelper::zip(tracks, it)) < region.maxZip);
- }
-
- __device__ __forceinline__ bool strictCut(const TrackSoAConstView &tracks, int it) const {
- auto roughLog = [](float x) {
- // max diff [0.5,12] at 1.25 0.16143
- // average diff 0.0662998
- union IF {
- uint32_t i;
- float f;
- };
- IF z;
- z.f = x;
- uint32_t lsb = 1 < 21;
- z.i += lsb;
- z.i >>= 21;
- auto f = z.i & 3;
- int ex = int(z.i >> 2) - 127;
-
- // log2(1+0.25*f)
- // averaged over bins
- const float frac[4] = {0.160497f, 0.452172f, 0.694562f, 0.901964f};
- return float(ex) + frac[f];
- };
-
- float pt = std::min(tracks.pt(it), chi2MaxPt);
- float chi2Cut = chi2Scale * (chi2Coeff[0] + roughLog(pt) * chi2Coeff[1]);
- if (tracks.chi2(it) >= chi2Cut) {
-#ifdef NTUPLE_FIT_DEBUG
- printf("Bad chi2 %d pt %f eta %f chi2 %f\n", it, tracks.pt(it), tracks.eta(it), tracks.chi2(it));
-#endif
- return true;
- }
- return false;
- }
- };
-
- template
- struct QualityCutsT> {
- using TrackSoAView = typename TrackSoA::template TrackSoALayout<>::View;
- using TrackSoAConstView = typename TrackSoA::template TrackSoALayout<>::ConstView;
- using tracksHelper = TracksUtilities;
-
- float maxChi2;
- float minPt;
- float maxTip;
- float maxZip;
-
- __device__ __forceinline__ bool isHP(const TrackSoAConstView &tracks, int nHits, int it) const {
- return (std::abs(tracksHelper::tip(tracks, it)) < maxTip) and (tracks.pt(it) > minPt) and
- (std::abs(tracksHelper::zip(tracks, it)) < maxZip);
- }
- __device__ __forceinline__ bool strictCut(const TrackSoAConstView &tracks, int it) const {
- return tracks.chi2(it) >= maxChi2;
- }
- };
-
-} // namespace pixelTrack
-
-template
-using TrackLayout = typename TrackSoA::template TrackSoALayout<>;
-template
-using TrackSoAView = typename TrackSoA::template TrackSoALayout<>::View;
-template
-using TrackSoAConstView = typename TrackSoA::template TrackSoALayout<>::ConstView;
-
-template struct TracksUtilities;
-template struct TracksUtilities;
-
-#endif
diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h
deleted file mode 100644
index 04d286a767ab0..0000000000000
--- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h
+++ /dev/null
@@ -1,37 +0,0 @@
-#ifndef CUDADataFormats_Track_TrackHeterogeneousDevice_H
-#define CUDADataFormats_Track_TrackHeterogeneousDevice_H
-
-#include
-
-#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
-#include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"
-
-#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
-
-// TODO: The class is created via inheritance of the PortableDeviceCollection.
-// This is generally discouraged, and should be done via composition.
-// See: https://github.com/cms-sw/cmssw/pull/40465#discussion_r1067364306
-template
-class TrackSoAHeterogeneousDevice : public cms::cuda::PortableDeviceCollection> {
-public:
- using cms::cuda::PortableDeviceCollection>::view;
- using cms::cuda::PortableDeviceCollection>::const_view;
- using cms::cuda::PortableDeviceCollection>::buffer;
- using cms::cuda::PortableDeviceCollection>::bufferSize;
-
- TrackSoAHeterogeneousDevice() = default; // cms::cuda::Product needs this
-
- // Constructor which specifies the SoA size
- explicit TrackSoAHeterogeneousDevice(cudaStream_t stream)
- : cms::cuda::PortableDeviceCollection>(TrackerTraits::maxNumberOfTuples, stream) {}
-};
-
-namespace pixelTrack {
-
- using TrackSoADevicePhase1 = TrackSoAHeterogeneousDevice;
- using TrackSoADevicePhase2 = TrackSoAHeterogeneousDevice;
- using TrackSoADeviceHIonPhase1 = TrackSoAHeterogeneousDevice;
-
-} // namespace pixelTrack
-
-#endif // CUDADataFormats_Track_TrackHeterogeneousT_H
diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h
deleted file mode 100644
index 39e83491e1769..0000000000000
--- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h
+++ /dev/null
@@ -1,35 +0,0 @@
-#ifndef CUDADataFormats_Track_TrackHeterogeneousHost_H
-#define CUDADataFormats_Track_TrackHeterogeneousHost_H
-
-#include
-
-#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
-#include "CUDADataFormats/Common/interface/PortableHostCollection.h"
-
-// TODO: The class is created via inheritance of the PortableHostCollection.
-// This is generally discouraged, and should be done via composition.
-// See: https://github.com/cms-sw/cmssw/pull/40465#discussion_r1067364306
-template
-class TrackSoAHeterogeneousHost : public cms::cuda::PortableHostCollection> {
-public:
- static constexpr int32_t S = TrackerTraits::maxNumberOfTuples; //TODO: this could be made configurable at runtime
- explicit TrackSoAHeterogeneousHost() : cms::cuda::PortableHostCollection>(S) {}
-
- using cms::cuda::PortableHostCollection>::view;
- using cms::cuda::PortableHostCollection>::const_view;
- using cms::cuda::PortableHostCollection>::buffer;
- using cms::cuda::PortableHostCollection>::bufferSize;
-
- // Constructor which specifies the SoA size
- explicit TrackSoAHeterogeneousHost(cudaStream_t stream)
- : cms::cuda::PortableHostCollection>(S, stream) {}
-};
-
-namespace pixelTrack {
-
- using TrackSoAHostPhase1 = TrackSoAHeterogeneousHost;
- using TrackSoAHostPhase2 = TrackSoAHeterogeneousHost;
- using TrackSoAHostHIonPhase1 = TrackSoAHeterogeneousHost;
-} // namespace pixelTrack
-
-#endif // CUDADataFormats_Track_TrackHeterogeneousT_H
diff --git a/CUDADataFormats/Track/test/BuildFile.xml b/CUDADataFormats/Track/test/BuildFile.xml
deleted file mode 100644
index 32256c87ed577..0000000000000
--- a/CUDADataFormats/Track/test/BuildFile.xml
+++ /dev/null
@@ -1,22 +0,0 @@
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cpp b/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cpp
deleted file mode 100644
index dafa75e2e18d7..0000000000000
--- a/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cpp
+++ /dev/null
@@ -1,73 +0,0 @@
-/**
- Simple test for the pixelTrack::TrackSoA data structure
- which inherits from PortableDeviceCollection.
-
- Creates an instance of the class (automatically allocates
- memory on device), passes the view of the SoA data to
- the CUDA kernels which:
- - Fill the SoA with data.
- - Verify that the data written is correct.
-
- Then, the SoA data are copied back to Host, where
- a temporary host-side view (tmp_view) is created using
- the same Layout to access the data on host and print it.
- */
-
-#include
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h"
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
-
-#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
-
-namespace testTrackSoA {
-
- template
- void runKernels(TrackSoAView &tracks_view, cudaStream_t stream);
-}
-
-int main() {
- cms::cudatest::requireDevices();
-
- cudaStream_t stream;
- cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
-
- // Inner scope to deallocate memory before destroying the stream
- {
- // Instantiate tracks on device. PortableDeviceCollection allocates
- // SoA on device automatically.
- TrackSoAHeterogeneousDevice tracks_d(stream);
- testTrackSoA::runKernels(tracks_d.view(), stream);
-
- // Instantate tracks on host. This is where the data will be
- // copied to from device.
- TrackSoAHeterogeneousHost tracks_h(stream);
-
- cudaCheck(cudaMemcpyAsync(
- tracks_h.buffer().get(), tracks_d.const_buffer().get(), tracks_d.bufferSize(), cudaMemcpyDeviceToHost, stream));
- cudaCheck(cudaStreamSynchronize(stream));
-
- // Print results
- std::cout << "pt"
- << "\t"
- << "eta"
- << "\t"
- << "chi2"
- << "\t"
- << "quality"
- << "\t"
- << "nLayers"
- << "\t"
- << "hitIndices off" << std::endl;
-
- for (int i = 0; i < 10; ++i) {
- std::cout << tracks_h.view()[i].pt() << "\t" << tracks_h.view()[i].eta() << "\t" << tracks_h.view()[i].chi2()
- << "\t" << (int)tracks_h.view()[i].quality() << "\t" << (int)tracks_h.view()[i].nLayers() << "\t"
- << tracks_h.view().hitIndices().off[i] << std::endl;
- }
- }
- cudaCheck(cudaStreamDestroy(stream));
-
- return 0;
-}
diff --git a/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cu b/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cu
deleted file mode 100644
index 8e8595eb43e94..0000000000000
--- a/CUDADataFormats/Track/test/TrackSoAHeterogeneous_test.cu
+++ /dev/null
@@ -1,63 +0,0 @@
-#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/OneToManyAssoc.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
-
-namespace testTrackSoA {
-
- // Kernel which fills the TrackSoAView with data
- // to test writing to it
- template
- __global__ void fill(TrackSoAView tracks_view) {
- int i = threadIdx.x;
- if (i == 0) {
- tracks_view.nTracks() = 420;
- }
-
- for (int j = i; j < tracks_view.metadata().size(); j += blockDim.x) {
- tracks_view[j].pt() = (float)j;
- tracks_view[j].eta() = (float)j;
- tracks_view[j].chi2() = (float)j;
- tracks_view[j].quality() = (pixelTrack::Quality)(j % 256);
- tracks_view[j].nLayers() = j % 128;
- tracks_view.hitIndices().off[j] = j;
- }
- }
-
- // Kernel which reads from the TrackSoAView to verify
- // that it was written correctly from the fill kernel
- template
- __global__ void verify(TrackSoAConstView tracks_view) {
- int i = threadIdx.x;
-
- if (i == 0) {
- printf("SoA size: % d, block dims: % d\n", tracks_view.metadata().size(), blockDim.x);
- assert(tracks_view.nTracks() == 420);
- }
- for (int j = i; j < tracks_view.metadata().size(); j += blockDim.x) {
- assert(abs(tracks_view[j].pt() - (float)j) < .0001);
- assert(abs(tracks_view[j].eta() - (float)j) < .0001);
- assert(abs(tracks_view[j].chi2() - (float)j) < .0001);
- assert(tracks_view[j].quality() == (pixelTrack::Quality)(j % 256));
- assert(tracks_view[j].nLayers() == j % 128);
- assert(tracks_view.hitIndices().off[j] == j);
- }
- }
-
- // Host function which invokes the two kernels above
- template
- void runKernels(TrackSoAView& tracks_view, cudaStream_t stream) {
- fill<<<1, 1024, 0, stream>>>(tracks_view);
- cudaCheck(cudaGetLastError());
- cudaCheck(cudaDeviceSynchronize());
-
- verify<<<1, 1024, 0, stream>>>(tracks_view);
- cudaCheck(cudaGetLastError());
- cudaCheck(cudaDeviceSynchronize());
- }
-
- template void runKernels(TrackSoAView& tracks_view,
- cudaStream_t stream);
- template void runKernels(TrackSoAView& tracks_view,
- cudaStream_t stream);
-
-} // namespace testTrackSoA
diff --git a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cpp b/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cpp
deleted file mode 100644
index d6ff539a642b0..0000000000000
--- a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cpp
+++ /dev/null
@@ -1 +0,0 @@
-#include "TrajectoryStateSOA_t.h"
diff --git a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cu b/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cu
deleted file mode 100644
index d6ff539a642b0..0000000000000
--- a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.cu
+++ /dev/null
@@ -1 +0,0 @@
-#include "TrajectoryStateSOA_t.h"
diff --git a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.h b/CUDADataFormats/Track/test/TrajectoryStateSOA_t.h
deleted file mode 100644
index 6ba0eaa5c986e..0000000000000
--- a/CUDADataFormats/Track/test/TrajectoryStateSOA_t.h
+++ /dev/null
@@ -1,85 +0,0 @@
-#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
-#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h"
-
-using Vector5d = Eigen::Matrix;
-using Matrix5d = Eigen::Matrix;
-using helper = TracksUtilities;
-
-__host__ __device__ Matrix5d loadCov(Vector5d const& e) {
- Matrix5d cov;
- for (int i = 0; i < 5; ++i)
- cov(i, i) = e(i) * e(i);
- for (int i = 0; i < 5; ++i) {
- for (int j = 0; j < i; ++j) {
- double v = 0.3 * std::sqrt(cov(i, i) * cov(j, j)); // this makes the matrix pos defined
- cov(i, j) = (i + j) % 2 ? -0.4 * v : 0.1 * v;
- cov(j, i) = cov(i, j);
- }
- }
- return cov;
-}
-
-template
-__global__ void testTSSoA(TrackSoAView ts) {
- Vector5d par0;
- par0 << 0.2, 0.1, 3.5, 0.8, 0.1;
- Vector5d e0;
- e0 << 0.01, 0.01, 0.035, -0.03, -0.01;
- auto cov0 = loadCov(e0);
-
- int first = threadIdx.x + blockIdx.x * blockDim.x;
-
- for (int i = first; i < ts.metadata().size(); i += blockDim.x * gridDim.x) {
- helper::copyFromDense(ts, par0, cov0, i);
- Vector5d par1;
- Matrix5d cov1;
- helper::copyToDense(ts, par1, cov1, i);
- Vector5d delV = par1 - par0;
- Matrix5d delM = cov1 - cov0;
- for (int j = 0; j < 5; ++j) {
- assert(std::abs(delV(j)) < 1.e-5);
- for (auto k = j; k < 5; ++k) {
- assert(cov0(k, j) == cov0(j, k));
- assert(cov1(k, j) == cov1(j, k));
- assert(std::abs(delM(k, j)) < 1.e-5);
- }
- }
- }
-}
-
-#ifdef __CUDACC__
-#include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
-#endif
-
-int main() {
-#ifdef __CUDACC__
- cms::cudatest::requireDevices();
- cudaStream_t stream;
- cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
-#endif
-
-#ifdef __CUDACC__
- // Since we are going to copy data from ts_d to ts_h, we
- // need to initialize the Host collection with a stream.
- TrackSoAHeterogeneousHost ts_h(stream);
- TrackSoAHeterogeneousDevice ts_d(stream);
-#else
- // If CUDA is not available, Host collection must not be initialized
- // with a stream.
- TrackSoAHeterogeneousHost ts_h;
-#endif
-
-#ifdef __CUDACC__
- testTSSoA<<<1, 64, 0, stream>>>(ts_d.view());
- cudaCheck(cudaGetLastError());
- cudaCheck(cudaMemcpyAsync(
- ts_h.buffer().get(), ts_d.const_buffer().get(), ts_d.bufferSize(), cudaMemcpyDeviceToHost, stream));
- cudaCheck(cudaGetLastError());
- cudaCheck(cudaStreamSynchronize(stream));
-#else
- testTSSoA(ts_h.view());
-#endif
-}
diff --git a/CUDADataFormats/Vertex/BuildFile.xml b/CUDADataFormats/Vertex/BuildFile.xml
deleted file mode 100644
index c6b918ec4b12b..0000000000000
--- a/CUDADataFormats/Vertex/BuildFile.xml
+++ /dev/null
@@ -1,10 +0,0 @@
-
-
-
-
-
-
-
-
-
-
diff --git a/CUDADataFormats/Vertex/README.md b/CUDADataFormats/Vertex/README.md
deleted file mode 100644
index 3e495d15f776e..0000000000000
--- a/CUDADataFormats/Vertex/README.md
+++ /dev/null
@@ -1,45 +0,0 @@
-# Vertex CUDA Data Formats
-
-`CUDADataFormat`s meant to be used on Host (CPU) or Device (CUDA GPU) for
-storing information about vertices created during the Pixel-local Reconstruction
-chain. It stores data in an SoA manner. It contains the data that was previously
-contained in the deprecated `ZVertexSoA` class.
-
-The host format is inheriting from `CUDADataFormats/Common/interface/PortableHostCollection.h`,
-while the device format is inheriting from `CUDADataFormats/Common/interface/PortableDeviceCollection.h`
-
-Both formats use the same SoA Layout (`ZVertexSoAHeterogeneousLayout`) which is generated
-via the `GENERATE_SOA_LAYOUT` macro in the `ZVertexUtilities.h` file.
-
-## Notes
-
-- Initially, `ZVertexSoA` had distinct array sizes for each attribute (e.g. `zv` was `MAXVTX` elements
-long, `ndof` was `MAXTRACKS` elements long). All columns are now of uniform `MAXTRACKS` size,
-meaning that there will be some wasted space (appx. 190kB).
-- Host and Device classes should **not** be created via inheritance, as they're done here,
-but via composition. See [this discussion](https://github.com/cms-sw/cmssw/pull/40465#discussion_r1066039309).
-
-## ZVertexHeterogeneousHost
-
-The version of the data format to be used for storing vertex data on the CPU.
-Instances of this class are to be used for:
-
-- Having a place to copy data to host from device, via `cudaMemcpy`, or
-- Running host-side algorithms using data stored in an SoA manner.
-
-## ZVertexHeterogeneousDevice
-
-The version of the data format to be used for storing vertex data on the GPU.
-
-Instances of `ZVertexHeterogeneousDevice` are to be created on host and be
-used on device only. To do so, the instance's `view()` method is to be called
-to pass a `View` to any kernel launched. Accessing data from the `view()` is not
-possible on the host side.
-
-## Utilities
-
-Apart from `ZVertexSoAHeterogeneousLayout`, `ZVertexUtilities.h` also contains
-a collection of methods which were originally
-defined as class methods inside the `ZVertexSoA` class
-which have been adapted to operate on `View` instances, so that they are callable
-from within `__global__` kernels, on both CPU and CPU.
diff --git a/CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h b/CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h
deleted file mode 100644
index 417a960951fb1..0000000000000
--- a/CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h
+++ /dev/null
@@ -1,13 +0,0 @@
-#ifndef CUDADataFormatsVertexZVertexHeterogeneous_H
-#define CUDADataFormatsVertexZVertexHeterogeneous_H
-
-#include "CUDADataFormats/Vertex/interface/ZVertexSoA.h"
-#include "CUDADataFormats/Common/interface/HeterogeneousSoA.h"
-
-using ZVertexHeterogeneous = HeterogeneousSoA;
-#ifndef __CUDACC__
-#include "CUDADataFormats/Common/interface/Product.h"
-using ZVertexCUDAProduct = cms::cuda::Product;
-#endif
-
-#endif
diff --git a/CUDADataFormats/Vertex/interface/ZVertexSoA.h b/CUDADataFormats/Vertex/interface/ZVertexSoA.h
deleted file mode 100644
index 95106050f3d7a..0000000000000
--- a/CUDADataFormats/Vertex/interface/ZVertexSoA.h
+++ /dev/null
@@ -1,26 +0,0 @@
-#ifndef CUDADataFormats_Vertex_ZVertexSoA_h
-#define CUDADataFormats_Vertex_ZVertexSoA_h
-
-#include
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
-
-// SOA for vertices
-// These vertices are clusterized and fitted only along the beam line (z)
-// to obtain their global coordinate the beam spot position shall be added (eventually correcting for the beam angle as well)
-struct ZVertexSoA {
- static constexpr uint32_t MAXTRACKS = 128 * 1024;
- static constexpr uint32_t MAXVTX = 1024;
-
- int16_t idv[MAXTRACKS]; // vertex index for each associated (original) track (-1 == not associate)
- float zv[MAXVTX]; // output z-posistion of found vertices
- float wv[MAXVTX]; // output weight (1/error^2) on the above
- float chi2[MAXVTX]; // vertices chi2
- float ptv2[MAXVTX]; // vertices pt^2
- int32_t ndof[MAXTRACKS]; // vertices number of dof (reused as workspace for the number of nearest neighbours FIXME)
- uint16_t sortInd[MAXVTX]; // sorted index (by pt2) ascending
- uint32_t nvFinal; // the number of vertices
-
- __host__ __device__ void init() { nvFinal = 0; }
-};
-
-#endif // CUDADataFormats_Vertex_ZVertexSoA_h
diff --git a/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousDevice.h b/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousDevice.h
deleted file mode 100644
index ae662d7fd5f9a..0000000000000
--- a/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousDevice.h
+++ /dev/null
@@ -1,22 +0,0 @@
-#ifndef CUDADataFormats_Vertex_ZVertexHeterogeneousDevice_H
-#define CUDADataFormats_Vertex_ZVertexHeterogeneousDevice_H
-
-#include "CUDADataFormats/Vertex/interface/ZVertexUtilities.h"
-#include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"
-
-// TODO: The class is created via inheritance of the PortableDeviceCollection.
-// This is generally discouraged, and should be done via composition.
-// See: https://github.com/cms-sw/cmssw/pull/40465#discussion_r1067364306
-template
-class ZVertexSoAHeterogeneousDevice : public cms::cuda::PortableDeviceCollection> {
-public:
- ZVertexSoAHeterogeneousDevice() = default; // cms::cuda::Product needs this
-
- // Constructor which specifies the SoA size
- explicit ZVertexSoAHeterogeneousDevice(cudaStream_t stream)
- : PortableDeviceCollection>(S, stream) {}
-};
-
-using ZVertexSoADevice = ZVertexSoAHeterogeneousDevice;
-
-#endif // CUDADataFormats_Vertex_ZVertexHeterogeneousDevice_H
diff --git a/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousHost.h b/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousHost.h
deleted file mode 100644
index 6b62d615e1d11..0000000000000
--- a/CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousHost.h
+++ /dev/null
@@ -1,24 +0,0 @@
-#ifndef CUDADataFormats_Vertex_ZVertexHeterogeneousHost_H
-#define CUDADataFormats_Vertex_ZVertexHeterogeneousHost_H
-
-#include
-
-#include "CUDADataFormats/Vertex/interface/ZVertexUtilities.h"
-#include "CUDADataFormats/Common/interface/PortableHostCollection.h"
-
-// TODO: The class is created via inheritance of the PortableHostCollection.
-// This is generally discouraged, and should be done via composition.
-// See: https://github.com/cms-sw/cmssw/pull/40465#discussion_r1067364306
-template
-class ZVertexSoAHeterogeneousHost : public cms::cuda::PortableHostCollection> {
-public:
- explicit ZVertexSoAHeterogeneousHost() : cms::cuda::PortableHostCollection>(S) {}
-
- // Constructor which specifies the SoA size and CUDA stream
- explicit ZVertexSoAHeterogeneousHost(cudaStream_t stream)
- : PortableHostCollection>(S, stream) {}
-};
-
-using ZVertexSoAHost = ZVertexSoAHeterogeneousHost;
-
-#endif // CUDADataFormats_Vertex_ZVertexHeterogeneousHost_H
diff --git a/CUDADataFormats/Vertex/interface/ZVertexUtilities.h b/CUDADataFormats/Vertex/interface/ZVertexUtilities.h
deleted file mode 100644
index 2403652377971..0000000000000
--- a/CUDADataFormats/Vertex/interface/ZVertexUtilities.h
+++ /dev/null
@@ -1,35 +0,0 @@
-#ifndef CUDADataFormats_Vertex_ZVertexUtilities_h
-#define CUDADataFormats_Vertex_ZVertexUtilities_h
-
-#include
-#include "DataFormats/SoATemplate/interface/SoALayout.h"
-
-GENERATE_SOA_LAYOUT(ZVertexSoAHeterogeneousLayout,
- SOA_COLUMN(int16_t, idv),
- SOA_COLUMN(float, zv),
- SOA_COLUMN(float, wv),
- SOA_COLUMN(float, chi2),
- SOA_COLUMN(float, ptv2),
- SOA_COLUMN(int32_t, ndof),
- SOA_COLUMN(uint16_t, sortInd),
- SOA_SCALAR(uint32_t, nvFinal))
-
-// Previous ZVertexSoA class methods.
-// They operate on View and ConstView of the ZVertexSoA.
-namespace zVertex {
- // Common types for both Host and Device code
- using ZVertexSoALayout = ZVertexSoAHeterogeneousLayout<>;
- using ZVertexSoAView = ZVertexSoAHeterogeneousLayout<>::View;
- using ZVertexSoAConstView = ZVertexSoAHeterogeneousLayout<>::ConstView;
-
- namespace utilities {
-
- static constexpr uint32_t MAXTRACKS = 128 * 1024;
- static constexpr uint32_t MAXVTX = 1024;
-
- __host__ __device__ inline void init(ZVertexSoAView &vertices) { vertices.nvFinal() = 0; }
-
- } // namespace utilities
-} // namespace zVertex
-
-#endif
diff --git a/RecoTauTag/HLTProducers/BuildFile.xml b/RecoTauTag/HLTProducers/BuildFile.xml
index 79cdf947dd105..865bdb508e891 100644
--- a/RecoTauTag/HLTProducers/BuildFile.xml
+++ b/RecoTauTag/HLTProducers/BuildFile.xml
@@ -1,20 +1,18 @@
-
-
-
+
-
-
-
+
+
+
-
-
-
+
+
+
+
+
+
-
-
-
-
+
diff --git a/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc b/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc
deleted file mode 100644
index 051766a05606b..0000000000000
--- a/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc
+++ /dev/null
@@ -1,822 +0,0 @@
-/*
- * \class L2TauTagProducer
- *
- * L2Tau identification using Convolutional NN.
- *
- * \author Valeria D'Amante, Università di Siena and INFN Pisa
- * Konstantin Androsov, EPFL and ETHZ
-*/
-#include
-#include
-#include
-
-#include "FWCore/Framework/interface/stream/EDProducer.h"
-#include "FWCore/Framework/interface/ESHandle.h"
-#include "FWCore/Framework/interface/Event.h"
-#include "FWCore/Framework/interface/EventSetup.h"
-#include "FWCore/Framework/interface/Frameworkfwd.h"
-#include "FWCore/MessageLogger/interface/MessageLogger.h"
-#include "DataFormats/Math/interface/deltaR.h"
-#include "DataFormats/Common/interface/Handle.h"
-#include "FWCore/Utilities/interface/InputTag.h"
-#include "FWCore/Utilities/interface/isFinite.h"
-#include "FWCore/ParameterSet/interface/ParameterSet.h"
-#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
-#include "PhysicsTools/TensorFlow/interface/TensorFlow.h"
-#include "Geometry/CaloGeometry/interface/CaloCellGeometry.h"
-#include "Geometry/CaloGeometry/interface/CaloGeometry.h"
-#include "Geometry/CaloTopology/interface/HcalTopology.h"
-#include "Geometry/Records/interface/CaloGeometryRecord.h"
-#include "DataFormats/CaloRecHit/interface/CaloRecHit.h"
-#include "DataFormats/EcalRecHit/interface/EcalRecHit.h"
-#include "DataFormats/EcalRecHit/interface/EcalRecHitCollections.h"
-#include "DataFormats/EcalDetId/interface/EcalDetIdCollections.h"
-#include "DataFormats/HcalDetId/interface/HcalDetId.h"
-#include "DataFormats/HcalRecHit/interface/HBHERecHit.h"
-#include "DataFormats/HcalRecHit/interface/HcalRecHitDefs.h"
-#include "DataFormats/HcalRecHit/interface/HFRecHit.h"
-#include "DataFormats/HcalRecHit/interface/HORecHit.h"
-#include "DataFormats/HLTReco/interface/TriggerTypeDefs.h"
-#include "DataFormats/HLTReco/interface/TriggerFilterObjectWithRefs.h"
-#include "TrackingTools/TrajectoryParametrization/interface/CurvilinearTrajectoryError.h"
-#include "RecoTracker/PixelTrackFitting/interface/FitUtils.h"
-#include "TrackingTools/TrajectoryParametrization/interface/GlobalTrajectoryParameters.h"
-#include "DataFormats/TrackReco/interface/HitPattern.h"
-#include "TrackingTools/AnalyticalJacobians/interface/JacobianLocalToCurvilinear.h"
-#include "DataFormats/TrajectoryState/interface/LocalTrajectoryParameters.h"
-#include "DataFormats/GeometrySurface/interface/Plane.h"
-#include "DataFormats/BeamSpot/interface/BeamSpot.h"
-#include "MagneticField/Records/interface/IdealMagneticFieldRecord.h"
-
-#include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
-#include "CUDADataFormats/Vertex/interface/ZVertexUtilities.h"
-#include "CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousHost.h"
-
-namespace L2TauTagNNv1 {
- constexpr int nCellEta = 5;
- constexpr int nCellPhi = 5;
- constexpr int nVars = 31;
- constexpr float dR_max = 0.5;
- enum class NNInputs {
- nVertices = 0,
- l1Tau_pt,
- l1Tau_eta,
- l1Tau_hwIso,
- EcalEnergySum,
- EcalSize,
- EcalEnergyStdDev,
- EcalDeltaEta,
- EcalDeltaPhi,
- EcalChi2,
- EcalEnergySumForPositiveChi2,
- EcalSizeForPositiveChi2,
- HcalEnergySum,
- HcalSize,
- HcalEnergyStdDev,
- HcalDeltaEta,
- HcalDeltaPhi,
- HcalChi2,
- HcalEnergySumForPositiveChi2,
- HcalSizeForPositiveChi2,
- PatatrackPtSum,
- PatatrackSize,
- PatatrackSizeWithVertex,
- PatatrackPtSumWithVertex,
- PatatrackChargeSum,
- PatatrackDeltaEta,
- PatatrackDeltaPhi,
- PatatrackChi2OverNdof,
- PatatrackNdof,
- PatatrackDxy,
- PatatrackDz
- };
-
- const std::map varNameMap = {
- {NNInputs::nVertices, "nVertices"},
- {NNInputs::l1Tau_pt, "l1Tau_pt"},
- {NNInputs::l1Tau_eta, "l1Tau_eta"},
- {NNInputs::l1Tau_hwIso, "l1Tau_hwIso"},
- {NNInputs::EcalEnergySum, "EcalEnergySum"},
- {NNInputs::EcalSize, "EcalSize"},
- {NNInputs::EcalEnergyStdDev, "EcalEnergyStdDev"},
- {NNInputs::EcalDeltaEta, "EcalDeltaEta"},
- {NNInputs::EcalDeltaPhi, "EcalDeltaPhi"},
- {NNInputs::EcalChi2, "EcalChi2"},
- {NNInputs::EcalEnergySumForPositiveChi2, "EcalEnergySumForPositiveChi2"},
- {NNInputs::EcalSizeForPositiveChi2, "EcalSizeForPositiveChi2"},
- {NNInputs::HcalEnergySum, "HcalEnergySum"},
- {NNInputs::HcalSize, "HcalSize"},
- {NNInputs::HcalEnergyStdDev, "HcalEnergyStdDev"},
- {NNInputs::HcalDeltaEta, "HcalDeltaEta"},
- {NNInputs::HcalDeltaPhi, "HcalDeltaPhi"},
- {NNInputs::HcalChi2, "HcalChi2"},
- {NNInputs::HcalEnergySumForPositiveChi2, "HcalEnergySumForPositiveChi2"},
- {NNInputs::HcalSizeForPositiveChi2, "HcalSizeForPositiveChi2"},
- {NNInputs::PatatrackPtSum, "PatatrackPtSum"},
- {NNInputs::PatatrackSize, "PatatrackSize"},
- {NNInputs::PatatrackSizeWithVertex, "PatatrackSizeWithVertex"},
- {NNInputs::PatatrackPtSumWithVertex, "PatatrackPtSumWithVertex"},
- {NNInputs::PatatrackChargeSum, "PatatrackChargeSum"},
- {NNInputs::PatatrackDeltaEta, "PatatrackDeltaEta"},
- {NNInputs::PatatrackDeltaPhi, "PatatrackDeltaPhi"},
- {NNInputs::PatatrackChi2OverNdof, "PatatrackChi2OverNdof"},
- {NNInputs::PatatrackNdof, "PatatrackNdof"},
- {NNInputs::PatatrackDxy, "PatatrackDxy"},
- {NNInputs::PatatrackDz, "PatatrackDz"}};
-} // namespace L2TauTagNNv1
-namespace {
- inline float& getCellImpl(
- tensorflow::Tensor& cellGridMatrix, int tau_idx, int phi_idx, int eta_idx, L2TauTagNNv1::NNInputs NNInput_idx) {
- return cellGridMatrix.tensor()(tau_idx, phi_idx, eta_idx, static_cast(NNInput_idx));
- }
-} // namespace
-struct normDictElement {
- float mean;
- float std;
- float min;
- float max;
-};
-
-struct L2TauNNProducerCacheData {
- L2TauNNProducerCacheData() : graphDef(nullptr), session(nullptr) {}
- tensorflow::GraphDef* graphDef;
- tensorflow::Session* session;
- std::vector normVec;
-};
-
-class L2TauNNProducer : public edm::stream::EDProducer> {
-public:
- using TrackSoAHost = pixelTrack::TrackSoAHostPhase1;
-
- struct caloRecHitCollections {
- const HBHERecHitCollection* hbhe;
- const HORecHitCollection* ho;
- const EcalRecHitCollection* eb;
- const EcalRecHitCollection* ee;
- const CaloGeometry* geometry;
- };
-
- struct InputDescTau {
- std::string CollectionName;
- edm::EDGetTokenT inputToken_;
- };
-
- static constexpr float dR2_max = L2TauTagNNv1::dR_max * L2TauTagNNv1::dR_max;
- static constexpr float dEta_width = 2 * L2TauTagNNv1::dR_max / static_cast(L2TauTagNNv1::nCellEta);
- static constexpr float dPhi_width = 2 * L2TauTagNNv1::dR_max / static_cast(L2TauTagNNv1::nCellPhi);
-
- explicit L2TauNNProducer(const edm::ParameterSet&, const L2TauNNProducerCacheData*);
- static void fillDescriptions(edm::ConfigurationDescriptions&);
- static std::unique_ptr initializeGlobalCache(const edm::ParameterSet&);
- static void globalEndJob(L2TauNNProducerCacheData*);
-
-private:
- void checknan(tensorflow::Tensor& tensor, int debugLevel);
- void standardizeTensor(tensorflow::Tensor& tensor);
- std::vector getTauScore(const tensorflow::Tensor& cellGridMatrix);
- void produce(edm::Event& event, const edm::EventSetup& eventsetup) override;
- void fillL1TauVars(tensorflow::Tensor& cellGridMatrix, const std::vector& allTaus);
- void fillCaloRecHits(tensorflow::Tensor& cellGridMatrix,
- const std::vector& allTaus,
- const caloRecHitCollections& caloRecHits);
- void fillPatatracks(tensorflow::Tensor& cellGridMatrix,
- const std::vector& allTaus,
- const TrackSoAHost& patatracks_tsoa,
- const ZVertexSoAHost& patavtx_soa,
- const reco::BeamSpot& beamspot,
- const MagneticField* magfi);
- void selectGoodTracksAndVertices(const ZVertexSoAHost& patavtx_soa,
- const TrackSoAHost& patatracks_tsoa,
- std::vector& trkGood,
- std::vector& vtxGood);
-
- std::pair impactParameter(int it,
- const TrackSoAHost& patatracks_tsoa,
- float patatrackPhi,
- const reco::BeamSpot& beamspot,
- const MagneticField* magfi);
- template
- std::tuple getEtaPhiIndices(const VPos& position, const LVec& tau_p4);
- template
- std::tuple getEtaPhiIndices(float eta, float phi, const LVec& tau_p4);
-
-private:
- const int debugLevel_;
- const edm::EDGetTokenT tauTriggerToken_;
- std::vector L1TauDesc_;
- const edm::EDGetTokenT hbheToken_;
- const edm::EDGetTokenT hoToken_;
- const edm::EDGetTokenT ebToken_;
- const edm::EDGetTokenT eeToken_;
- const edm::ESGetToken geometryToken_;
- const edm::ESGetToken bFieldToken_;
- const edm::EDGetTokenT pataVerticesToken_;
- const edm::EDGetTokenT pataTracksToken_;
- const edm::EDGetTokenT beamSpotToken_;
- const unsigned int maxVtx_;
- const float fractionSumPt2_;
- const float minSumPt2_;
- const float trackPtMin_;
- const float trackPtMax_;
- const float trackChi2Max_;
- std::string inputTensorName_;
- std::string outputTensorName_;
- const L2TauNNProducerCacheData* L2cacheData_;
-};
-
-std::unique_ptr L2TauNNProducer::initializeGlobalCache(const edm::ParameterSet& cfg) {
- std::unique_ptr cacheData = std::make_unique();
- cacheData->normVec.reserve(L2TauTagNNv1::nVars);
-
- auto const graphPath = edm::FileInPath(cfg.getParameter("graphPath")).fullPath();
-
- cacheData->graphDef = tensorflow::loadGraphDef(graphPath);
- cacheData->session = tensorflow::createSession(cacheData->graphDef);
-
- boost::property_tree::ptree loadPtreeRoot;
- auto const normalizationDict = edm::FileInPath(cfg.getParameter("normalizationDict")).fullPath();
- boost::property_tree::read_json(normalizationDict, loadPtreeRoot);
- for (const auto& [key, val] : L2TauTagNNv1::varNameMap) {
- boost::property_tree::ptree var = loadPtreeRoot.get_child(val);
- normDictElement current_element;
- current_element.mean = var.get_child("mean").get_value();
- current_element.std = var.get_child("std").get_value();
- current_element.min = var.get_child("min").get_value();
- current_element.max = var.get_child("max").get_value();
- cacheData->normVec.push_back(current_element);
- }
- return cacheData;
-}
-void L2TauNNProducer::globalEndJob(L2TauNNProducerCacheData* cacheData) {
- if (cacheData->graphDef != nullptr) {
- delete cacheData->graphDef;
- }
- tensorflow::closeSession(cacheData->session);
-}
-void L2TauNNProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
- edm::ParameterSetDescription desc;
- desc.add("debugLevel", 0)->setComment("set debug level for printing out info");
- edm::ParameterSetDescription l1TausPset;
- l1TausPset.add("L1CollectionName", "DoubleTau")->setComment("Name of collections");
- l1TausPset.add("L1TauTrigger", edm::InputTag("hltL1sDoubleTauBigOR"))
- ->setComment("Which trigger should the L1 Taus collection pass");
- edm::ParameterSet l1TausPSetDefault;
- l1TausPSetDefault.addParameter("L1CollectionName", "DoubleTau");
- l1TausPSetDefault.addParameter("L1TauTrigger", edm::InputTag("hltL1sDoubleTauBigOR"));
- desc.addVPSet("L1Taus", l1TausPset, {l1TausPSetDefault});
- desc.add("hbheInput", edm::InputTag("hltHbhereco"))->setComment("HBHE recHit collection");
- desc.add("hoInput", edm::InputTag("hltHoreco"))->setComment("HO recHit Collection");
- desc.add("ebInput", edm::InputTag("hltEcalRecHit:EcalRecHitsEB"))->setComment("EB recHit Collection");
- desc.add("eeInput", edm::InputTag("hltEcalRecHit:EcalRecHitsEE"))->setComment("EE recHit Collection");
- desc.add("pataVertices", edm::InputTag("hltPixelVerticesSoA"))
- ->setComment("patatrack vertices collection");
- desc.add("pataTracks", edm::InputTag("hltPixelTracksSoA"))->setComment("patatrack collection");
- desc.add("BeamSpot", edm::InputTag("hltOnlineBeamSpot"))->setComment("BeamSpot Collection");
- desc.add("maxVtx", 100)->setComment("max output collection size (number of accepted vertices)");
- desc.add("fractionSumPt2", 0.3)->setComment("threshold on sumPt2 fraction of the leading vertex");
- desc.add("minSumPt2", 0.)->setComment("min sumPt2");
- desc.add("track_pt_min", 1.0)->setComment("min track p_T");
- desc.add("track_pt_max", 10.0)->setComment("max track p_T");
- desc.add("track_chi2_max", 99999.)->setComment("max track chi2");
- desc.add("graphPath", "RecoTauTag/TrainingFiles/data/L2TauNNTag/L2TauTag_Run3v1.pb")
- ->setComment("path to the saved CNN");
- desc.add("normalizationDict", "RecoTauTag/TrainingFiles/data/L2TauNNTag/NormalizationDict.json")
- ->setComment("path to the dictionary for variable standardization");
- descriptions.addWithDefaultLabel(desc);
-}
-
-L2TauNNProducer::L2TauNNProducer(const edm::ParameterSet& cfg, const L2TauNNProducerCacheData* cacheData)
- : debugLevel_(cfg.getParameter("debugLevel")),
- hbheToken_(consumes(cfg.getParameter("hbheInput"))),
- hoToken_(consumes(cfg.getParameter("hoInput"))),
- ebToken_(consumes(cfg.getParameter("ebInput"))),
- eeToken_(consumes(cfg.getParameter("eeInput"))),
- geometryToken_(esConsumes()),
- bFieldToken_(esConsumes()),
- pataVerticesToken_(consumes(cfg.getParameter("pataVertices"))),
- pataTracksToken_(consumes(cfg.getParameter("pataTracks"))),
- beamSpotToken_(consumes(cfg.getParameter("BeamSpot"))),
- maxVtx_(cfg.getParameter("maxVtx")),
- fractionSumPt2_(cfg.getParameter("fractionSumPt2")),
- minSumPt2_(cfg.getParameter("minSumPt2")),
- trackPtMin_(cfg.getParameter("track_pt_min")),
- trackPtMax_(cfg.getParameter("track_pt_max")),
- trackChi2Max_(cfg.getParameter("track_chi2_max")) {
- if (cacheData->graphDef == nullptr) {
- throw cms::Exception("InvalidCacheData") << "Invalid Cache Data.";
- }
- inputTensorName_ = cacheData->graphDef->node(0).name();
- outputTensorName_ = cacheData->graphDef->node(cacheData->graphDef->node_size() - 1).name();
- L2cacheData_ = cacheData;
- std::vector L1TauCollections = cfg.getParameter>("L1Taus");
- L1TauDesc_.reserve(L1TauCollections.size());
- for (const auto& l1TauInput : L1TauCollections) {
- InputDescTau toInsert;
- toInsert.CollectionName = l1TauInput.getParameter("L1CollectionName");
- toInsert.inputToken_ =
- consumes(l1TauInput.getParameter("L1TauTrigger"));
- L1TauDesc_.push_back(toInsert);
- }
- for (const auto& desc : L1TauDesc_)
- produces>(desc.CollectionName);
-}
-
-void L2TauNNProducer::checknan(tensorflow::Tensor& tensor, int debugLevel) {
- using NNInputs = L2TauTagNNv1::NNInputs;
- std::vector tensor_shape(tensor.shape().dims());
- for (int d = 0; d < tensor.shape().dims(); d++) {
- tensor_shape.at(d) = tensor.shape().dim_size(d);
- }
- if (tensor_shape.size() != 4) {
- throw cms::Exception("InvalidTensor") << "Tensor shape does not have 4 dimensions!";
- }
- for (int tau_idx = 0; tau_idx < tensor_shape.at(0); tau_idx++) {
- for (int phi_idx = 0; phi_idx < tensor_shape.at(1); phi_idx++) {
- for (int eta_idx = 0; eta_idx < tensor_shape.at(2); eta_idx++) {
- for (int var_idx = 0; var_idx < tensor_shape.at(3); var_idx++) {
- auto getCell = [&](NNInputs input) -> float& {
- return getCellImpl(tensor, tau_idx, phi_idx, eta_idx, input);
- };
- auto nonstd_var = getCell(static_cast(var_idx));
- if (edm::isNotFinite(nonstd_var)) {
- edm::LogWarning("InputVar") << "var is nan \nvar name= "
- << L2TauTagNNv1::varNameMap.at(static_cast(var_idx))
- << "\t var_idx = " << var_idx << "\t eta_idx = " << eta_idx
- << "\t phi_idx = " << phi_idx << "\t tau_idx = " << tau_idx;
- if (debugLevel > 2) {
- edm::LogWarning("InputVar") << "other vars in same cell \n";
- if (var_idx + 1 < tensor_shape.at(3))
- edm::LogWarning("InputVar") << L2TauTagNNv1::varNameMap.at(static_cast(var_idx + 1))
- << "\t = " << getCell(static_cast(var_idx + 1));
- if (var_idx + 2 < tensor_shape.at(3))
- edm::LogWarning("InputVar") << L2TauTagNNv1::varNameMap.at(static_cast(var_idx + 2))
- << "\t = " << getCell(static_cast(var_idx + 2));
- if (var_idx + 3 < tensor_shape.at(3))
- edm::LogWarning("InputVar") << L2TauTagNNv1::varNameMap.at(static_cast(var_idx + 3))
- << "\t = " << getCell(static_cast(var_idx + 3));
- if (var_idx + 4 < tensor_shape.at(3))
- edm::LogWarning("InputVar") << L2TauTagNNv1::varNameMap.at(static_cast(var_idx + 4))
- << "\t = " << getCell(static_cast(var_idx + 4));
- }
- }
- }
- }
- }
- }
-}
-
-void L2TauNNProducer::standardizeTensor(tensorflow::Tensor& tensor) {
- using NNInputs = L2TauTagNNv1::NNInputs;
- std::vector tensor_shape(tensor.shape().dims());
- for (int d = 0; d < tensor.shape().dims(); d++) {
- tensor_shape.at(d) = tensor.shape().dim_size(d);
- }
- if (tensor_shape.size() != 4) {
- throw cms::Exception("InvalidTensor") << "Tensor shape does not have 4 dimensions!";
- }
- for (int tau_idx = 0; tau_idx < tensor_shape.at(0); tau_idx++) {
- for (int phi_idx = 0; phi_idx < tensor_shape.at(1); phi_idx++) {
- for (int eta_idx = 0; eta_idx < tensor_shape.at(2); eta_idx++) {
- for (int var_idx = 0; var_idx < tensor_shape.at(3); var_idx++) {
- auto getCell = [&](NNInputs input) -> float& {
- return getCellImpl(tensor, tau_idx, phi_idx, eta_idx, input);
- };
- float mean = L2cacheData_->normVec.at(var_idx).mean;
- float std = L2cacheData_->normVec.at(var_idx).std;
- float min = L2cacheData_->normVec.at(var_idx).min;
- float max = L2cacheData_->normVec.at(var_idx).max;
- float nonstd_var = getCell(static_cast(var_idx));
- float std_var = static_cast((nonstd_var - mean) / std);
- if (std_var > max) {
- std_var = static_cast(max);
- } else if (std_var < min) {
- std_var = static_cast(min);
- }
- getCell(static_cast(var_idx)) = std_var;
- }
- }
- }
- }
-}
-
-void L2TauNNProducer::fillL1TauVars(tensorflow::Tensor& cellGridMatrix, const std::vector& allTaus) {
- using NNInputs = L2TauTagNNv1::NNInputs;
-
- const int nTaus = allTaus.size();
- for (int tau_idx = 0; tau_idx < nTaus; tau_idx++) {
- for (int eta_idx = 0; eta_idx < L2TauTagNNv1::nCellEta; eta_idx++) {
- for (int phi_idx = 0; phi_idx < L2TauTagNNv1::nCellPhi; phi_idx++) {
- auto getCell = [&](NNInputs input) -> float& {
- return getCellImpl(cellGridMatrix, tau_idx, phi_idx, eta_idx, input);
- };
- getCell(NNInputs::l1Tau_pt) = allTaus[tau_idx]->pt();
- getCell(NNInputs::l1Tau_eta) = allTaus[tau_idx]->eta();
- getCell(NNInputs::l1Tau_hwIso) = allTaus[tau_idx]->hwIso();
- }
- }
- }
-}
-
-template
-std::tuple L2TauNNProducer::getEtaPhiIndices(float eta, float phi, const LVec& tau_p4) {
- const float deta = eta - tau_p4.eta();
- const float dphi = reco::deltaPhi(phi, tau_p4.phi());
- const int eta_idx = static_cast(floor((deta + L2TauTagNNv1::dR_max) / dEta_width));
- const int phi_idx = static_cast(floor((dphi + L2TauTagNNv1::dR_max) / dPhi_width));
- return std::make_tuple(deta, dphi, eta_idx, phi_idx);
-}
-
-template
-std::tuple L2TauNNProducer::getEtaPhiIndices(const VPos& position, const LVec& tau_p4) {
- return getEtaPhiIndices(position.eta(), position.phi(), tau_p4);
-}
-
-void L2TauNNProducer::fillCaloRecHits(tensorflow::Tensor& cellGridMatrix,
- const std::vector& allTaus,
- const caloRecHitCollections& caloRecHits) {
- using NNInputs = L2TauTagNNv1::NNInputs;
-
- const int nTaus = allTaus.size();
- float deta, dphi;
- int eta_idx = 0;
- int phi_idx = 0;
- int tau_idx = 0;
-
- auto getCell = [&](NNInputs input) -> float& {
- return getCellImpl(cellGridMatrix, tau_idx, phi_idx, eta_idx, input);
- };
- for (tau_idx = 0; tau_idx < nTaus; tau_idx++) {
- // calorechit_EE
- for (const auto& caloRecHit_ee : *caloRecHits.ee) {
- if (caloRecHit_ee.energy() <= 0)
- continue;
- const auto& position = caloRecHits.geometry->getGeometry(caloRecHit_ee.id())->getPosition();
- const float eeCalEn = caloRecHit_ee.energy();
- const float eeCalChi2 = caloRecHit_ee.chi2();
- if (reco::deltaR2(position, allTaus[tau_idx]->polarP4()) < dR2_max) {
- std::tie(deta, dphi, eta_idx, phi_idx) = getEtaPhiIndices(position, allTaus[tau_idx]->polarP4());
- getCell(NNInputs::EcalEnergySum) += eeCalEn;
- getCell(NNInputs::EcalSize) += 1.;
- getCell(NNInputs::EcalEnergyStdDev) += eeCalEn * eeCalEn;
- getCell(NNInputs::EcalDeltaEta) += deta * eeCalEn;
- getCell(NNInputs::EcalDeltaPhi) += dphi * eeCalEn;
- if (eeCalChi2 >= 0) {
- getCell(NNInputs::EcalChi2) += eeCalChi2 * eeCalEn;
- getCell(NNInputs::EcalEnergySumForPositiveChi2) += eeCalEn;
- getCell(NNInputs::EcalSizeForPositiveChi2) += 1.;
- }
- }
- }
-
- // calorechit_EB
- for (const auto& caloRecHit_eb : *caloRecHits.eb) {
- if (caloRecHit_eb.energy() <= 0)
- continue;
- const auto& position = caloRecHits.geometry->getGeometry(caloRecHit_eb.id())->getPosition();
- const float ebCalEn = caloRecHit_eb.energy();
- const float ebCalChi2 = caloRecHit_eb.chi2();
- if (reco::deltaR2(position, allTaus[tau_idx]->polarP4()) < dR2_max) {
- std::tie(deta, dphi, eta_idx, phi_idx) = getEtaPhiIndices(position, allTaus[tau_idx]->polarP4());
- getCell(NNInputs::EcalEnergySum) += ebCalEn;
- getCell(NNInputs::EcalSize) += 1.;
- getCell(NNInputs::EcalEnergyStdDev) += ebCalEn * ebCalEn;
- getCell(NNInputs::EcalDeltaEta) += deta * ebCalEn;
- getCell(NNInputs::EcalDeltaPhi) += dphi * ebCalEn;
- if (ebCalChi2 >= 0) {
- getCell(NNInputs::EcalChi2) += ebCalChi2 * ebCalEn;
- getCell(NNInputs::EcalEnergySumForPositiveChi2) += ebCalEn;
- getCell(NNInputs::EcalSizeForPositiveChi2) += 1.;
- }
- }
- }
-
- // calorechit_HBHE
- for (const auto& caloRecHit_hbhe : *caloRecHits.hbhe) {
- if (caloRecHit_hbhe.energy() <= 0)
- continue;
- const auto& position = caloRecHits.geometry->getGeometry(caloRecHit_hbhe.id())->getPosition();
- const float hbheCalEn = caloRecHit_hbhe.energy();
- const float hbheCalChi2 = caloRecHit_hbhe.chi2();
- if (reco::deltaR2(position, allTaus[tau_idx]->polarP4()) < dR2_max) {
- std::tie(deta, dphi, eta_idx, phi_idx) = getEtaPhiIndices(position, allTaus[tau_idx]->polarP4());
- getCell(NNInputs::HcalEnergySum) += hbheCalEn;
- getCell(NNInputs::HcalEnergyStdDev) += hbheCalEn * hbheCalEn;
- getCell(NNInputs::HcalSize) += 1.;
- getCell(NNInputs::HcalDeltaEta) += deta * hbheCalEn;
- getCell(NNInputs::HcalDeltaPhi) += dphi * hbheCalEn;
- if (hbheCalChi2 >= 0) {
- getCell(NNInputs::HcalChi2) += hbheCalChi2 * hbheCalEn;
- getCell(NNInputs::HcalEnergySumForPositiveChi2) += hbheCalEn;
- getCell(NNInputs::HcalSizeForPositiveChi2) += 1.;
- }
- }
- }
-
- // calorechit_HO
- for (const auto& caloRecHit_ho : *caloRecHits.ho) {
- if (caloRecHit_ho.energy() <= 0)
- continue;
- const auto& position = caloRecHits.geometry->getGeometry(caloRecHit_ho.id())->getPosition();
- const float hoCalEn = caloRecHit_ho.energy();
- if (reco::deltaR2(position, allTaus[tau_idx]->polarP4()) < dR2_max) {
- std::tie(deta, dphi, eta_idx, phi_idx) = getEtaPhiIndices(position, allTaus[tau_idx]->polarP4());
- getCell(NNInputs::HcalEnergySum) += hoCalEn;
- getCell(NNInputs::HcalEnergyStdDev) += hoCalEn * hoCalEn;
- getCell(NNInputs::HcalSize) += 1.;
- getCell(NNInputs::HcalDeltaEta) += deta * hoCalEn;
- getCell(NNInputs::HcalDeltaPhi) += dphi * hoCalEn;
- }
- }
-
- // normalize to sum and define stdDev
- for (eta_idx = 0; eta_idx < L2TauTagNNv1::nCellEta; eta_idx++) {
- for (phi_idx = 0; phi_idx < L2TauTagNNv1::nCellPhi; phi_idx++) {
- /* normalize eCal vars*/
- if (getCell(NNInputs::EcalEnergySum) > 0.) {
- getCell(NNInputs::EcalDeltaEta) /= getCell(NNInputs::EcalEnergySum);
- getCell(NNInputs::EcalDeltaPhi) /= getCell(NNInputs::EcalEnergySum);
- }
- if (getCell(NNInputs::EcalEnergySumForPositiveChi2) > 0.) {
- getCell(NNInputs::EcalChi2) /= getCell(NNInputs::EcalEnergySumForPositiveChi2);
- }
- if (getCell(NNInputs::EcalSize) > 1.) {
- // (stdDev - (enSum*enSum)/size) / (size-1)
- getCell(NNInputs::EcalEnergyStdDev) =
- (getCell(NNInputs::EcalEnergyStdDev) -
- (getCell(NNInputs::EcalEnergySum) * getCell(NNInputs::EcalEnergySum)) / getCell(NNInputs::EcalSize)) /
- (getCell(NNInputs::EcalSize) - 1);
- } else {
- getCell(NNInputs::EcalEnergyStdDev) = 0.;
- }
- /* normalize hCal Vars */
- if (getCell(NNInputs::HcalEnergySum) > 0.) {
- getCell(NNInputs::HcalDeltaEta) /= getCell(NNInputs::HcalEnergySum);
- getCell(NNInputs::HcalDeltaPhi) /= getCell(NNInputs::HcalEnergySum);
- }
- if (getCell(NNInputs::HcalEnergySumForPositiveChi2) > 0.) {
- getCell(NNInputs::HcalChi2) /= getCell(NNInputs::HcalEnergySumForPositiveChi2);
- }
- if (getCell(NNInputs::HcalSize) > 1.) {
- // (stdDev - (enSum*enSum)/size) / (size-1)
- getCell(NNInputs::HcalEnergyStdDev) =
- (getCell(NNInputs::HcalEnergyStdDev) -
- (getCell(NNInputs::HcalEnergySum) * getCell(NNInputs::HcalEnergySum)) / getCell(NNInputs::HcalSize)) /
- (getCell(NNInputs::HcalSize) - 1);
- } else {
- getCell(NNInputs::HcalEnergyStdDev) = 0.;
- }
- }
- }
- }
-}
-
-void L2TauNNProducer::selectGoodTracksAndVertices(const ZVertexSoAHost& patavtx_soa,
- const TrackSoAHost& patatracks_tsoa,
- std::vector& trkGood,
- std::vector& vtxGood) {
- using patatrackHelpers = TracksUtilities;
- const auto maxTracks = patatracks_tsoa.view().metadata().size();
- const int nv = patavtx_soa.view().nvFinal();
- trkGood.clear();
- trkGood.reserve(maxTracks);
- vtxGood.clear();
- vtxGood.reserve(nv);
- auto const quality = patatracks_tsoa.view().quality();
-
- // No need to sort either as the algorithms is just using the max (not even the location, just the max value of pt2sum).
- std::vector pTSquaredSum(nv, 0);
- std::vector nTrkAssociated(nv, 0);
-
- for (int32_t trk_idx = 0; trk_idx < maxTracks; ++trk_idx) {
- auto nHits = patatrackHelpers::nHits(patatracks_tsoa.view(), trk_idx);
- if (nHits == 0) {
- break;
- }
- int vtx_ass_to_track = patavtx_soa.view()[trk_idx].idv();
- if (vtx_ass_to_track >= 0 && vtx_ass_to_track < nv) {
- auto patatrackPt = patatracks_tsoa.view()[trk_idx].pt();
- ++nTrkAssociated[vtx_ass_to_track];
- if (patatrackPt >= trackPtMin_ && patatracks_tsoa.const_view()[trk_idx].chi2() <= trackChi2Max_) {
- patatrackPt = std::min(patatrackPt, trackPtMax_);
- pTSquaredSum[vtx_ass_to_track] += patatrackPt * patatrackPt;
- }
- }
- if (nHits > 0 and quality[trk_idx] >= pixelTrack::Quality::loose) {
- trkGood.push_back(trk_idx);
- }
- }
- if (nv > 0) {
- const auto minFOM_fromFrac = (*std::max_element(pTSquaredSum.begin(), pTSquaredSum.end())) * fractionSumPt2_;
- for (int j = nv - 1; j >= 0 && vtxGood.size() < maxVtx_; --j) {
- auto vtx_idx = patavtx_soa.view()[j].sortInd();
- assert(vtx_idx < nv);
- if (nTrkAssociated[vtx_idx] >= 2 && pTSquaredSum[vtx_idx] >= minFOM_fromFrac &&
- pTSquaredSum[vtx_idx] > minSumPt2_) {
- vtxGood.push_back(vtx_idx);
- }
- }
- }
-}
-
-std::pair L2TauNNProducer::impactParameter(int it,
- const TrackSoAHost& patatracks_tsoa,
- float patatrackPhi,
- const reco::BeamSpot& beamspot,
- const MagneticField* magfi) {
- /* dxy and dz */
- riemannFit::Vector5d ipar, opar;
- riemannFit::Matrix5d icov, ocov;
- TracksUtilities::copyToDense(patatracks_tsoa.view(), ipar, icov, it);
- riemannFit::transformToPerigeePlane(ipar, icov, opar, ocov);
- LocalTrajectoryParameters lpar(opar(0), opar(1), opar(2), opar(3), opar(4), 1.);
- float sp = std::sin(patatrackPhi);
- float cp = std::cos(patatrackPhi);
- Surface::RotationType Rotation(sp, -cp, 0, 0, 0, -1.f, cp, sp, 0);
- GlobalPoint BeamSpotPoint(beamspot.x0(), beamspot.y0(), beamspot.z0());
- Plane impPointPlane(BeamSpotPoint, Rotation);
- GlobalTrajectoryParameters gp(
- impPointPlane.toGlobal(lpar.position()), impPointPlane.toGlobal(lpar.momentum()), lpar.charge(), magfi);
- GlobalPoint vv = gp.position();
- math::XYZPoint pos(vv.x(), vv.y(), vv.z());
- GlobalVector pp = gp.momentum();
- math::XYZVector mom(pp.x(), pp.y(), pp.z());
- auto lambda = M_PI_2 - pp.theta();
- auto phi = pp.phi();
- float patatrackDxy = -vv.x() * std::sin(phi) + vv.y() * std::cos(phi);
- float patatrackDz =
- (vv.z() * std::cos(lambda) - (vv.x() * std::cos(phi) + vv.y() * std::sin(phi)) * std::sin(lambda)) /
- std::cos(lambda);
- return std::make_pair(patatrackDxy, patatrackDz);
-}
-
-void L2TauNNProducer::fillPatatracks(tensorflow::Tensor& cellGridMatrix,
- const std::vector& allTaus,
- const TrackSoAHost& patatracks_tsoa,
- const ZVertexSoAHost& patavtx_soa,
- const reco::BeamSpot& beamspot,
- const MagneticField* magfi) {
- using NNInputs = L2TauTagNNv1::NNInputs;
- using patatrackHelpers = TracksUtilities;
- float deta, dphi;
- int eta_idx = 0;
- int phi_idx = 0;
- int tau_idx = 0;
-
- auto getCell = [&](NNInputs input) -> float& {
- return getCellImpl(cellGridMatrix, tau_idx, phi_idx, eta_idx, input);
- };
-
- std::vector trkGood;
- std::vector vtxGood;
-
- selectGoodTracksAndVertices(patavtx_soa, patatracks_tsoa, trkGood, vtxGood);
-
- const int nTaus = allTaus.size();
- for (tau_idx = 0; tau_idx < nTaus; tau_idx++) {
- const float tauEta = allTaus[tau_idx]->eta();
- const float tauPhi = allTaus[tau_idx]->phi();
-
- for (const auto it : trkGood) {
- const float patatrackPt = patatracks_tsoa.const_view()[it].pt();
- if (patatrackPt <= 0)
- continue;
- const float patatrackPhi = patatrackHelpers::phi(patatracks_tsoa.const_view(), it);
- const float patatrackEta = patatracks_tsoa.const_view()[it].eta();
- const float patatrackCharge = patatrackHelpers::charge(patatracks_tsoa.const_view(), it);
- const float patatrackChi2OverNdof = patatracks_tsoa.view()[it].chi2();
- const auto nHits = patatrackHelpers::nHits(patatracks_tsoa.const_view(), it);
- if (nHits <= 0)
- continue;
- const int patatrackNdof = 2 * std::min(6, nHits) - 5;
-
- const int vtx_idx_assTrk = patavtx_soa.view()[it].idv();
- if (reco::deltaR2(patatrackEta, patatrackPhi, tauEta, tauPhi) < dR2_max) {
- std::tie(deta, dphi, eta_idx, phi_idx) =
- getEtaPhiIndices(patatrackEta, patatrackPhi, allTaus[tau_idx]->polarP4());
- getCell(NNInputs::PatatrackPtSum) += patatrackPt;
- getCell(NNInputs::PatatrackSize) += 1.;
- getCell(NNInputs::PatatrackChargeSum) += patatrackCharge;
- getCell(NNInputs::PatatrackDeltaEta) += deta * patatrackPt;
- getCell(NNInputs::PatatrackDeltaPhi) += dphi * patatrackPt;
- getCell(NNInputs::PatatrackChi2OverNdof) += patatrackChi2OverNdof * patatrackPt;
- getCell(NNInputs::PatatrackNdof) += patatrackNdof * patatrackPt;
- std::pair impactParameters = impactParameter(it, patatracks_tsoa, patatrackPhi, beamspot, magfi);
- getCell(NNInputs::PatatrackDxy) += impactParameters.first * patatrackPt;
- getCell(NNInputs::PatatrackDz) += impactParameters.second * patatrackPt;
- if ((std::find(vtxGood.begin(), vtxGood.end(), vtx_idx_assTrk) != vtxGood.end())) {
- getCell(NNInputs::PatatrackPtSumWithVertex) += patatrackPt;
- getCell(NNInputs::PatatrackSizeWithVertex) += 1.;
- }
- }
- }
-
- // normalize to sum and define stdDev
- for (eta_idx = 0; eta_idx < L2TauTagNNv1::nCellEta; eta_idx++) {
- for (phi_idx = 0; phi_idx < L2TauTagNNv1::nCellPhi; phi_idx++) {
- getCell(NNInputs::nVertices) = vtxGood.size();
- if (getCell(NNInputs::PatatrackPtSum) > 0.) {
- getCell(NNInputs::PatatrackDeltaEta) /= getCell(NNInputs::PatatrackPtSum);
- getCell(NNInputs::PatatrackDeltaPhi) /= getCell(NNInputs::PatatrackPtSum);
- getCell(NNInputs::PatatrackChi2OverNdof) /= getCell(NNInputs::PatatrackPtSum);
- getCell(NNInputs::PatatrackNdof) /= getCell(NNInputs::PatatrackPtSum);
- getCell(NNInputs::PatatrackDxy) /= getCell(NNInputs::PatatrackPtSum);
- getCell(NNInputs::PatatrackDz) /= getCell(NNInputs::PatatrackPtSum);
- }
- }
- }
- }
-}
-
-std::vector L2TauNNProducer::getTauScore(const tensorflow::Tensor& cellGridMatrix) {
- const int nTau = cellGridMatrix.shape().dim_size(0);
- std::vector pred_vector(nTau);
- if (nTau > 0) {
- // Only run the inference if there are taus to process
- std::vector pred_tensor;
- tensorflow::run(L2cacheData_->session, {{inputTensorName_, cellGridMatrix}}, {outputTensorName_}, &pred_tensor);
- for (int tau_idx = 0; tau_idx < nTau; ++tau_idx) {
- pred_vector[tau_idx] = pred_tensor[0].matrix()(tau_idx, 0);
- }
- }
- return pred_vector;
-}
-
-void L2TauNNProducer::produce(edm::Event& event, const edm::EventSetup& eventsetup) {
- std::vector> TauCollectionMap(L1TauDesc_.size());
- l1t::TauVectorRef allTaus;
-
- for (size_t inp_idx = 0; inp_idx < L1TauDesc_.size(); inp_idx++) {
- l1t::TauVectorRef l1Taus;
- auto const& l1TriggeredTaus = event.get(L1TauDesc_[inp_idx].inputToken_);
- l1TriggeredTaus.getObjects(trigger::TriggerL1Tau, l1Taus);
- TauCollectionMap.at(inp_idx).resize(l1Taus.size());
-
- for (size_t l1_idx = 0; l1_idx < l1Taus.size(); l1_idx++) {
- size_t tau_idx;
- const auto iter = std::find(allTaus.begin(), allTaus.end(), l1Taus[l1_idx]);
- if (iter != allTaus.end()) {
- tau_idx = std::distance(allTaus.begin(), iter);
- } else {
- allTaus.push_back(l1Taus[l1_idx]);
- tau_idx = allTaus.size() - 1;
- }
- TauCollectionMap.at(inp_idx).at(l1_idx) = tau_idx;
- }
- }
- const auto ebCal = event.getHandle(ebToken_);
- const auto eeCal = event.getHandle(eeToken_);
- const auto hbhe = event.getHandle(hbheToken_);
- const auto ho = event.getHandle(hoToken_);
- auto const& patatracks_SoA = event.get(pataTracksToken_);
- auto const& vertices_SoA = event.get(pataVerticesToken_);
- const auto bsHandle = event.getHandle(beamSpotToken_);
-
- auto const fieldESH = eventsetup.getHandle(bFieldToken_);
- auto const geometry = eventsetup.getHandle(geometryToken_);
-
- caloRecHitCollections caloRecHits;
- caloRecHits.hbhe = &*hbhe;
- caloRecHits.ho = &*ho;
- caloRecHits.eb = &*ebCal;
- caloRecHits.ee = &*eeCal;
- caloRecHits.geometry = &*geometry;
-
- const int nTaus = allTaus.size();
- tensorflow::Tensor cellGridMatrix(tensorflow::DT_FLOAT,
- {nTaus, L2TauTagNNv1::nCellEta, L2TauTagNNv1::nCellPhi, L2TauTagNNv1::nVars});
- const int n_inputs = nTaus * L2TauTagNNv1::nCellEta * L2TauTagNNv1::nCellPhi * L2TauTagNNv1::nVars;
- for (int input_idx = 0; input_idx < n_inputs; ++input_idx) {
- cellGridMatrix.flat()(input_idx) = 0;
- }
- fillL1TauVars(cellGridMatrix, allTaus);
-
- fillCaloRecHits(cellGridMatrix, allTaus, caloRecHits);
-
- fillPatatracks(cellGridMatrix, allTaus, patatracks_SoA, vertices_SoA, *bsHandle, fieldESH.product());
-
- standardizeTensor(cellGridMatrix);
-
- if (debugLevel_ > 0) {
- checknan(cellGridMatrix, debugLevel_);
- }
-
- std::vector tau_score = getTauScore(cellGridMatrix);
-
- for (size_t inp_idx = 0; inp_idx < L1TauDesc_.size(); inp_idx++) {
- const size_t nTau = TauCollectionMap[inp_idx].size();
- auto tau_tags = std::make_unique>(nTau);
- for (size_t tau_pos = 0; tau_pos < nTau; ++tau_pos) {
- const auto tau_idx = TauCollectionMap[inp_idx][tau_pos];
- if (debugLevel_ > 0) {
- edm::LogInfo("DebugInfo") << event.id().event() << " \t " << (allTaus[tau_idx])->pt() << " \t "
- << tau_score.at(tau_idx) << std::endl;
- }
- (*tau_tags)[tau_pos] = tau_score.at(tau_idx);
- }
- event.put(std::move(tau_tags), L1TauDesc_[inp_idx].CollectionName);
- }
-}
-//define this as a plug-in
-#include "FWCore/Framework/interface/MakerMacros.h"
-DEFINE_FWK_MODULE(L2TauNNProducer);
diff --git a/RecoTracker/PixelTrackFitting/plugins/BuildFile.xml b/RecoTracker/PixelTrackFitting/plugins/BuildFile.xml
index 6c8c102293651..eae412ec800cc 100644
--- a/RecoTracker/PixelTrackFitting/plugins/BuildFile.xml
+++ b/RecoTracker/PixelTrackFitting/plugins/BuildFile.xml
@@ -1,10 +1,7 @@
-
-
-
diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc
deleted file mode 100644
index 6bff9a7c42292..0000000000000
--- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc
+++ /dev/null
@@ -1,104 +0,0 @@
-#include
-#include // needed here by soa layout
-
-#include "CUDADataFormats/Common/interface/Product.h"
-#include "DataFormats/Common/interface/Handle.h"
-#include "FWCore/Framework/interface/ConsumesCollector.h"
-#include "FWCore/Framework/interface/Event.h"
-#include "FWCore/Framework/interface/EventSetup.h"
-#include "FWCore/Framework/interface/MakerMacros.h"
-#include "FWCore/Framework/interface/global/EDAnalyzer.h"
-#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
-#include "FWCore/ParameterSet/interface/ParameterSet.h"
-#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
-#include "FWCore/PluginManager/interface/ModuleDef.h"
-#include "FWCore/Utilities/interface/EDGetToken.h"
-#include "FWCore/Utilities/interface/InputTag.h"
-#include "FWCore/Utilities/interface/RunningAverage.h"
-#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
-#include "RecoTracker/TkMSParametrization/interface/PixelRecoUtilities.h"
-
-#include "CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousHost.h"
-#include "CUDADataFormats/Vertex/interface/ZVertexSoAHeterogeneousDevice.h"
-
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h"
-#include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
-
-template
-class PixelTrackDumpCUDAT : public edm::global::EDAnalyzer<> {
-public:
- using TrackSoAHost = TrackSoAHeterogeneousHost;
- using TrackSoADevice = TrackSoAHeterogeneousDevice;
-
- using VertexSoAHost = ZVertexSoAHost;
- using VertexSoADevice = ZVertexSoADevice;
-
- explicit PixelTrackDumpCUDAT(const edm::ParameterSet& iConfig);
- ~PixelTrackDumpCUDAT() override = default;
-
- static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
-
-private:
- void analyze(edm::StreamID streamID, edm::Event const& iEvent, const edm::EventSetup& iSetup) const override;
- const bool m_onGPU;
- edm::EDGetTokenT> tokenGPUTrack_;
- edm::EDGetTokenT> tokenGPUVertex_;
- edm::EDGetTokenT tokenSoATrack_;
- edm::EDGetTokenT tokenSoAVertex_;
-};
-
-template
-PixelTrackDumpCUDAT::PixelTrackDumpCUDAT(const edm::ParameterSet& iConfig)
- : m_onGPU(iConfig.getParameter("onGPU")) {
- if (m_onGPU) {
- tokenGPUTrack_ = consumes(iConfig.getParameter("pixelTrackSrc"));
- tokenGPUVertex_ = consumes(iConfig.getParameter