Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Refactor grid-stride loop #4190

Open
wants to merge 8 commits into
base: development
Choose a base branch
from

Conversation

WeiqunZhang
Copy link
Member

Move grid-stride loop out of GPU kernels. @ashesh2512 noticed performance issues with grid-stride loops on AMD GPUs in PelePhyscis's large kernels.

Thank @AlexanderSinn for the suggestion implemented in this PR.

@WeiqunZhang WeiqunZhang force-pushed the grid-stride-loop branch 2 times, most recently from 168b2c4 to f2af094 Compare October 12, 2024 18:46
Move grid-stride loop out of GPU kernels. @ashesh2512 noticed performance
issues with grid-stride loops on AMD GPUs in PelePhyscis's large kernels.

Thank @AlexanderSinn for the suggestion implemented in this PR.
@AlexanderSinn
Copy link
Member

Can you also replace the remaining blockDim.x in that file with MT or AMREX_GPU_MAX_THREADS 

@WeiqunZhang
Copy link
Member Author

Can you also replace the remaining blockDim.x in that file with MT or AMREX_GPU_MAX_THREADS

I have been thinking about going through all uses of blockDim and doing it in another PR.

Co-authored-by: Alexander Sinn <64009254+AlexanderSinn@users.noreply.github.com>
@WeiqunZhang WeiqunZhang marked this pull request as ready for review October 16, 2024 17:49
WeiqunZhang and others added 2 commits October 16, 2024 15:19
Co-authored-by: Alexander Sinn <64009254+AlexanderSinn@users.noreply.github.com>
WeiqunZhang and others added 2 commits October 16, 2024 16:02
Co-authored-by: Alexander Sinn <64009254+AlexanderSinn@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants