Remove all remaining calls to cudaStreamSynchronize() #109
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Certainly. |
…avoid implicit cudaStreamSynchronize
…Synchronize() We need the number of hits/clusters in RecHit acquire(), but fortunately we can do the necessary calculation already in Raw2Cluster.
|
Rebased. |
Validation summaryReference release CMSSW_10_2_0_pre6 at a674e1f
|
|
Workflows 10824.5 are broken also before this PR, probably due to #105. |
|
Hmm, there is still something synchronizing the CUDA streams... Screenshots from the testing 10824 profile Raw2cluster Rechits In both the acquire phase includes all the queued GPU activity. In contrast, here is the raw2cluster picture from development where clearly the acquire ends before the GPU activity ends. |
|
And the reason is likely here (still to be verified)
So all |
|
Do we have copies from the gpu to the stack on the cpu, or only to the heap ? |
|
I'm going through them, so far only heap (e.g. normal class member variable). I think we cleaned up the stack ones in the last hackathon. |
|
Yes, of course - sorry for the confusion. |
|
#103 alone does not fix the issue. |
|
Fix is in #112. |
fwyzard
pushed a commit
that referenced
this pull request
Oct 8, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Oct 19, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Oct 20, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Oct 23, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Nov 6, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Nov 16, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Dec 25, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Dec 29, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
fwyzard
pushed a commit
that referenced
this pull request
Dec 29, 2020
Remove all of the remaining calls to cudaStreamSynchronize() from the pixel "raw to cluster" workflow. Replace thrust::inclusive_scan with cub::DeviceScan::InclusiveSum to avoid implicit cudaStreamSynchronize and per-event buffer allocations Avoid a data dependency on the number of hits: - in raw2cluster, always transfer the errors for the maximum number of modules. - in rechits, replace the calculation of the total number of hits with the total number of clusters Copy the phase1PixelTopology::layerStart array to the GPU to avoid an extra copy back and forth from the CPU.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR removes all of the remaining calls to
cudaStreamSynchronize():thrust::inclusive_scanwithcub::DeviceScan::InclusiveSumin rechits (though moved to raw2cluster within the last bullet)thrust::inclusive_scanseems to use thecubalgorithm internally, and in addition to the implicitcudaStreamSynchronize()it also allocates+frees the buffer on each call (!)In addition, I noticed that rechits was accessing a host memory after a
cudaMemcpyAsyncto there without synchronization. First I added the synchronization, and then (following a comment) moved the number shuffling to GPU. It feels a bit dumb though to transfer the 11 elements of aconstexpr uint32_tarray to device memory, but apparently something like that is needed.Resolves #79.
No changes expected.
@fwyzard @felicepantaleo @VinInn