GPU: better hits #81

VinInn · 2018-06-12T14:03:36Z

Migrated PixelRecHit producer to Heterogeneous (including a cpu product)
Data structures on gpu now include everything needed for Doublets, CA and fit.
Layer splitting done: phi sorting (or partial sorting) requires #69.
took the opportunity for some cleanup and bug fixes.

cmsbot · 2018-06-12T14:03:54Z

A new Pull Request was created by @VinInn (Vincenzo Innocente) for CMSSW_10_2_X_Patatrack.

It involves the following packages:

Geometry/TrackerGeometryBuilder
RecoLocalTracker/Configuration
RecoLocalTracker/SiPixelRecHits

@cmsbot, @fwyzard can you please review it and eventually sign? Thanks.

cms-bot commands are listed here

makortel · 2018-06-12T14:10:08Z

RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu

+
+    cudaCheck(cudaMalloc((void**) & gpu_d, sizeof(HitsOnGPU)));
+    cudaCheck(cudaMemcpy(gpu_d, &gpu_, sizeof(HitsOnGPU), cudaMemcpyDefault));
+    cudaCheck(cudaDeviceSynchronize());


Minor detail, but the cudaMemcpy could be made asynchronous by passing the CUDA stream from here

cmssw/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitHeterogeneous.cc

Lines 153 to 155 in 67bbc84

void SiPixelRecHitHeterogeneous::beginStreamGPUCuda(edm::StreamID streamId, cuda::stream_t<>& cudaStream) {

gpuAlgo_ = std::make_unique<pixelgpudetails::PixelRecHitGPUKernel>();

}

makortel · 2018-06-12T14:13:01Z

RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu

+    for (int i=0;i<10;++i) std::cout << phase1PixelTopology::layerName[i] << ':' << hitsLayerStart[i] << ' ';
+    std::cout << "end:" << hitsLayerStart[10] << std::endl;
+
+    cudaCheck(cudaMemcpyAsync(gpu_.hitsLayerStart_d, hitsLayerStart, (11) * sizeof(uint32_t), cudaMemcpyDefault, stream.id()));


I believe this is not safe because hitsLayerStart is a local array. It should have a life time longer than the async copy takes (effectively to be a member of PixelRecHitGPUKernel).

makortel · 2018-06-12T14:15:33Z

RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitHeterogeneous.cc


-  output->shrink_to_fit();
+  output->collection.shrink_to_fit();
  iEvent.put(std::move(output));


This should be changed to iEvent.put<Output>(std::move(output)) to produce the HeterogeneousProduct for CPU.

makortel · 2018-06-12T14:18:51Z

RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py

 siPixelRecHitsPreSplitting = siPixelRecHits.clone(
    src = 'siPixelClustersPreSplitting'
 )
+
+
+from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import siPixelRecHitHeterogeneous


In #77 when doing the same for raw2cluster I put this import to RecoLocalTracker_cff.py.

makortel · 2018-06-12T14:25:22Z

RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitHeterogeneous.cc


-  iEvent.put(std::move(output));
+  iEvent.put<Output>(std::move(output), [this, hits, hclusters](const GPUProduct&, CPUProduct& cpu) {
+      this->convertGPUtoCPU(hclusters, hits, cpu);


You could avoid capturing hits by using the HitsOnCPU as the GPUProduct.

what is the lifetime of the lambda? in principle in this way once the lambda is gone hits are gone as well
while as part of GPUProduct will stay in the event

The lambda is store as part of the HeterogeneousProduct' (as is the GPUProduct`), so it will stay to the end of the event anyway.

makortel · 2018-06-12T14:30:18Z

RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py

+from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import siPixelRecHitHeterogeneous
+
+from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneousConverter_cfi import siPixelRecHitHeterogeneousConverter as _siPixelRecHitHeterogeneousConverter
+gpu.toReplaceWith(siPixelRecHitsPreSplitting, _siPixelRecHitHeterogeneousConverter.clone())


Replacing siPixelRecHits before the clone() to ...PreSplitting would work as well (then in principle the default src should be kept as siPixelClusters).

Combining with my comment above, this file could be left untouched.

makortel · 2018-06-12T14:37:47Z

Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h

+  constexpr char const * layerName[10] = {"BL1","BL2","BL3","BL4",
+   		  	                  "E+1", "E+2", "E+3",
+			                  "E-1", "E-2", "E-3"
+                                          };


Just to note that this introduces yet another convention for naming pixel layers. It seems to be only for debug prints, so I don't object.

@makortel what are the "usual" pixel layer names, and where are they defined ?

We (=CMS) have many conventions in different places

Seeding layers are specified with BPixN and FPixN_pos/FPixN_neg

cmssw/RecoTracker/TkSeedingLayers/python/PixelLayerTriplets_cfi.py

Lines 6 to 11 in 96559f3

PixelLayerTriplets.layerList = cms.vstring('BPix1+BPix2+BPix3',

'BPix1+BPix2+FPix1_pos',

'BPix1+BPix2+FPix1_neg',

'BPix1+FPix1_pos+FPix2_pos',

'BPix1+FPix1_neg+FPix2_neg'

)

PixelSubdetector uses enums PixelBarrel and PixelForward

cmssw/DataFormats/SiPixelDetId/interface/PixelSubdetector.h

Line 11 in 96559f3

enum SubDetector {PixelBarrel=1,PixelEndcap=2};

GeomDetEnumerators uses enums PixelBarrel, PixelEndcap, P1PXB, P1PXEC, P2PXB, P2PXEC

cmssw/Geometry/CommonDetUnit/interface/GeomDetEnumerators.h

Line 11 in 96559f3

enum SubDetector {PixelBarrel, PixelEndcap, TIB, TOB, TID, TEC, CSC, DT, RPCBarrel, RPCEndcap, GEM, ME0, P2OTB, P2OTEC, P1PXB, P1PXEC, P2PXB, P2PXEC, TimingBarrel, TimingEndcap, invalidDet};

TrackerTopology uses pxb and pxf in function prefixes

cmssw/DataFormats/TrackerCommon/interface/TrackerTopology.h

Lines 160 to 163 in 96559f3

unsigned int pxbModule(const DetId &id) const {

return ((id.rawId()>>pbVals_.moduleStartBit_)& pbVals_.moduleMask_);

}

unsigned int pxfModule(const DetId &id) const {

TrackingNtuple python library uses BPixN and FPixN+/FPixN- in printouts

cmssw/Validation/RecoTrack/python/plotting/ntupleDataFormat.py

Lines 100 to 112 in 96559f3

if subdet in [SubDet.FPix, SubDet.TID, SubDet.TEC] or isPhase2OTBarrel:

sideNum = get("side")

if sideNum == 1:

side = "-"

elif sideNum == 2:

side = "+"

elif isPhase2OTBarrel and sideNum == 3:

side = ""

else:

side = "?"

return "%s%d%s" % (SubDet.toString(subdet),

getattr(self._tree, self._prefix+"_layer")[self._index],

side)

I don't think none of these is authoritative enough to suggest a change in here (since they're for printouts only, for configuration input's I'd probably suggest the seeding layers' convention). Anyway these have the nice property (on purpose?) that the length of BPix and FPix strings are the same (that none of the other conventions have).

fwyzard · 2018-06-14T14:27:56Z

Apart from the requested changes and merge conflicts, I have the impression that the GPU workflows (10824.8) do not produce any tracks any more ?

makortel · 2018-06-15T07:03:35Z

Apart from the requested changes and merge conflicts, I have the impression that the GPU workflows (10824.8) do not produce any tracks any more ?

From https://fwyzard.web.cern.ch/fwyzard/patatrack/pulls/42f6f40e4973cddd77d4138005948039b75ff1ca/RelValTTbar_13-CMSSW_10_2_0_pre3-PU25ns_101X_upgrade2018_realistic_v7-v1/10824.8/plots_summary.html
there are 303 tracks (all fake) with this PR while the "development" gives 42766.

VinInn · 2018-06-15T07:10:26Z

impressive!
I suspect I know why... (local vs global coords)

VinInn · 2018-06-15T07:11:04Z

btw what is the new baseline w/r/t I have to sync?

fwyzard · 2018-06-15T08:40:14Z

CMSSW_10_2_0_pre5_Patatrack - I think the instructions [here](at https://patatrack.web.cern.ch/patatrack/wiki/patatrackdevelopment.html) are up to date.

VinInn · 2018-06-15T09:02:12Z

I started from CMSSW_10_2_0_pre5_Patatrack ...
will have a look to the conflicts...

cmsbot · 2018-06-16T08:49:48Z

Pull request #81 was updated. @cmsbot, @fwyzard can you please check and sign again.

cmsbot · 2018-06-16T08:59:17Z

Pull request #81 was updated. @cmsbot, @fwyzard can you please check and sign again.

fwyzard · 2018-06-16T13:57:49Z

Indeed, now this PR reproduces the results of the development branch.

fwyzard · 2018-06-29T08:17:43Z

Validation summary

Reference release CMSSW_10_2_0_pre5 at 30c7b03
Development branch CMSSW_10_2_X_Patatrack at 10d59f2
Testing PRs:

GPU: better hits #81 at 4620fdf

`makeTrackValidationPlots.py` plots

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

tracking validation plots for workflow 10824.5
tracking validation plots for workflow 10824.8
tracking validation plots for workflow 10824.7
tracking validation plots for workflow 10824.9

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

tracking validation plots for workflow 10824.5
tracking validation plots for workflow 10824.8 are missing
tracking validation plots for workflow 10824.7
tracking validation plots for workflow 10824.9

DQM GUI plots

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

reference DQM plots for reference release, workflow 10824.5
DQM plots for development release, workflow 10824.5
DQM plots for development release, workflow 10824.8
DQM plots for development release, workflow 10824.7
DQM plots for development release, workflow 10824.9
DQM plots for testing release, workflow 10824.5
DQM plots for testing release, workflow 10824.8
DQM plots for testing release, workflow 10824.7
DQM plots for testing release, workflow 10824.9
DQM comparison for reference workflow 10824.5
DQM comparison for workflow 10824.8
DQM comparison for workflow 10824.7
DQM comparison for workflow 10824.9

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

reference DQM plots for reference release, workflow 10824.5
DQM plots for development release, workflow 10824.5
DQM plots for development release, workflow 10824.8 are missing
DQM plots for development release, workflow 10824.7
DQM plots for development release, workflow 10824.9
DQM plots for testing release, workflow 10824.5
DQM plots for testing release, workflow 10824.8 are missing
DQM plots for testing release, workflow 10824.7
DQM plots for testing release, workflow 10824.9
DQM comparison for reference workflow 10824.5
DQM comparison for workflow 10824.8
DQM comparison for workflow 10824.7
DQM comparison for workflow 10824.9

logs and `nvprof/nvvp` profiles

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

reference log, visual profile and summary for workflow 10824.5
development log, visual profile and summary for workflow 10824.5
development log, visual profile and summary for workflow 10824.8
development log, visual profile and summary for workflow 10824.7
development log, visual profile and summary for workflow 10824.9
testing log, visual profile and summary for workflow 10824.5
testing log, visual profile and summary for workflow 10824.8
testing log, visual profile and summary for workflow 10824.7
testing log, visual profile and summary for workflow 10824.9

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

reference log, visual profile and summary for workflow 10824.5
development log, visual profile and summary for workflow 10824.5
development log, visual profile and summary for workflow 10824.8
development log, visual profile and summary for workflow 10824.7
development log, visual profile and summary for workflow 10824.9
testing log, visual profile and summary for workflow 10824.5
testing log, visual profile and summary for workflow 10824.8
testing log, visual profile and summary for workflow 10824.7
testing log, visual profile and summary for workflow 10824.9

Logs

The full log is available at https://fwyzard.web.cern.ch/fwyzard/patatrack/pulls/d9fd2c79bc3226ea452e21c0d0431bed4b281e91/log .

Migrate PixelRecHit EDProducer to HeterogeneousEDProducer, including the cpu product. Data structures on gpu now include everything needed for Doublets, CA and fit. Layer splitting done: phi sorting (or partial sorting) requires #69. Includes some cleanup and bug fixes.

makortel · 2018-06-29T19:15:13Z

RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu

+
+    std::cout << "hit layerStart "; 
+    for (int i=0;i<10;++i) std::cout << phase1PixelTopology::layerName[i] << ':' << hitsLayerStart_[i] << ' ';
+    std::cout << "end:" << hitsLayerStart_[10] << std::endl;


By the way, now we have a per-event (non-thread safe) printout from here.

will be removed at next iteration when the sorting will be implemented as well

plots for DQM tau

Migrate PixelRecHit EDProducer to HeterogeneousEDProducer, including the cpu product. Data structures on gpu now include everything needed for Doublets, CA and fit. Layer splitting done: phi sorting (or partial sorting) requires #69. Includes some cleanup and bug fixes.

VinInn added 8 commits June 11, 2018 17:00

hits improved to support CA and fitting

5412113

move self pointer out of struct

777a84d

move product in its own header

d713b2a

moved to produce HeterogeneousPixelRecHits

65a94a1

add converter to legacy

fd12ed6

works

b7ee911

use minCol to match clusters

72c5952

fix indexing logic

67bbc84

VinInn requested a review from makortel June 12, 2018 14:03

cmsbot added comparison-pending labels Jun 12, 2018

makortel requested changes Jun 12, 2018

View reviewed changes

This comment has been minimized.

Sign in to view

Merge branch 'CMSSW_10_2_X_Patatrack' into gpuMoreHits

2fafecf

VinInn added 2 commits June 16, 2018 10:50

Merged gpuMoreHits from repository VinInn with cms-merge-topic

7a2edc2

transfer local coord to cpu

dcedee7

This comment has been minimized.

Sign in to view

fwyzard merged commit a9f0289 into cms-patatrack:CMSSW_10_2_X_Patatrack Jun 29, 2018

makortel mentioned this pull request Jun 29, 2018

Various fixes and cleanup #87

Merged

makortel reviewed Jun 29, 2018

View reviewed changes

fwyzard removed comparison-pending labels Jul 5, 2018

fwyzard added this to the CMSSW_10_2_0_pre6_Patatrack milestone Jul 5, 2018

fwyzard pushed a commit that referenced this pull request Nov 1, 2018

Merge pull request #81 from aspiezia/CMSSW_10_3_X_tau_pog_DQMupdate

74a413a

plots for DQM tau

fwyzard mentioned this pull request Oct 8, 2020

Patatrack integration - Pixel local reconstruction (9/N) cms-sw/cmssw#31721

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

GPU: better hits #81

GPU: better hits #81

VinInn commented Jun 12, 2018

cmsbot commented Jun 12, 2018

makortel Jun 12, 2018

makortel Jun 12, 2018

makortel Jun 12, 2018

makortel Jun 12, 2018

makortel Jun 12, 2018

VinInn Jun 16, 2018

makortel Jun 18, 2018

makortel Jun 12, 2018

makortel Jun 12, 2018

fwyzard Jun 16, 2018

makortel Jun 18, 2018

This comment has been minimized.

fwyzard commented Jun 14, 2018

makortel commented Jun 15, 2018

VinInn commented Jun 15, 2018

VinInn commented Jun 15, 2018

fwyzard commented Jun 15, 2018

VinInn commented Jun 15, 2018

cmsbot commented Jun 16, 2018

cmsbot commented Jun 16, 2018

This comment has been minimized.

fwyzard commented Jun 16, 2018

fwyzard commented Jun 29, 2018

makortel Jun 29, 2018

VinInn Jun 30, 2018

	void SiPixelRecHitHeterogeneous::beginStreamGPUCuda(edm::StreamID streamId, cuda::stream_t<>& cudaStream) {
	gpuAlgo_ = std::make_unique<pixelgpudetails::PixelRecHitGPUKernel>();
	}

	PixelLayerTriplets.layerList = cms.vstring('BPix1+BPix2+BPix3',
	'BPix1+BPix2+FPix1_pos',
	'BPix1+BPix2+FPix1_neg',
	'BPix1+FPix1_pos+FPix2_pos',
	'BPix1+FPix1_neg+FPix2_neg'
	)

	unsigned int pxbModule(const DetId &id) const {
	return ((id.rawId()>>pbVals_.moduleStartBit_)& pbVals_.moduleMask_);
	}
	unsigned int pxfModule(const DetId &id) const {

	if subdet in [SubDet.FPix, SubDet.TID, SubDet.TEC] or isPhase2OTBarrel:
	sideNum = get("side")
	if sideNum == 1:
	side = "-"
	elif sideNum == 2:
	side = "+"
	elif isPhase2OTBarrel and sideNum == 3:
	side = ""
	else:
	side = "?"
	return "%s%d%s" % (SubDet.toString(subdet),
	getattr(self._tree, self._prefix+"_layer")[self._index],
	side)

GPU: better hits #81

GPU: better hits #81

Conversation

VinInn commented Jun 12, 2018

cmsbot commented Jun 12, 2018

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

This comment has been minimized.

fwyzard commented Jun 14, 2018

makortel commented Jun 15, 2018

VinInn commented Jun 15, 2018

VinInn commented Jun 15, 2018

fwyzard commented Jun 15, 2018

VinInn commented Jun 15, 2018

cmsbot commented Jun 16, 2018

cmsbot commented Jun 16, 2018

This comment has been minimized.

fwyzard commented Jun 16, 2018

fwyzard commented Jun 29, 2018

Validation summary

makeTrackValidationPlots.py plots

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

DQM GUI plots

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

logs and nvprof/nvvp profiles

/RelValTTbar_13/CMSSW_10_2_0_pre5-PU25ns_102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

/RelValZMM_13/CMSSW_10_2_0_pre5-102X_upgrade2018_realistic_v1-v1/GEN-SIM-DIGI-RAW

Logs

Choose a reason for hiding this comment

Choose a reason for hiding this comment

`makeTrackValidationPlots.py` plots

logs and `nvprof/nvvp` profiles