GPU Shared AtPoints Bases #1711
Merged
Conversation
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
jeremylt
force-pushed
the
jeremy/shared-at-points
branch
6 times, most recently
from
37e9e49 to
514184d
Compare
Member
Author
|
Yay - it works locally but fails on Nother. Those are always fun to debug |
Member
Author
|
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
Member
Author
|
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
Member
Author
|
Ok, the kernels work now. We can optimize them in the future |
Member
Author
|
Local testing for Ratel passes |
jeremylt
commented
Nov 26, 2024
include/ceed/jit-source/cuda/cuda-shared-basis-tensor-at-points.h
Outdated
Show resolved
Hide resolved
Open
zatkins-dev
approved these changes
Dec 2, 2024
Collaborator
zatkins-dev
left a comment
zatkins-dev
left a comment
There was a problem hiding this comment.
Overall looks good! I have a few clarifying questions, but I think this is good to merge
include/ceed/jit-source/cuda/cuda-shared-basis-tensor-at-points-templates.h
Show resolved
Hide resolved
co-authored-by: zatkins-dev <zach.atkins@colorado.edu>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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).