Fix GPU segfaults in PW USPP calculations#7645
Closed
chengleizheng wants to merge 6 commits into
Closed
Conversation
…t2ia lookup arrays and added #pragma omp parallel for guarded by #ifdef _OPENMP in runner() coord building, runner() force assignment, and type_map() atype assignment.
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
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.
Summary
Fix GPU segfaults in PW USPP calculations caused by CPU code accessing GPU
memory pointers. Added CPU intermediate buffers and D2H/H2D sync patterns
in three USPP-specific code paths.
Changes
hamilt_pw: sPsi reads qq_nt from CPU backup (
qq_nt.ptr) instead ofGPU pointer, fills a CPU buffer, then syncs to GPU for cuBLAS.
elecstate_pw (cal_becsum): becp computed on GPU via gemm then synced
to CPU; auxk1/auxk2 use CPU vectors; aux_gk computed on GPU then synced
to CPU; becsum allocated on CPU (consumed by CPU-only force/stress code).
elecstate_pw (add_usrho): Added D2H/H2D wrappers around real2recip
and recip2real FFT calls, which read/write via CPU loops.
elecstate_pw (addusdens_g): Rewritten to use pure CPU computation
(CPU BLAS zgemm_, CPU radial_fft_q); ylmk0 computed on GPU with H2D'd
gcar then D2H'd back.
Test plan