Fix synchronisation problems in the clusterizer #102
Merged
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
fwyzard
force-pushed
the
various_fixes_and_cleanup_part2
branch
from
be5c775 to
d448e03
Compare
makortel
reviewed
Jul 24, 2018
| moduleStart_d, | ||
| clusInModule_d, moduleId_d, | ||
| clus_d, | ||
| wordCounter | ||
| ); | ||
| cudaDeviceSynchronize(); |
There was a problem hiding this comment.
Is this cudaDeviceSynchronize() leftover from the debugging?
makortel
reviewed
Jul 24, 2018
| #endif | ||
|
|
||
| first += threadIdx.x; | ||
| if (first>= numElements) | ||
| return; |
There was a problem hiding this comment.
For my education, why is this replaced with the
bool active = (first < numElements);
if (active) {?
There was a problem hiding this comment.
I am unsure how __syncthreads is affected by threads that have returned, so I have removed all return statements.
I can see if putting them back reintroduces a problem or not...
|
Looks like we just pushed back the crash to event 400+... back to debugging. |
fwyzard
force-pushed
the
various_fixes_and_cleanup_part2
branch
2 times, most recently
from
177c055 to
0260b3e
Compare
|
@VinInn can you help us looking into a related crash ?
It should crash after few events with :-( |
|
ok. will need to understand your changes |
Validation summaryReference release CMSSW_10_2_0_pre6 at a674e1f
|
|
Not OK, there is still something very wrong with the workflow Zmumu GPU workflow:
|
Improve documentation.
fwyzard
force-pushed
the
various_fixes_and_cleanup_part2
branch
from
0260b3e to
3c139fc
Compare
|
I think I found the error in this PR: was not a very useful loop :-( |
fwyzard
changed the title
makortel
reviewed
Jul 25, 2018
| countModules<<<blocks, threadsPerBlock, 0, stream.id()>>>(moduleInd_d, moduleStart_d, clus_d, wordCounter); | ||
| cudaCheck(cudaGetLastError()); | ||
|
|
||
| // read the number of modules into a data memeber, used by getProduct()) |
|
|
||
| // clusters | ||
| cudaCheck(cudaMemcpyAsync(clus_h, clus_d, wordCounter*sizeof(uint32_t), cudaMemcpyDefault, stream.id())); | ||
| cudaCheck(cudaGetLastError()); |
There was a problem hiding this comment.
I believe the last cudaGetLastError() is not needed after removing the cudaStreamSynchronize(), possible errors in queueing the async copy should get caught already by its return value.
makortel
reviewed
Jul 25, 2018
| break; | ||
| } | ||
| default: if (debug) printf("Cabling check returned unexpected result, status = %i\n", status); | ||
| case(1) : { |
There was a problem hiding this comment.
Now that you "point" to them, I find these parentheses weird.
Validation summaryReference release CMSSW_10_2_0_pre6 at a674e1f
|
Validation summaryReference release CMSSW_10_2_0_pre6 at a674e1f
|
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.
Fix synchronisation problems in the
findCluskernel, and improve its documentation.Avoid unnecessary
cudaStreamSynchronizeinmakeClustersAsync.