GPU Shared AtPoints Bases #1711
Open
+1,570
−36
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
jeremylt
force-pushed
the
jeremy/shared-at-points
branch
6 times, most recently
from
37e9e49
to
514184d
Compare
|
Yay - it works locally but fails on Nother. Those are always fun to debug |
|
Cuda is more betterer now. Need to check thread block sizes setup for Hip for 3D |
jeremylt
force-pushed
the
jeremy/shared-at-points
branch
7 times, most recently
from
dcfc06f
to
ff3d054
Compare
|
Ok, now every element other than the first in 3D is wrong for HIP. Progress, but super bizarre how it doesn't seem to behave as I expect. Almost the same logic as CUDA so I have to miss some subtle change between HIP shared and CUDA shared as those two have slightly diverged. |
jeremylt
force-pushed
the
jeremy/shared-at-points
branch
from
88a3d5f
to
ebcc04c
Compare
jeremylt
force-pushed
the
jeremy/shared-at-points
branch
from
ebcc04c
to
1b3d9bd
Compare
|
Ok, the kernels work now. We can optimize them in the future |
|
Local testing for Ratel passes |
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 adds AtPoints to
/gpu/[cuda,hip]/shared, which is a blocker for a/gpu/[cuda,hip]/genAtPoints capability.Work in progress. Thet35*series tests pass, but thet59*series don't all pass.There seems to be an issue with the 2D/3D transpose interp and grad with multiple elements.Passing for libCEED t* tests.
petsc-bpsswarmisn't passing yet though. There is also some issue with a few Ratel tests where they get stuck, probably on a__syncthreads(). (MMS tests it seems).