Skip to content

Commit

Permalink
Synchronise GPU code with CPU updates (cms-sw#576)
Browse files Browse the repository at this point in the history 
Update GPU code following cms-sw#32146.
  • Loading branch information
@mariadalfonso
mariadalfonso committed Nov 22, 2020
1 parent a0dfe83 commit 3ef20e5
Show file tree
Hide file tree
Showing 6 changed files with 60 additions and 50 deletions.
18 changes: 9 additions & 9 deletions RecoLocalCalo/HcalRecAlgos/interface/PulseShapeFunctor.h
View file
Expand Up @@ -59,20 +59,20 @@ namespace FitterFuncs {
}

// getters
inline std::vector<float> const &get_acc25nsVec() const { return acc25nsVec; }
inline std::vector<float> const &get_diff25nsItvlVec() const { return diff25nsItvlVec; }
inline std::vector<float> const &get_accVarLenIdxZEROVec() const { return accVarLenIdxZEROVec; }
inline std::vector<float> const &get_diffVarItvlIdxZEROVec() const { return diffVarItvlIdxZEROVec; }
inline std::vector<float> const &get_accVarLenIdxMinusOneVec() const { return accVarLenIdxMinusOneVec; }
inline std::vector<float> const &get_diffVarItvlIdxMinusOneVec() const { return diffVarItvlIdxMinusOneVec; }
inline std::vector<float> const &acc25nsVec() const { return acc25nsVec_; }
inline std::vector<float> const &diff25nsItvlVec() const { return diff25nsItvlVec_; }
inline std::vector<float> const &accVarLenIdxZEROVec() const { return accVarLenIdxZEROVec_; }
inline std::vector<float> const &diffVarItvlIdxZEROVec() const { return diffVarItvlIdxZEROVec_; }
inline std::vector<float> const &accVarLenIdxMinusOneVec() const { return accVarLenIdxMinusOneVec_; }
inline std::vector<float> const &diffVarItvlIdxMinusOneVec() const { return diffVarItvlIdxMinusOneVec_; }

private:
std::array<float, hcal::constants::maxPSshapeBin> pulse_hist;

int cntNANinfit;
std::vector<float> acc25nsVec, diff25nsItvlVec;
std::vector<float> accVarLenIdxZEROVec, diffVarItvlIdxZEROVec;
std::vector<float> accVarLenIdxMinusOneVec, diffVarItvlIdxMinusOneVec;
std::vector<float> acc25nsVec_, diff25nsItvlVec_;
std::vector<float> accVarLenIdxZEROVec_, diffVarItvlIdxZEROVec_;
std::vector<float> accVarLenIdxMinusOneVec_, diffVarItvlIdxMinusOneVec_;

void funcShape(std::array<double, hcal::constants::maxSamples> &ntmpbin,
const double pulseTime,
Expand Down
28 changes: 14 additions & 14 deletions RecoLocalCalo/HcalRecAlgos/src/HcalRecoParamsWithPulseShapesGPU.cc
View file
Expand Up @@ -55,20 +55,20 @@ HcalRecoParamsWithPulseShapesGPU::HcalRecoParamsWithPulseShapesGPU(HcalRecoParam

// precompute and get values from the functor
auto const& pulseShape = pulseShapes.getShape(pulseShapeId);
FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, 10};
FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, hcal::constants::maxSamples};
auto const offset256 = newId * hcal::constants::maxPSshapeBin;
auto const offset25 = newId * hcal::constants::nsPerBX;
auto const numShapes = newId;
for (int i = 0; i < hcal::constants::maxPSshapeBin; i++) {
acc25nsVec_[offset256 * numShapes + i] = functor.get_acc25nsVec()[i];
diff25nsItvlVec_[offset256 * numShapes + i] = functor.get_diff25nsItvlVec()[i];
acc25nsVec_[offset256 * numShapes + i] = functor.acc25nsVec()[i];
diff25nsItvlVec_[offset256 * numShapes + i] = functor.diff25nsItvlVec()[i];
}

for (int i = 0; i < hcal::constants::nsPerBX; i++) {
accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.get_accVarLenIdxMinusOneVec()[i];
diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.get_diffVarItvlIdxMinusOneVec()[i];
accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.get_accVarLenIdxZEROVec()[i];
diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.get_diffVarItvlIdxZEROVec()[i];
accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.accVarLenIdxMinusOneVec()[i];
diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxMinusOneVec()[i];
accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.accVarLenIdxZEROVec()[i];
diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxZEROVec()[i];
}
} else {
// already recorded this pulse shape, just set id
Expand Down Expand Up @@ -113,20 +113,20 @@ HcalRecoParamsWithPulseShapesGPU::HcalRecoParamsWithPulseShapesGPU(HcalRecoParam

// precompute and get values from the functor
auto const& pulseShape = pulseShapes.getShape(pulseShapeId);
FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, 10};
FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, hcal::constants::maxSamples};
auto const offset256 = newId * hcal::constants::maxPSshapeBin;
auto const offset25 = newId * hcal::constants::nsPerBX;
auto const numShapes = newId;
for (int i = 0; i < hcal::constants::maxPSshapeBin; i++) {
acc25nsVec_[offset256 * numShapes + i] = functor.get_acc25nsVec()[i];
diff25nsItvlVec_[offset256 * numShapes + i] = functor.get_diff25nsItvlVec()[i];
acc25nsVec_[offset256 * numShapes + i] = functor.acc25nsVec()[i];
diff25nsItvlVec_[offset256 * numShapes + i] = functor.diff25nsItvlVec()[i];
}

for (int i = 0; i < hcal::constants::nsPerBX; i++) {
accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.get_accVarLenIdxMinusOneVec()[i];
diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.get_diffVarItvlIdxMinusOneVec()[i];
accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.get_accVarLenIdxZEROVec()[i];
diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.get_diffVarItvlIdxZEROVec()[i];
accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.accVarLenIdxMinusOneVec()[i];
diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxMinusOneVec()[i];
accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.accVarLenIdxZEROVec()[i];
diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxZEROVec()[i];
}
} else {
// already recorded this pulse shape, just set id
Expand Down
6 changes: 5 additions & 1 deletion RecoLocalCalo/HcalRecAlgos/src/MahiFit.cc
View file
Expand Up @@ -84,7 +84,7 @@ void MahiFit::phase1Apply(const HBHEChannelInfo& channelData,
tstrig *= channelData.tsGain(0);

useTriple = false;
if (tstrig >= ts4Thresh_ && tsTOT > 0) {
if (tstrig > ts4Thresh_ && tsTOT > 0) {
//Average pedestal width (for covariance matrix constraint)
nnlsWork_.pedVal = 0.25f * (channelData.tsPedestalWidth(0) * channelData.tsPedestalWidth(0) +
channelData.tsPedestalWidth(1) * channelData.tsPedestalWidth(1) +
Expand Down Expand Up @@ -255,6 +255,10 @@ void MahiFit::updatePulseShape(const float itQ,
FullSampleVector& pulseShape,
FullSampleVector& pulseDeriv,
FullSampleMatrix& pulseCov) const {
// set a null pulse shape for negative / or null TS
if (itQ <= 0.f)
return;

float t0 = meanTime_;

if (applyTimeSlew_) {
Expand Down
38 changes: 19 additions & 19 deletions RecoLocalCalo/HcalRecAlgos/src/PulseShapeFunctor.cc
View file
Expand Up @@ -16,36 +16,36 @@ namespace FitterFuncs {
double iPedMean,
unsigned nSamplesToFit)
: cntNANinfit(0),
acc25nsVec(hcal::constants::maxPSshapeBin),
diff25nsItvlVec(hcal::constants::maxPSshapeBin),
accVarLenIdxZEROVec(hcal::constants::nsPerBX),
diffVarItvlIdxZEROVec(hcal::constants::nsPerBX),
accVarLenIdxMinusOneVec(hcal::constants::nsPerBX),
diffVarItvlIdxMinusOneVec(hcal::constants::nsPerBX) {
acc25nsVec_(hcal::constants::maxPSshapeBin),
diff25nsItvlVec_(hcal::constants::maxPSshapeBin),
accVarLenIdxZEROVec_(hcal::constants::nsPerBX),
diffVarItvlIdxZEROVec_(hcal::constants::nsPerBX),
accVarLenIdxMinusOneVec_(hcal::constants::nsPerBX),
diffVarItvlIdxMinusOneVec_(hcal::constants::nsPerBX) {
//The raw pulse
for (int i = 0; i < hcal::constants::maxPSshapeBin; ++i)
pulse_hist[i] = pulse(i);
// Accumulate 25ns for each starting point of 0, 1, 2, 3...
for (int i = 0; i < hcal::constants::maxPSshapeBin; ++i) {
for (int j = i; j < i + hcal::constants::nsPerBX; ++j) { //sum over hcal::constants::nsPerBXns from point i
acc25nsVec[i] +=
acc25nsVec_[i] +=
(j < hcal::constants::maxPSshapeBin ? pulse_hist[j] : pulse_hist[hcal::constants::maxPSshapeBin - 1]);
}
diff25nsItvlVec[i] = (i + hcal::constants::nsPerBX < hcal::constants::maxPSshapeBin
? pulse_hist[i + hcal::constants::nsPerBX] - pulse_hist[i]
: pulse_hist[hcal::constants::maxPSshapeBin - 1] - pulse_hist[i]);
diff25nsItvlVec_[i] = (i + hcal::constants::nsPerBX < hcal::constants::maxPSshapeBin
? pulse_hist[i + hcal::constants::nsPerBX] - pulse_hist[i]
: pulse_hist[hcal::constants::maxPSshapeBin - 1] - pulse_hist[i]);
}
// Accumulate different ns for starting point of index either 0 or -1
for (int i = 0; i < hcal::constants::nsPerBX; ++i) {
if (i == 0) {
accVarLenIdxZEROVec[0] = pulse_hist[0];
accVarLenIdxMinusOneVec[i] = pulse_hist[0];
accVarLenIdxZEROVec_[0] = pulse_hist[0];
accVarLenIdxMinusOneVec_[i] = pulse_hist[0];
} else {
accVarLenIdxZEROVec[i] = accVarLenIdxZEROVec[i - 1] + pulse_hist[i];
accVarLenIdxMinusOneVec[i] = accVarLenIdxMinusOneVec[i - 1] + pulse_hist[i - 1];
accVarLenIdxZEROVec_[i] = accVarLenIdxZEROVec_[i - 1] + pulse_hist[i];
accVarLenIdxMinusOneVec_[i] = accVarLenIdxMinusOneVec_[i - 1] + pulse_hist[i - 1];
}
diffVarItvlIdxZEROVec[i] = pulse_hist[i + 1] - pulse_hist[0];
diffVarItvlIdxMinusOneVec[i] = pulse_hist[i] - pulse_hist[0];
diffVarItvlIdxZEROVec_[i] = pulse_hist[i + 1] - pulse_hist[0];
diffVarItvlIdxMinusOneVec_[i] = pulse_hist[i] - pulse_hist[0];
}
for (int i = 0; i < hcal::constants::maxSamples; i++) {
psFit_x[i] = 0;
Expand Down Expand Up @@ -103,12 +103,12 @@ namespace FitterFuncs {
ntmpbin[iTS_start] =
(bin_0_start == -1
? // Initial bin (I'm assuming this is ok)
accVarLenIdxMinusOneVec[distTo25ns_start] + factor * diffVarItvlIdxMinusOneVec[distTo25ns_start]
: accVarLenIdxZEROVec[distTo25ns_start] + factor * diffVarItvlIdxZEROVec[distTo25ns_start]);
accVarLenIdxMinusOneVec_[distTo25ns_start] + factor * diffVarItvlIdxMinusOneVec_[distTo25ns_start]
: accVarLenIdxZEROVec_[distTo25ns_start] + factor * diffVarItvlIdxZEROVec_[distTo25ns_start]);
//Fill the rest of the bins
for (int iTS = iTS_start + 1; iTS < hcal::constants::maxSamples; ++iTS) {
int bin_idx = distTo25ns_start + 1 + (iTS - iTS_start - 1) * ns_per_bx + bin_0_start;
ntmpbin[iTS] = acc25nsVec[bin_idx] + factor * diff25nsItvlVec[bin_idx];
ntmpbin[iTS] = acc25nsVec_[bin_idx] + factor * diff25nsItvlVec_[bin_idx];
}
//Scale the pulse
if (scalePulse) {
Expand Down
1 change: 1 addition & 0 deletions RecoLocalCalo/HcalRecProducers/src/KernelHelpers.h
View file
Expand Up @@ -146,6 +146,7 @@ namespace hcal {
}

// TODO: remove what's not needed
// originally from from RecoLocalCalo/HcalRecAlgos/src/PulseShapeFunctor.cc
__forceinline__ __device__ float compute_pulse_shape_value(float const pulse_time,
int const sample,
int const shift,
Expand Down
19 changes: 12 additions & 7 deletions RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
View file
Expand Up @@ -568,7 +568,7 @@ namespace hcal {
}

if (sample == 0 && ipulse == 0) {
for (int i = 0; i < 10; i++) {
for (int i = 0; i < hcal::constants::maxSamples; i++) {
auto const value = hcal::reconstruction::compute_pulse_shape_value(t0,
i,
0,
Expand All @@ -581,7 +581,7 @@ namespace hcal {
printf("pulse(%d) = %f\n", i, value);
}
printf("\n");
for (int i = 0; i < 10; i++) {
for (int i = 0; i < hcal::constants::maxSamples; i++) {
auto const value = hcal::reconstruction::compute_pulse_shape_value(t0p,
i,
0,
Expand All @@ -594,7 +594,7 @@ namespace hcal {
printf("pulseP(%d) = %f\n", i, value);
}
printf("\n");
for (int i = 0; i < 10; i++) {
for (int i = 0; i < hcal::constants::maxSamples; i++) {
auto const value = hcal::reconstruction::compute_pulse_shape_value(t0m,
i,
0,
Expand Down Expand Up @@ -651,10 +651,15 @@ namespace hcal {
: 0;

// store to global
pulseMatrix[ipulse * nsamples + sample] = value;
;
pulseMatrixM[ipulse * nsamples + sample] = value_t0m;
pulseMatrixP[ipulse * nsamples + sample] = value_t0p;
if (amplitude > 0.f) {
pulseMatrix[ipulse * nsamples + sample] = value;
pulseMatrixM[ipulse * nsamples + sample] = value_t0m;
pulseMatrixP[ipulse * nsamples + sample] = value_t0p;
} else {
pulseMatrix[ipulse * nsamples + sample] = 0.f;
pulseMatrixM[ipulse * nsamples + sample] = 0.f;
pulseMatrixP[ipulse * nsamples + sample] = 0.f;
}
}

template <int NSAMPLES, int NPULSES>
Expand Down

0 comments on commit 3ef20e5

Please sign in to comment.