Fix setting the data pointer of error SimpleVector #236

makortel · 2018-12-19T20:26:35Z

While working on Raw2Cluster migration to #100, I realized that #109 removed re-setting the data pointer of the erro SimpleVector to point to the host memory after the cudaMemcpyAsync() in

cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu

Lines 586 to 587 in 6f55d70

    
           cudaCheck(cudaMemcpyAsync(digis_clusters_h.error.get(), error_d.get(), vsize, cudaMemcpyDefault, stream.id())); 
        
           cudaCheck(cudaMemcpyAsync(digis_clusters_h.data.get(), data_d.get(), MAX_ERROR_SIZE, cudaMemcpyDefault, stream.id()));

We haven't hit on the problem because on the consumer side

cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterHeterogeneous.cc

Lines 644 to 646 in 6f55d70

    
           auto size = digis_clusters_h.error->size(); 
        
           for (auto i = 0; i < size; i++) { 
        
             pixelgpudetails::error_obj err = (*digis_clusters_h.error)[i];

the data is accessed only if size() > 0, and for MC the size appears to be 0.

This PR sets the data pointer in a proper place (after the async copies have completed).

fwyzard · 2018-12-20T08:39:53Z

Hi @makortel,
you mean the fact that we dropped error_h->set_data(data_h) because we did not need to get error_h->size() any more - but we should have it back to copy the errors from GPU to CPU ?

makortel · 2018-12-20T15:11:15Z

@fwyzard Something like that (even though error_h->size() is independent of error_h->set_data(data_h)).

VinInn · 2018-12-20T15:25:49Z

OK! I was just debugging that section of code...

makortel · 2018-12-20T15:32:08Z

And yeah, maybe we should add a data workflow to the tests (to at least ensure that it runs for starters).

VinInn · 2018-12-20T15:41:01Z

In any case I think we should drop GPUSimpleVector in favor of the GPUVecArray that requires much less gymnastics.
GPUSimpleVector is useful only if the capacity needs to be set at runtime

makortel · 2018-12-20T16:02:45Z

In any case I think we should drop GPUSimpleVector in favor of the GPUVecArray that requires much less gymnastics.
GPUSimpleVector is useful only if the capacity needs to be set at runtime

I fully agree (and even for runtime-known capacity we should make a class with a safer interface).

VinInn · 2018-12-20T18:03:44Z

at the moment git cms-merge-topic (git in general for CMSSW I suppose) does not work @SimonS so cannot test on V100

fwyzard · 2019-01-07T15:45:29Z

Validation summary

Reference release CMSSW_10_4_0_pre4 at d74dd18
Development branch CMSSW_10_4_X_Patatrack at ebc280c
Testing PRs:

Fix setting the data pointer of error SimpleVector #236 at 5aeaba2

`makeTrackValidationPlots.py` plots

/RelValTTbar_13/CMSSW_10_4_0_pre3-PU25ns_103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

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

/RelValZMM_13/CMSSW_10_4_0_pre3-103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

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

logs and `nvprof`/`nvvp` profiles

/RelValTTbar_13/CMSSW_10_4_0_pre3-PU25ns_103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

reference release, workflow 10824.5
- step3.py: log, visual profile and summary
development release, workflow 10824.5
- step3.py: log, visual profile and summary
development release, workflow 10824.8
- step3.py: log, visual profile and summary
- profile.py: log, visual profile and summary
- ✔️ cuda-memcheck --tool initcheck --track-unused-memory no (report, log) did not find any errors
- ✔️ cuda-memcheck --tool memcheck --leak-check full --report-api-errors all (report, log) did not find any errors
- ✔️ cuda-memcheck --tool synccheck (report, log) did not find any errors
development release, workflow 10824.7
- step3.py: log, visual profile and summary
development release, workflow 10824.9
- step3.py: log, visual profile and summary
testing release, workflow 10824.5
- step3.py: log, visual profile and summary
testing release, workflow 10824.8
- step3.py: log, visual profile and summary
- profile.py: log, visual profile and summary
- ✔️ cuda-memcheck --tool initcheck --track-unused-memory no (report, log) did not find any errors
- ✔️ cuda-memcheck --tool memcheck --leak-check full --report-api-errors all (report, log) did not find any errors
- ✔️ cuda-memcheck --tool synccheck (report, log) did not find any errors
testing release, workflow 10824.7
- step3.py: log, visual profile and summary
testing release, workflow 10824.9
- step3.py: log, visual profile and summary

/RelValZMM_13/CMSSW_10_4_0_pre3-103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

reference release, workflow 10824.5
- step3.py: log, visual profile and summary
development release, workflow 10824.5
- step3.py: log, visual profile and summary
development release, workflow 10824.8
- step3.py: log, visual profile and summary
- profile.py: log, visual profile and summary
- ✔️ cuda-memcheck --tool initcheck --track-unused-memory no (report, log) did not find any errors
- ✔️ cuda-memcheck --tool memcheck --leak-check full --report-api-errors all (report, log) did not find any errors
- ✔️ cuda-memcheck --tool synccheck (report, log) did not find any errors
development release, workflow 10824.7
- step3.py: log, visual profile and summary
development release, workflow 10824.9
- step3.py: log, visual profile and summary
testing release, workflow 10824.5
- step3.py: log, visual profile and summary
testing release, workflow 10824.8
- step3.py: log, visual profile and summary
- profile.py: log, visual profile and summary
- ✔️ cuda-memcheck --tool initcheck --track-unused-memory no (report, log) did not find any errors
- ✔️ cuda-memcheck --tool memcheck --leak-check full --report-api-errors all (report, log) did not find any errors
- ✔️ cuda-memcheck --tool synccheck (report, log) did not find any errors
testing release, workflow 10824.7
- step3.py: log, visual profile and summary
testing release, workflow 10824.9
- step3.py: log, visual profile and summary

Logs

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

fwyzard · 2019-01-07T23:06:23Z

No impact on physics performance (as expected):

RelValTTbar_13	reference-10824.5	development-10824.5	development-10824.8	testing-10824.8
Efficiency	0.4818	0.4824	0.4842	0.4842
Number of TrackingParticles (after cuts)	5556	5556	5556	5556
Number of matched TrackingParticles	2677	2680	2690	2690
Fake rate	0.0519	0.0517	0.0337	0.0337
Duplicate rate	0.0168	0.0175	0.0176	0.0179
Number of tracks	32452	32480	31796	31800
Number of true tracks	30769	30801	30724	30727
Number of fake tracks	1683	1679	1072	1073
Number of pileup tracks	27093	27118	27000	27002
Number of duplicate tracks	546	567	559	569

fwyzard · 2019-01-07T23:24:31Z

No impact on the throughput:

before

Warming up
Running 4 times over 4200 events with 1 jobs, each with 8 threads, 8 streams and 1 GPUs
  1725.6 ±   2.1 ev/s (4000 events)
  1728.5 ±   2.1 ev/s (4000 events)
  1731.0 ±   2.0 ev/s (4000 events)
  1736.1 ±   1.4 ev/s (4000 events)

after

Warming up
Running 4 times over 4200 events with 1 jobs, each with 8 threads, 8 streams and 1 GPUs
  1736.0 ±   2.3 ev/s (4000 events)
  1735.5 ±   2.4 ev/s (4000 events)
  1730.9 ±   2.0 ev/s (4000 events)
  1723.8 ±   2.3 ev/s (4000 events)

Fix setting the data pointer of error SimpleVector

5aeaba2

This was referenced Dec 21, 2018

Full workflow on GPU #197

Closed

Speed up in clusterizer and doubletFinder #238

Closed

fwyzard added the bug label Jan 7, 2019

fwyzard merged commit eae2fbc into cms-patatrack:CMSSW_10_4_X_Patatrack Jan 7, 2019

fwyzard mentioned this pull request Jan 8, 2019

Make a local copy of CUB's CachingDeviceAllocator and patch it for the race condition #240

Merged

fwyzard added this to the CMSSW_10_4_X_Patatrack milestone Jan 8, 2019

fwyzard added the fixed label Jan 8, 2019

fwyzard pushed a commit that referenced this pull request Oct 8, 2020

Set the data pointer of error's SimpleVector (#236)

37b5849

fwyzard mentioned this pull request Oct 8, 2020

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

Merged

fwyzard pushed a commit that referenced this pull request Oct 19, 2020

Set the data pointer of error's SimpleVector (#236)

99b0d0c

fwyzard pushed a commit that referenced this pull request Oct 20, 2020

Set the data pointer of error's SimpleVector (#236)

aed142f

fwyzard pushed a commit that referenced this pull request Oct 23, 2020

Set the data pointer of error's SimpleVector (#236)

4e9319a

fwyzard pushed a commit that referenced this pull request Nov 6, 2020

Set the data pointer of error's SimpleVector (#236)

096398c

fwyzard pushed a commit that referenced this pull request Nov 16, 2020

Set the data pointer of error's SimpleVector (#236)

8f17ede

fwyzard pushed a commit that referenced this pull request Dec 25, 2020

Set the data pointer of error's SimpleVector (#236)

c00d365

fwyzard pushed a commit that referenced this pull request Dec 29, 2020

Set the data pointer of error's SimpleVector (#236)

fbdfdb0

fwyzard pushed a commit that referenced this pull request Dec 29, 2020

Set the data pointer of error's SimpleVector (#236)

64c573c

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix setting the data pointer of error SimpleVector #236

Fix setting the data pointer of error SimpleVector #236

makortel commented Dec 19, 2018

fwyzard commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

fwyzard commented Jan 7, 2019 •

edited

Loading

fwyzard commented Jan 7, 2019 •

edited

Loading

fwyzard commented Jan 7, 2019

	cudaCheck(cudaMemcpyAsync(digis_clusters_h.error.get(), error_d.get(), vsize, cudaMemcpyDefault, stream.id()));
	cudaCheck(cudaMemcpyAsync(digis_clusters_h.data.get(), data_d.get(), MAX_ERROR_SIZE, cudaMemcpyDefault, stream.id()));

	auto size = digis_clusters_h.error->size();
	for (auto i = 0; i < size; i++) {
	pixelgpudetails::error_obj err = (*digis_clusters_h.error)[i];

Fix setting the data pointer of error SimpleVector #236

Fix setting the data pointer of error SimpleVector #236

Conversation

makortel commented Dec 19, 2018

fwyzard commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

makortel commented Dec 20, 2018

VinInn commented Dec 20, 2018

fwyzard commented Jan 7, 2019 • edited Loading

Validation summary

makeTrackValidationPlots.py plots

/RelValTTbar_13/CMSSW_10_4_0_pre3-PU25ns_103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

/RelValZMM_13/CMSSW_10_4_0_pre3-103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

logs and nvprof/nvvp profiles

/RelValTTbar_13/CMSSW_10_4_0_pre3-PU25ns_103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

/RelValZMM_13/CMSSW_10_4_0_pre3-103X_upgrade2018_realistic_v8-v1/GEN-SIM-DIGI-RAW

Logs

fwyzard commented Jan 7, 2019 • edited Loading

fwyzard commented Jan 7, 2019

before

after

fwyzard commented Jan 7, 2019 •

edited

Loading

`makeTrackValidationPlots.py` plots

logs and `nvprof`/`nvvp` profiles

fwyzard commented Jan 7, 2019 •

edited

Loading