Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Code cleanliness improvements in follow-up to #31722 #33865

Closed
wants to merge 10 commits into from
Closed

Code cleanliness improvements in follow-up to #31722 #33865

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
07c25b6
Fixed include guards naming. (#614)
ericcano Apr 6, 2021
7467e17
Unified member access methods for class TrackSoAHeterogeneousT. (#614)
ericcano Apr 6, 2021
7222e8a
Fixed indlude guards naming. (#614)
ericcano Apr 6, 2021
c52e732
Fixed C-style cast. (#614)
ericcano Apr 6, 2021
01043c8
Updated comments after methods renaming. (#614)
ericcano Apr 8, 2021
e783810
Modified PixelTrackProducer class to use esConsume(). (#614)
ericcano Apr 8, 2021
323eca7
Returned complex strcuture by reference instead of value.
ericcano Apr 12, 2021
4ed484d
Fused class, header and SealModule.cc file.
ericcano Apr 13, 2021
0eb287c
Used a wider NaN test to also cover possible Inf with -ofast
ericcano Apr 15, 2021
e1541ce
Fixed code format
ericcano May 27, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
75 changes: 48 additions & 27 deletions CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#ifndef CUDADataFormats_Track_TrackHeterogeneousT_H
#define CUDADataFormats_Track_TrackHeterogeneousT_H
#ifndef CUDADataFormats_Track_PixelTrackHeterogeneousT_h
#define CUDADataFormats_Track_PixelTrackHeterogeneousT_h
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

why PIXEL?
is not specific to pixel

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Indeed, this is a mix-up.

Comment on lines +1 to +2
Copy link
Contributor

@fwyzard fwyzard Jun 18, 2021 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
#ifndef CUDADataFormats_Track_PixelTrackHeterogeneousT_h
#define CUDADataFormats_Track_PixelTrackHeterogeneousT_h
#ifndef CUDADataFormats_Track_interface_TrackSoAHeterogeneousT_h
#define CUDADataFormats_Track_interface_TrackSoAHeterogeneousT_h


#include "CUDADataFormats/Track/interface/TrajectoryStateSoAT.h"
#include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h"
Expand All @@ -19,39 +19,60 @@ class TrackSoAHeterogeneousT {
using hindex_type = uint32_t;
using HitContainer = cms::cuda::OneToManyAssoc<hindex_type, S + 1, 5 * S>;

// quality accessors
constexpr Quality quality(int32_t i) const { return reinterpret_cast<Quality>(quality_(i)); }
constexpr Quality &quality(int32_t i) { return reinterpret_cast<Quality &>(quality_(i)); }
constexpr Quality const *qualityData() const { return reinterpret_cast<Quality const *>(quality_.data()); }
constexpr Quality *qualityData() { return reinterpret_cast<Quality *>(quality_.data()); }

// chi2 accessors
constexpr auto &chi2(int32_t i) { return chi2_(i); }
constexpr auto chi2(int32_t i) const { return chi2_(i); }

// stateAtBS accessors
constexpr TrajectoryStateSoAT<S> &stateAtBS() { return stateAtBS_; }
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

already discussed. I strongly disagree

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This was the consequence of applying the coding rule 4.25 which requires only one of each public, protected and private section. (It can be waived under special circumstances, though).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is not a question of "special circumstances". Coding rules were written at a time of naive OOD. This is DDD. This is a struct composed of SoAs. A new type of beast. Old rules do not apply.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This was the consequence of applying the coding rule 4.25 which requires only one of each public, protected and private section. (It can be waived under special circumstances, though).

4.25 is Each class may have only one each of public, protected, and private sections, which must be declared in that order.

was it something from section 7 instead?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, also. Rule 7.4 stipulates that structs should be real PODs with no getters and setters, and that applied here.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Never understood why a POD looses its "podness" with trivial getters or setters:
what's the difference btw constepxr float A::x() const; and constexpr float x(A const&); ?
Different are constructors or destructor that definitively defy Podness.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

To me 7.4 (and in particular the linked C.131 from the C++ Core Guidelines) does not instruct against a struct with public data members and non-setter/getter member functions.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Never understood why a POD looses its "podness" with trivial getters or setters:

With C++ standard definitionsof POD (or now TrivialType) "podness" is not lost with regular member functions.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ok. So no reason CMS shall be more restrictive than Standard.

constexpr TrajectoryStateSoAT<S> const &stateAtBS() const { return stateAtBS_; }
// eta accessors
constexpr auto &eta(int32_t i) { return eta_(i); }
constexpr auto eta(int32_t i) const { return eta_(i); }
// pt accessors
constexpr auto &pt(int32_t i) { return pt_(i); }
constexpr auto pt(int32_t i) const { return pt_(i); }

constexpr float charge(int32_t i) const { return std::copysign(1.f, stateAtBS_.state(i)(2)); }
constexpr float phi(int32_t i) const { return stateAtBS_.state(i)(0); }
constexpr float tip(int32_t i) const { return stateAtBS_.state(i)(1); }
constexpr float zip(int32_t i) const { return stateAtBS_.state(i)(4); }

// hitIndices accessors
constexpr auto &hitIndices() { return hitIndices_; }
constexpr auto const &hitIndices() const { return hitIndices_; }

// detInndices accessor
constexpr int nHits(int i) const { return detIndices_.size(i); }
constexpr auto &detIndices() { return detIndices_; }
constexpr auto const &detIndices() const { return detIndices_; }

// state at the detector of the outermost hit
// representation to be decided...
// not yet filled on GPU
// TrajectoryStateSoA<S> stateAtOuterDet;
private:
// Always check quality is at least loose!
// CUDA does not support enums in __lgc ...
private:
eigenSoA::ScalarSoA<uint8_t, S> quality_;

public:
constexpr Quality quality(int32_t i) const { return (Quality)(quality_(i)); }
constexpr Quality &quality(int32_t i) { return (Quality &)(quality_(i)); }
constexpr Quality const *qualityData() const { return (Quality const *)(quality_.data()); }
constexpr Quality *qualityData() { return (Quality *)(quality_.data()); }

// this is chi2/ndof as not necessarely all hits are used in the fit
eigenSoA::ScalarSoA<float, S> chi2;

constexpr int nHits(int i) const { return detIndices.size(i); }
eigenSoA::ScalarSoA<float, S> chi2_;

// State at the Beam spot
// phi,tip,1/pt,cotan(theta),zip
TrajectoryStateSoAT<S> stateAtBS;
eigenSoA::ScalarSoA<float, S> eta;
eigenSoA::ScalarSoA<float, S> pt;
constexpr float charge(int32_t i) const { return std::copysign(1.f, stateAtBS.state(i)(2)); }
constexpr float phi(int32_t i) const { return stateAtBS.state(i)(0); }
constexpr float tip(int32_t i) const { return stateAtBS.state(i)(1); }
constexpr float zip(int32_t i) const { return stateAtBS.state(i)(4); }

// state at the detector of the outermost hit
// representation to be decided...
// not yet filled on GPU
// TrajectoryStateSoA<S> stateAtOuterDet;
TrajectoryStateSoAT<S> stateAtBS_;
eigenSoA::ScalarSoA<float, S> eta_;
eigenSoA::ScalarSoA<float, S> pt_;

HitContainer hitIndices;
HitContainer detIndices;
HitContainer hitIndices_;
HitContainer detIndices_;
};

namespace pixelTrack {
Expand All @@ -70,4 +91,4 @@ namespace pixelTrack {

} // namespace pixelTrack

#endif // CUDADataFormats_Track_TrackHeterogeneousT_H
#endif // CUDADataFormats_Track_PixelTrackHeterogeneousT_h
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
#endif // CUDADataFormats_Track_PixelTrackHeterogeneousT_h
#endif // CUDADataFormats_Track_interface_TrackSoAHeterogeneousT_h

12 changes: 6 additions & 6 deletions CUDADataFormats/Track/interface/TrajectoryStateSoAT.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#ifndef CUDADataFormats_Track_TrajectoryStateSOAT_H
#define CUDADataFormats_Track_TrajectoryStateSOAT_H
#ifndef CUDADataFormats_Track_TrajectoryStateSoAT_h
#define CUDADataFormats_Track_TrajectoryStateSoAT_h

#include <Eigen/Dense>
#include "HeterogeneousCore/CUDAUtilities/interface/eigenSoA.h"
Expand All @@ -25,12 +25,12 @@ struct TrajectoryStateSoAT {
auto cov = covariance(i);
cov(0) = ccov(0, 0);
cov(1) = ccov(0, 1);
cov(2) = b * float(ccov(0, 2));
cov(2) = b * static_cast<float>(ccov(0, 2));
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

sorry, do not see the reason to add static_cast

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is purely switching to C++ syntax for the cast.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

just adding useless characters. I claim that for intrinsic type "C-type cast" is more readable and more appropriate.
There is NO other type of cast that can be applied here. The semantic MUST be the of "C-type cast". Not of any C++ cast.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there any use case where it should translate to a reinterpret_cast? There is no dynamic cast for sure with intrisic types and no const correctness issue with an rvalue.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, reinterpret_cast between double and float is invalid.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

just adding useless characters. I claim that for intrinsic type "C-type cast" is more readable and more appropriate.
There is NO other type of cast that can be applied here. The semantic MUST be the of "C-type cast". Not of any C++ cast.

I propose to get an exception in the rule clarified after some discussion in the core group.
@cms-sw/core-l2

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The rule 4.15 already allows exceptions, and admittedly explicit narrowing conversions within floating point types falls on the low-risk side of C-style casts.

The C++ Core Guidelines advocates for gsl::narrow_cast (from the guidelines support library). While that describes the intent well, it wouldn't help towards "adding useless characters".

cov(4) = cov(3) = 0;
cov(5) = ccov(1, 1);
cov(6) = b * float(ccov(1, 2));
cov(6) = b * static_cast<float>(ccov(1, 2));
cov(8) = cov(7) = 0;
cov(9) = b * b * float(ccov(2, 2));
cov(9) = b * b * static_cast<float>(ccov(2, 2));
cov(11) = cov(10) = 0;
cov(12) = lcov(0, 0);
cov(13) = lcov(0, 1);
Expand All @@ -56,4 +56,4 @@ struct TrajectoryStateSoAT {
}
};

#endif // CUDADataFormats_Track_TrajectoryStateSOAT_H
#endif // CUDADataFormats_Track_TrajectoryStateSoAT_h
10 changes: 5 additions & 5 deletions RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -535,9 +535,9 @@ namespace brokenline {

/*!
\brief Helix fit by three step:
-fast pre-fit (see Fast_fit() for further info); \n
-circle fit of the hits projected in the transverse plane by Broken Line algorithm (see BL_Circle_fit() for further info); \n
-line fit of the hits projected on the (pre-fitted) cilinder surface by Broken Line algorithm (see BL_Line_fit() for further info); \n
-fast pre-fit (see fastFit() for further info); \n
-circle fit of the hits projected in the transverse plane by Broken Line algorithm (see circleFit() for further info); \n
-line fit of the hits projected on the (pre-fitted) cilinder surface by Broken Line algorithm (see lineFit() for further info); \n
Points must be passed ordered (from inner to outer layer).

\param hits Matrix3xNd hits coordinates in this form: \n
Expand All @@ -561,9 +561,9 @@ namespace brokenline {
|(x1,z4)|(x2,z4)|(x3,z4)|(x4,z4)|.|(y1,z4)|(y2,z4)|(y3,z4)|(y4,z4)|.|(z1,z4)|(z2,z4)|(z3,z4)|(z4,z4)|
\param bField magnetic field in the center of the detector in Gev/cm/c, in order to perform the p_t calculation.

\warning see BL_Circle_fit(), BL_Line_fit() and Fast_fit() warnings.
\warning see circleFit(), lineFit() and fastFit() warnings.

\bug see BL_Circle_fit(), BL_Line_fit() and Fast_fit() bugs.
\bug see circleFit(), lineFit() and fastFit() bugs.

\return (phi,Tip,p_t,cot(theta)),Zip), their covariance matrix and the chi2's of the circle and line fits.
*/
Expand Down
68 changes: 45 additions & 23 deletions RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,38 +15,60 @@
#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
#include "Geometry/Records/interface/TrackerTopologyRcd.h"

#include "PixelTrackProducer.h"
#include "FWCore/Framework/interface/stream/EDProducer.h"
#include "RecoPixelVertexing/PixelTrackFitting/interface/PixelTrackReconstruction.h"
#include "Geometry/Records/interface/TrackerTopologyRcd.h"
#include "storeTracks.h"

#include "FWCore/PluginManager/interface/ModuleDef.h"
#include "FWCore/Framework/interface/MakerMacros.h"

namespace edm {
class Event;
class EventSetup;
class ParameterSet;
class ConfigurationDescriptions;
} // namespace edm
class TrackerTopology;

using namespace pixeltrackfitting;
using edm::ParameterSet;

PixelTrackProducer::PixelTrackProducer(const ParameterSet& cfg) : theReconstruction(cfg, consumesCollector()) {
edm::LogInfo("PixelTrackProducer") << " construction...";
produces<reco::TrackCollection>();
produces<TrackingRecHitCollection>();
produces<reco::TrackExtraCollection>();
}
class PixelTrackProducer : public edm::stream::EDProducer<> {
public:
explicit PixelTrackProducer(const edm::ParameterSet& cfg)
: theReconstruction(cfg, consumesCollector()), htTopoToken_(esConsumes()) {
edm::LogInfo("PixelTrackProducer") << " construction...";
produces<reco::TrackCollection>();
produces<TrackingRecHitCollection>();
produces<reco::TrackExtraCollection>();
}

~PixelTrackProducer() override {}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
~PixelTrackProducer() override {}
~PixelTrackProducer() override = default;


static void fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
edm::ParameterSetDescription desc;

PixelTrackProducer::~PixelTrackProducer() {}
desc.add<std::string>("passLabel", "pixelTracks"); // What is this? It is not used anywhere in this code.
PixelTrackReconstruction::fillDescriptions(desc);

void PixelTrackProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
edm::ParameterSetDescription desc;
descriptions.add("pixelTracks", desc);
}

desc.add<std::string>("passLabel", "pixelTracks"); // What is this? It is not used anywhere in this code.
PixelTrackReconstruction::fillDescriptions(desc);
void produce(edm::Event& ev, const edm::EventSetup& es) override {
LogDebug("PixelTrackProducer, produce") << "event# :" << ev.id();

descriptions.add("pixelTracks", desc);
}
TracksWithTTRHs tracks;
theReconstruction.run(tracks, ev, es);
auto htTopo = es.getData(htTopoToken_);

void PixelTrackProducer::produce(edm::Event& ev, const edm::EventSetup& es) {
LogDebug("PixelTrackProducer, produce") << "event# :" << ev.id();
// store tracks
storeTracks(ev, tracks, htTopo);
}

TracksWithTTRHs tracks;
theReconstruction.run(tracks, ev, es);
edm::ESHandle<TrackerTopology> httopo;
es.get<TrackerTopologyRcd>().get(httopo);
private:
PixelTrackReconstruction theReconstruction;
const edm::ESGetToken<TrackerTopology, TrackerTopologyRcd> htTopoToken_;
};

// store tracks
storeTracks(ev, tracks, *httopo);
}
DEFINE_FWK_MODULE(PixelTrackProducer);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I already said I was takling care of this in my PR that is coming soon

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I can keep this PR on hold and apply on top of yours when it is in.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think that this PR should be closed. None of the modification in here is neither essential nor an improvement.

29 changes: 0 additions & 29 deletions RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h
View file
Open in desktop

This file was deleted.

4 changes: 2 additions & 2 deletions RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,8 +133,8 @@ void PixelTrackProducerFromSoA::produce(edm::StreamID streamID,
const auto &tsoa = *iEvent.get(tokenTrack_);

auto const *quality = tsoa.qualityData();
auto const &fit = tsoa.stateAtBS;
auto const &hitIndices = tsoa.hitIndices;
auto const &fit = tsoa.stateAtBS();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

as noted before, (and discussed already also in CORE, there is no agreement that this is a better syntax for SOA.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This was also a consequence of the 4.25 rule which forces to choose between direct member access in struct style or using accessors, and this class was a in-between.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There are very good reason why some members are directly accessed and other through a get functon. For instance it let user to think was is going on! It avoids obfuscations.

Copy link
Contributor

@fwyzard fwyzard Jun 8, 2021 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the question meant: is there any chance we want this to ever be a reinterpret_cast ?
If not, then using a static_cast is more accurate than a C cast (that can apply all of static_cast, reinterpret_cast and const_cast at once).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

in that specific case we want the compiler to emit a conversion (by value!). We want a temporary float created there!
equivalent to float tmp = f; no more, no less. nothing fancy. As in good old C code (or assembler).
I prefer to introduce such an intermediate statement than use a C++ cast that has not place in math expressions.

Copy link
Contributor

@VinInn VinInn Jun 8, 2021 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The only reinterpret_cast accepted is reinterpret_cast<long long int&>(x) and this is funny as it is equivalent to type punning that is undef behavior. in C++20 they have introduced (at last, after having made undef behavior all C-style idioms!) bit_cast that is equivalent to a memcpy which is the correct semantic. In short nobody uses reinterpret_cast among intrinsic type.

auto const &hitIndices = tsoa.hitIndices();
auto maxTracks = tsoa.stride();

int32_t nt = 0;
Expand Down
5 changes: 0 additions & 5 deletions RecoPixelVertexing/PixelTrackFitting/plugins/SealModule.cc
View file
Open in desktop

This file was deleted.

2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -155,7 +155,7 @@ __global__ void kernel_BLFit(caConstants::TupleMultiplicity const *__restrict__
brokenline::lineFit(hits_ge, fast_fit, bField, data, line);
brokenline::circleFit(hits, hits_ge, fast_fit, bField, data, circle);

results->stateAtBS.copyFromCircle(circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid);
results->stateAtBS().copyFromCircle(circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid);
results->pt(tkid) = float(bField) / float(std::abs(circle.par(2)));
results->eta(tkid) = asinhf(line.par(0));
results->chi2(tkid) = (circle.chi2 + line.chi2) / (2 * N - 5);
Expand Down
6 changes: 3 additions & 3 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ void CAHitNtupletGeneratorKernelsCPU::printCounters(Counters const *counters) {

template <>
void CAHitNtupletGeneratorKernelsCPU::fillHitDetIndices(HitsView const *hv, TkSoA *tracks_d, cudaStream_t) {
kernel_fillHitDetIndices(&tracks_d->hitIndices, hv, &tracks_d->detIndices);
kernel_fillHitDetIndices(&tracks_d->hitIndices(), hv, &tracks_d->detIndices());
}

template <>
Expand Down Expand Up @@ -75,7 +75,7 @@ void CAHitNtupletGeneratorKernelsCPU::buildDoublets(HitsOnCPU const &hh, cudaStr

template <>
void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) {
auto *tuples_d = &tracks_d->hitIndices;
auto *tuples_d = &tracks_d->hitIndices();
auto *quality_d = tracks_d->qualityData();

assert(tuples_d && quality_d);
Expand Down Expand Up @@ -154,7 +154,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA *

template <>
void CAHitNtupletGeneratorKernelsCPU::classifyTuples(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) {
auto const *tuples_d = &tracks_d->hitIndices;
auto const *tuples_d = &tracks_d->hitIndices();
auto *quality_d = tracks_d->qualityData();

// classify tracks based on kinematics
Expand Down
6 changes: 3 additions & 3 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ void CAHitNtupletGeneratorKernelsGPU::fillHitDetIndices(HitsView const *hv, TkSo
auto numberOfBlocks = (HitContainer::ctCapacity() + blockSize - 1) / blockSize;

kernel_fillHitDetIndices<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
&tracks_d->hitIndices, hv, &tracks_d->detIndices);
&tracks_d->hitIndices(), hv, &tracks_d->detIndices());
cudaCheck(cudaGetLastError());
#ifdef GPU_DEBUG
cudaDeviceSynchronize();
Expand All @@ -17,7 +17,7 @@ void CAHitNtupletGeneratorKernelsGPU::fillHitDetIndices(HitsView const *hv, TkSo
template <>
void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) {
// these are pointer on GPU!
auto *tuples_d = &tracks_d->hitIndices;
auto *tuples_d = &tracks_d->hitIndices();
auto *quality_d = tracks_d->qualityData();

// zero tuples
Expand Down Expand Up @@ -224,7 +224,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(HitsOnCPU const &hh, cudaStr
template <>
void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) {
// these are pointer on GPU!
auto const *tuples_d = &tracks_d->hitIndices;
auto const *tuples_d = &tracks_d->hitIndices();
auto *quality_d = tracks_d->qualityData();

int32_t nhits = hh.nHits();
Expand Down
7 changes: 4 additions & 3 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
#include "FWCore/Utilities/interface/isFinite.h"

#include "CAConstants.h"
#include "CAHitNtupletGeneratorKernels.h"
Expand Down Expand Up @@ -359,11 +360,11 @@ __global__ void kernel_classifyTracks(HitContainer const *__restrict__ tuples,
continue;

// if the fit has any invalid parameters, mark it as bad
bool isNaN = false;
bool isNotFinite = false;
for (int i = 0; i < 5; ++i) {
isNaN |= std::isnan(tracks->stateAtBS.state(it)(i));
isNotFinite |= edm::isNotFinite(tracks->stateAtBS().state(it)(i));
}
if (isNaN) {
if (isNotFinite) {
#ifdef NTUPLE_DEBUG
printf("NaN in fit %d size %d chi2 %f\n", it, tuples->size(it), tracks->chi2(it));
#endif
Expand Down
4 changes: 2 additions & 2 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -192,7 +192,7 @@ PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuplesAsync(TrackingRecH
kernels.fillHitDetIndices(hits_d.view(), soa, stream); // in principle needed only if Hits not "available"

HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
fitter.allocateOnGPU(&(soa->hitIndices()), kernels.tupleMultiplicity(), soa);
if (m_params.useRiemannFit_) {
fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream);
} else {
Expand Down Expand Up @@ -228,7 +228,7 @@ PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples(TrackingRecHit2DC

// now fit
HelixFitOnGPU fitter(bfield, m_params.fit5as4_);
fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa);
fitter.allocateOnGPU(&(soa->hitIndices()), kernels.tupleMultiplicity(), soa);

if (m_params.useRiemannFit_) {
fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets);
Expand Down
2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -159,7 +159,7 @@ __global__ void kernel_LineFit(caConstants::TupleMultiplicity const *__restrict_

riemannFit::fromCircleToPerigee(circle_fit[local_idx]);

results->stateAtBS.copyFromCircle(
results->stateAtBS().copyFromCircle(
circle_fit[local_idx].par, circle_fit[local_idx].cov, line_fit.par, line_fit.cov, 1.f / float(bField), tkid);
results->pt(tkid) = bField / std::abs(circle_fit[local_idx].par(2));
results->eta(tkid) = asinhf(line_fit.par(0));
Expand Down
Loading