Skip to content

Commit

Permalink
CTD19 developments for review and merging (#324)
Browse files Browse the repository at this point in the history
  • Loading branch information
@rovere @fwyzard
rovere authored and fwyzard committed May 2, 2019
1 parent 71ab195 commit 20adc27
Show file tree
Hide file tree
Showing 10 changed files with 51 additions and 191 deletions.
60 changes: 30 additions & 30 deletions Configuration/AlCa/python/autoCond.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,38 +2,38 @@

### NEW KEYS ###
# GlobalTag for MC production with perfectly aligned and calibrated detector for Run1
'run1_design' : '103X_mcRun1_design_v1',
'run1_design' : '105X_mcRun1_design_v1',
# GlobalTag for MC production (pp collisions) with realistic alignment and calibrations for Run1
'run1_mc' : '103X_mcRun1_realistic_v1',
'run1_mc' : '105X_mcRun1_realistic_v1',
# GlobalTag for MC production (Heavy Ions collisions) with realistic alignment and calibrations for Run1
'run1_mc_hi' : '103X_mcRun1_HeavyIon_v1',
'run1_mc_hi' : '105X_mcRun1_HeavyIon_v1',
# GlobalTag for MC production (p-Pb collisions) with realistic alignment and calibrations for Run1
'run1_mc_pa' : '103X_mcRun1_pA_v1',
'run1_mc_pa' : '105X_mcRun1_pA_v1',
# GlobalTag for MC production with pessimistic alignment and calibrations for Run2
'run2_mc_50ns' : '103X_mcRun2_startup_v3',
'run2_mc_50ns' : '105X_mcRun2_startup_v2',
# GlobalTag for MC production (L1 Trigger Stage1) with starup-like alignment and calibrations for Run2, L1 trigger in Stage1 mode
'run2_mc_l1stage1' : '103X_mcRun2_asymptotic_l1stage1_v1',
'run2_mc_l1stage1' : '105X_mcRun2_asymptotic_l1stage1_v2',
# GlobalTag for MC production with perfectly aligned and calibrated detector for Run2
'run2_design' : '105X_mcRun2_design_v1',
'run2_design' : '105X_mcRun2_design_v3',
#GlobalTag for MC production with optimistic alignment and calibrations for Run2
'run2_mc' : '105X_mcRun2_asymptotic_v2',
'run2_mc' : '105X_mcRun2_asymptotic_v4',
# GlobalTag for MC production (cosmics) with starup-like alignment and calibrations for Run2, Strip tracker in peak mode
'run2_mc_cosmics' : '105X_mcRun2cosmics_startup_deco_v2',
'run2_mc_cosmics' : '105X_mcRun2cosmics_startup_deco_v4',
# GlobalTag for MC production (Heavy Ions collisions) with optimistic alignment and calibrations for Run2
'run2_mc_hi' : '103X_mcRun2_HeavyIon_v3',
'run2_mc_hi' : '105X_mcRun2_HeavyIon_v2',
# GlobalTag for MC production (p-Pb collisions) with realistic alignment and calibrations for Run2
'run2_mc_pa' : '103X_mcRun2_pA_v3',
'run2_mc_pa' : '105X_mcRun2_pA_v2',
# GlobalTag for Run1 data reprocessing
'run1_data' : '105X_dataRun2_v6',
'run1_data' : '105X_dataRun2_v8',
# GlobalTag for Run2 data reprocessing
'run2_data' : '105X_dataRun2_v6',
'run2_data' : '105X_dataRun2_v8',
# GlobalTag for Run2 data relvals: allows customization to run with fixed L1 menu
'run2_data_relval' : '105X_dataRun2_relval_v6',
'run2_data_relval' : '105X_dataRun2_relval_v8',
# GlobalTag for Run2 data 2018B relvals only: HEM-15-16 fail
'run2_data_promptlike_HEfail' : '105X_dataRun2_PromptLike_HEfail_v5',
'run2_data_promptlike_HEfail' : '105X_dataRun2_PromptLike_HEfail_v7',
# GlobalTag for Run2 data 2016H relvals only: Prompt Conditions + fixed L1 menu (to be removed)
'run2_data_promptlike' : '105X_dataRun2_PromptLike_v6',
'run2_data_promptlike_hi' : '105X_dataRun2_PromptLike_HI_v1',
'run2_data_promptlike' : '105X_dataRun2_PromptLike_v8',
'run2_data_promptlike_hi' : '105X_dataRun2_PromptLike_HI_v3',
# GlobalTag for Run1 HLT: it points to the online GT
'run1_hlt' : '101X_dataRun2_HLT_frozen_v8',
# GlobalTag for Run2 HLT: it points to the online GT
Expand All @@ -44,31 +44,31 @@
# GlobalTag for Run2 HLT for HI (not 2018 HI): it points to the online GT
'run2_hlt_hi' : '101X_dataRun2_HLTHI_frozen_v8',
# GlobalTag for MC production with perfectly aligned and calibrated detector for Phase1 2017 (and 0,0,~0-centred beamspot)
'phase1_2017_design' : '105X_mc2017_design_IdealBS_v4',
'phase1_2017_design' : '105X_mc2017_design_IdealBS_v6',
# GlobalTag for MC production with realistic conditions for Phase1 2017 detector
'phase1_2017_realistic' : '105X_mc2017_realistic_v5',
'phase1_2017_realistic' : '105X_mc2017_realistic_v7',
# GlobalTag for MC production (cosmics) with realistic alignment and calibrations for Phase1 2017 detector, Strip tracker in DECO mode
'phase1_2017_cosmics' : '105X_mc2017cosmics_realistic_deco_v6',
'phase1_2017_cosmics' : '105X_mc2017cosmics_realistic_deco_v8',
# GlobalTag for MC production (cosmics) with realistic alignment and calibrations for Phase1 2017 detector, Strip tracker in PEAK mode
'phase1_2017_cosmics_peak' : '105X_mc2017cosmics_realistic_peak_v6',
'phase1_2017_cosmics_peak' : '105X_mc2017cosmics_realistic_peak_v8',
# GlobalTag for MC production with perfectly aligned and calibrated detector for full Phase1 2018 (and 0,0,0-centred beamspot)
'phase1_2018_design' : '105X_upgrade2018_design_v3',
'phase1_2018_design' : '105X_upgrade2018_design_v5',
# GlobalTag for MC production with realistic conditions for full Phase1 2018 detector
'phase1_2018_realistic' : '105X_upgrade2018_realistic_v4',
'phase1_2018_realistic' : '105X_upgrade2018_realistic_v6',
# GlobalTag for MC production with realistic conditions for full Phase1 2018 detector for Heavy Ion
'phase1_2018_realistic_hi' : '105X_upgrade2018_realistic_HI_v2',
'phase1_2018_realistic_hi' : '105X_upgrade2018_realistic_HI_v4',
# GlobalTag for MC production with realistic conditions for full Phase1 2018 detector: HEM-15-16 fail
'phase1_2018_realistic_HEfail' : '105X_upgrade2018_realistic_HEfail_v5',
'phase1_2018_realistic_HEfail' : '105X_upgrade2018_realistic_HEfail_v6',
# GlobalTag for MC production (cosmics) with realistic conditions for full Phase1 2018 detector, Strip tracker in DECO mode
'phase1_2018_cosmics' : '105X_upgrade2018cosmics_realistic_deco_v5',
'phase1_2018_cosmics' : '105X_upgrade2018cosmics_realistic_deco_v7',
# GlobalTag for MC production (cosmics) with realistic conditions for full Phase1 2018 detector, Strip tracker in PEAK mode
'phase1_2018_cosmics_peak' : '105X_upgrade2018cosmics_realistic_peak_v5',
'phase1_2018_cosmics_peak' : '105X_upgrade2018cosmics_realistic_peak_v7',
# GlobalTag for MC production with perfectly aligned and calibrated detector for Phase1 2019
'phase1_2019_design' : '105X_postLS2_design_v2', # GT containing design conditions for postLS2
'phase1_2019_design' : '105X_postLS2_design_v4', # GT containing design conditions for postLS2
# GlobalTag for MC production with realistic conditions for Phase1 2019
'phase1_2019_realistic' : '105X_postLS2_realistic_v4', # GT containing realistic conditions for postLS2
'phase1_2019_realistic' : '105X_postLS2_realistic_v6', # GT containing realistic conditions for postLS2
# GlobalTag for MC production with realistic conditions for Phase2 2023
'phase2_realistic' : '105X_upgrade2023_realistic_v3'
'phase2_realistic' : '105X_upgrade2023_realistic_v5'
}

aliases = {
Expand Down
1 change: 0 additions & 1 deletion RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,6 @@ namespace {
auto i = blockIdx.x * blockDim.x + threadIdx.x;

assert(0==hitsModuleStart[0]);

if(i < 11) {
hitsLayerStart[i] = hitsModuleStart[cpeParams->layerGeometry().layerStart[i]];
#ifdef GPU_DEBUG
Expand Down
2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelTrackFitting/interface/PixelNtupletsFitter.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@ class PixelNtupletsFitter final : public PixelFitterBase {
explicit PixelNtupletsFitter(float nominalB, const MagneticField *field, bool useRiemannFit);
~PixelNtupletsFitter() override = default;
std::unique_ptr<reco::Track> run(const std::vector<const TrackingRecHit *>& hits,
const TrackingRegion& region) const override;
const TrackingRegion& region, const edm::EventSetup& es) const override;

private:
float nominalB_;
Expand Down
3 changes: 1 addition & 2 deletions RecoPixelVertexing/PixelTrackFitting/python/PixelTracks_cff.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -94,5 +94,4 @@
_pixelTracksTask_ntupleFit.replace(pixelFitterByHelixProjections, pixelNtupletsFitter)
ntupleFit.toReplaceWith(pixelTracksTask, _pixelTracksTask_ntupleFit)

riemannFit.toReplaceWith(pixelTracksSequence, _pixelTracksSequence_riemannFit)
brokenLine.toReplaceWith(pixelTracksSequence, _pixelTracksSequence_brokenLine)
pixelTracksSequence = cms.Sequence(pixelTracksTask)
2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelTrackFitting/src/PixelNtupletsFitter.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ PixelNtupletsFitter::PixelNtupletsFitter(float nominalB, const MagneticField* fi
useRiemannFit_(useRiemannFit) {}

std::unique_ptr<reco::Track> PixelNtupletsFitter::run(
const std::vector<const TrackingRecHit*>& hits, const TrackingRegion& region) const {
const std::vector<const TrackingRecHit*>& hits, const TrackingRegion& region, const edm::EventSetup& ) const {

using namespace Rfit;

Expand Down
1 change: 1 addition & 0 deletions RecoPixelVertexing/PixelTrackFitting/test/BuildFile.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,7 @@
<use name="eigen"/>
<use name="cuda"/>
<use name="cuda-api-wrappers"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<flags CXXFLAGS="-g"/>
</bin>

Expand Down
161 changes: 10 additions & 151 deletions RecoPixelVertexing/PixelTrackFitting/test/testEigenGPU.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,23 +33,6 @@ namespace Rfit {
// fast fit
using Map4d = Eigen::Map<Vector4d,0,Eigen::InnerStride<stride()> >;

__global__
void kernelFullBrokenLineFastFitAndData(BrokenLine::Matrix3xNd * hits,
BrokenLine::Matrix3Nd * hits_cov,
BrokenLine::PreparedBrokenLineData * data,
Vector4d * fast_fit,
double B,
BrokenLine::helix_fit * helix_fit_resultsGPU,
BrokenLine::karimaki_circle_fit * circleGPU,
BrokenLine::line_fit * lineGPU,
Matrix3d * JacobGPU,
BrokenLine::MatrixNplusONEd * C_UGPU) {

BrokenLine::helix_fit& helix = (*helix_fit_resultsGPU);

helix.fast_fit=BrokenLine::BL_Fast_fit(*hits);

BrokenLine::PrepareBrokenLineData(*hits,*hits_cov,helix.fast_fit,B,*data);
}


Expand Down Expand Up @@ -157,71 +140,15 @@ if (0==i) {
printf("hits_cov(11,11): %f\n", (*hits_cov)(11,11));
printf("B: %f\n", B);
}

__global__
void kernelFullBrokenLineHelix(BrokenLine::Matrix3xNd * hits,
BrokenLine::Matrix3Nd * hits_cov,
BrokenLine::PreparedBrokenLineData * data,
Vector4d * fast_fit,
double B,
BrokenLine::helix_fit * helix_fit_resultsGPU,
BrokenLine::karimaki_circle_fit * circleGPU,
BrokenLine::line_fit * lineGPU,
Matrix3d * JacobGPU,
BrokenLine::MatrixNplusONEd * C_UGPU) {

BrokenLine::helix_fit& helix = (*helix_fit_resultsGPU);
BrokenLine::karimaki_circle_fit& circle = (*circleGPU);
BrokenLine::line_fit& line = (*lineGPU);

// the circle fit gives k, but here we want p_t, so let's change the parameter and the covariance matrix
Matrix3d& Jacob=(*JacobGPU);
Jacob << 1,0,0,
0,1,0,
0,0,-abs(circle.par(2))*B/(BrokenLine::sqr(circle.par(2))*circle.par(2));
circle.par(2)=B/abs(circle.par(2));
circle.cov=Jacob*circle.cov*Jacob.transpose();

helix.par << circle.par, line.par;
helix.cov=MatrixXd::Zero(5, 5);
helix.cov.block(0,0,3,3)=circle.cov;
helix.cov.block(3,3,2,2)=line.cov;
helix.q=circle.q;
helix.chi2_circle=circle.chi2;
helix.chi2_line=line.chi2;

//(*helix_fit_resultsGPU) = BrokenLine::Helix_fit(*hits, *hits_cov, B);
}

__global__
void kernelFastFit(Rfit::Matrix3xNd * hits, Vector4d * results) {
(*results) = Rfit::Fast_fit(*hits);
#endif
circle_fit_resultsGPU[i] =
Rfit::Circle_fit(hits.block(0,0,2,n), hits_cov,
fast_fit_input, rad, B, true);
#ifdef TEST_DEBUG
if (0==i) {
printf("Circle param %f,%f,%f\n",circle_fit_resultsGPU[i].par(0),circle_fit_resultsGPU[i].par(1),circle_fit_resultsGPU[i].par(2));
}

__global__
void kernelCircleFit(Rfit::Matrix3xNd * hits,
Rfit::Matrix3Nd * hits_cov, Vector4d * fast_fit_input, double B,
Rfit::circle_fit * circle_fit_resultsGPU) {
u_int n = hits->cols();
Rfit::VectorNd rad = (hits->block(0, 0, 2, n).colwise().norm());

#if TEST_DEBUG
printf("fast_fit_input(0): %f\n", (*fast_fit_input)(0));
printf("fast_fit_input(1): %f\n", (*fast_fit_input)(1));
printf("fast_fit_input(2): %f\n", (*fast_fit_input)(2));
printf("fast_fit_input(3): %f\n", (*fast_fit_input)(3));
printf("rad(0,0): %f\n", rad(0,0));
printf("rad(1,1): %f\n", rad(1,1));
printf("rad(2,2): %f\n", rad(2,2));
printf("hits_cov(0,0): %f\n", (*hits_cov)(0,0));
printf("hits_cov(1,1): %f\n", (*hits_cov)(1,1));
printf("hits_cov(2,2): %f\n", (*hits_cov)(2,2));
printf("hits_cov(11,11): %f\n", (*hits_cov)(11,11));
printf("B: %f\n", B);
#endif
(*circle_fit_resultsGPU) =
Rfit::Circle_fit(hits->block(0,0,2,n), hits_cov->block(0, 0, 2 * n, 2 * n),
*fast_fit_input, rad, B, false);
}

template<int N>
Expand Down Expand Up @@ -428,83 +355,15 @@ void testFit() {

}

void testFitOneGo(bool errors, double epsilon=1e-6) {
constexpr double B = 0.0113921;
Rfit::Matrix3xNd hits(3,4);
Rfit::Matrix3Nd hits_cov = MatrixXd::Zero(12,12);

fillHitsAndHitsCov(hits, hits_cov);

// FAST_FIT_CPU
Vector4d fast_fit_results = Rfit::Fast_fit(hits);
// CIRCLE_FIT CPU
u_int n = hits.cols();
Rfit::VectorNd rad = (hits.block(0, 0, 2, n).colwise().norm());

Rfit::circle_fit circle_fit_results = Rfit::Circle_fit(hits.block(0, 0, 2, n),
hits_cov.block(0, 0, 2 * n, 2 * n),
fast_fit_results, rad, B, errors);
// LINE_FIT CPU
Rfit::line_fit line_fit_results = Rfit::Line_fit(hits, hits_cov, circle_fit_results,
fast_fit_results, errors);

// FIT GPU
std::cout << "GPU FIT" << std::endl;
Rfit::Matrix3xNd * hitsGPU = nullptr; // new Rfit::Matrix3xNd(3,4);
Rfit::Matrix3Nd * hits_covGPU = nullptr;
Rfit::line_fit * line_fit_resultsGPU = nullptr;
Rfit::line_fit * line_fit_resultsGPUret = new Rfit::line_fit();
Rfit::circle_fit * circle_fit_resultsGPU = nullptr; // new Rfit::circle_fit();
Rfit::circle_fit * circle_fit_resultsGPUret = new Rfit::circle_fit();

cudaCheck(cudaMalloc((void **)&hitsGPU, sizeof(Rfit::Matrix3xNd(3,4))));
cudaCheck(cudaMalloc((void **)&hits_covGPU, sizeof(Rfit::Matrix3Nd(12,12))));
cudaCheck(cudaMalloc((void **)&line_fit_resultsGPU, sizeof(Rfit::line_fit)));
cudaCheck(cudaMalloc((void **)&circle_fit_resultsGPU, sizeof(Rfit::circle_fit)));
cudaCheck(cudaMemcpy(hitsGPU, &hits, sizeof(Rfit::Matrix3xNd(3,4)), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(hits_covGPU, &hits_cov, sizeof(Rfit::Matrix3Nd(12,12)), cudaMemcpyHostToDevice));

kernelFullFit<<<1, 1>>>(hitsGPU, hits_covGPU, B, errors,
circle_fit_resultsGPU, line_fit_resultsGPU);
cudaCheck(cudaDeviceSynchronize());

cudaCheck(cudaMemcpy(circle_fit_resultsGPUret, circle_fit_resultsGPU, sizeof(Rfit::circle_fit), cudaMemcpyDeviceToHost));
cudaCheck(cudaMemcpy(line_fit_resultsGPUret, line_fit_resultsGPU, sizeof(Rfit::line_fit), cudaMemcpyDeviceToHost));

std::cout << "Fitted values (CircleFit) CPU:\n" << circle_fit_results.par << std::endl;
std::cout << "Fitted values (LineFit): CPU\n" << line_fit_results.par << std::endl;
std::cout << "Fitted values (CircleFit) GPU:\n" << circle_fit_resultsGPUret->par << std::endl;
std::cout << "Fitted values (LineFit): GPU\n" << line_fit_resultsGPUret->par << std::endl;
assert(isEqualFuzzy(circle_fit_results.par, circle_fit_resultsGPUret->par, epsilon));
assert(isEqualFuzzy(line_fit_results.par, line_fit_resultsGPUret->par, epsilon));

cudaCheck(cudaFree(hitsGPU));
cudaCheck(cudaFree(hits_covGPU));
cudaCheck(cudaFree(line_fit_resultsGPU));
cudaCheck(cudaFree(circle_fit_resultsGPU));
delete line_fit_resultsGPUret;
delete circle_fit_resultsGPUret;

cudaDeviceReset();
}
int main (int argc, char * argv[]) {
exitSansCUDADevices();

testFit<4>();
testFit<3>();
testFit<5>();

std::cout << "TEST FIT, NO ERRORS" << std::endl;

int main (int argc, char * argv[]) {
// testFit();
/*std::cout << "TEST FIT, NO ERRORS" << std::endl;
testFitOneGo(false);
std::cout << "TEST FIT, ERRORS AND SCATTER" << std::endl;
testFitOneGo(true, 1e-5);*/

std::cout << "TEST BROKEN LINE" << std::endl;
testBrokenLineOneGo(1e-5);

return 0;
return 0;
}

9 changes: 5 additions & 4 deletions RecoPixelVertexing/PixelTrackFitting/test/testEigenGPUNoFit.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,8 @@
#include "test_common.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"

#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"

using namespace Eigen;

using Matrix5d = Matrix<double, 5, 5>;
Expand Down Expand Up @@ -218,14 +220,13 @@ std::cout << "*************************\n\n" << std::endl;


int main (int argc, char * argv[]) {

//cudaDeviceSetLimit(cudaLimitStackSize, 8500);
//cudaCheck(cudaDeviceSynchronize());
exitSansCUDADevices();

testEigenvalues();
testInverse3x3();
testInverse4x4();
testInverse5x5();
testInverse5x5();

testMultiply<1, 2, 2, 1>();
testMultiply<1, 2, 2, 2>();
testMultiply<1, 2, 2, 3>();
Expand Down
2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelVertexFinding/src/PixelVertexHeterogeneousProducer.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -206,7 +206,7 @@ void PixelVertexHeterogeneousProducer::produceGPUCuda(
assert(it< (*tuples_).indToEdm.size());
auto k = (*tuples_).indToEdm[it];
if (k>tracks.size()) {
std::cout << "oops track " << it << " does not exists on CPU " << k << std::endl;
edm::LogWarning("PixelVertexHeterogeneousProducer") << "oops track " << it << " does not exists on CPU " << k;
continue;
}
auto tk = reco::TrackRef(trackCollection, k);
Expand Down
1 change: 1 addition & 0 deletions SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationHeterogeneous.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,7 @@ class ClusterTPAssociationHeterogeneous : public HeterogeneousEDProducer<heterog
using CPUProduct = trackerHitAssociationHeterogeneousProduct::CPUProduct;
using Output = trackerHitAssociationHeterogeneousProduct::ClusterTPAHeterogeneousProduct;


using Clus2TP = ClusterSLGPU::Clus2TP;

explicit ClusterTPAssociationHeterogeneous(const edm::ParameterSet&);
Expand Down

0 comments on commit 20adc27

Please sign in to comment.