Fix race conditions in CUDA kernels causing incorrect energies #2883
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
ye-luo
reviewed
Feb 8, 2021
There was a problem hiding this comment.
Does this make deterministic tests passing more robustly?
https://cdash.qmcpack.org/CDash/viewTest.php?buildid=183717&onlydelta
| @@ -312,6 +312,7 @@ __global__ void find_core_electrons_PBC_kernel(T** R, | |||
| disp[tid][1] = r[tid][1] - i[ion][1]; | |||
| disp[tid][2] = r[tid][2] - i[ion][2]; | |||
| dist[tid] = min_dist<T>(disp[tid][0], disp[tid][1], disp[tid][2], L, Linv); | |||
| __syncthreads(); | |||
| @@ -1412,6 +1412,7 @@ __global__ void calc_ratio_grad_lapl(T** Ainv_list, | |||
| // ratio to make it w.r.t. new position | |||
| if (tid < 4) | |||
| ratio_prod[(tid + 1) * BS1] /= ratio_prod[0]; | |||
| __syncthreads(); | |||
|
I think the bug is not releated to Volta independent thread scheduling. Deterministic tests are failing on my Pascal as well. I think more efficient warp scheduling tricks the bug. The diamond test just has too low electron counts. |
ye-luo
approved these changes
Feb 8, 2021
|
As you mention, some of the deterministic tests are failing. I don't think any of these changes will help with them. An isolated water molecule in a periodic supercell tripped the problem. This has the same electron count as our diamond tests. However the wavefunction is large. Possibly we could make a workflow test out of it. There are some excess syncthreads calls to try to quiet racecheck. Since they didn't slow the code noticeably and they make the code a bit easier to "parse" mentally for threading issues, I prefer to keep them. |
|
Actually the diamondC_2x1x1 are working locally, see below. Hopefully the runs on bora will also improve: |
|
The bora failure is sporadic. let us hope your fix helps it. |
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.
Proposed changes
Fix reported bugs by Andrea Zen / Dario Alfe where the legacy CUDA code will give clearly wrong energies for certain calculations. The symptom is a total energy as well as several subcomponents of the total energy that vary significantly with walker count. This problem was initially discovered via non-sensical DMC results, but the problems also occur in VMC.
Bugs are traced to handful of potential race conditions in the kernels associated with Coulomb energy evaluation. These bugs are present in Volta generation and more recent GPUs due to their independent thread scheduling. Our regular testing for e.g. carbon diamond did not identify them.
A side effect of these fixes is that the legacy CUDA code now appears to be completely deterministic.
Despite the extra syncthreads calls, running the performance tests showed no change in runtime within measurement error.
I do not have a small reproducer for the problems.
Thanks to @jefflarkin for reminders of CUDA tools (cuda-memcheck) that helped track these problems down efficiently.
The racecheck tool still issues a warning in find_core_electrons_PBC_kernel in NLPP.cu . It looks to be a false warning but needs more study. Otherwise the code is clean of races in VMC for the standard spline+j1+j2 inputs that I tried.
The initcheck tool flags a couple of uninitialized memory references associated with determinant setup and update.
What type(s) of changes does this code introduce?
Does this introduce a breaking change?
What systems has this change been tested on?
nitrogen: AMD Rome + NVIDIA V100 32GB. LLVM dev + CUDA 11.2.0
Checklist
( Created the PR from within vscode, hence initial draft. )