HIP: add option to embed static blockDim #105
Closed
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
Using blockDim in hip kernels unfortunately incurs a large overhead, because this (dynamic) information is stored in the dispatch packet located in a host-coherent memory region. Since vkFFT always knows the work group size its going to use, just replace uses of blockDim with these values. This means the load from non-cached memory is avoided, the dispatch pointer doesn't have to be loaded which frees up 2 SGPRs, and some indexing calculations might constant fold better. The added option `useStaticWorkGroupSize` has three possible values: - -1: Disable embedding blockDim sizes, effectively the old behavior - 0: Automatically enable embedding when profitable (always except for RDNA2) - 1: Always enable RDNA is disabled by default because this can actually decrease performance sometimes with the reason not fully known, details at [1] [1]: ROCm/hipamd#53 Co-authored-by: [email protected]
|
Hello, Sorry for the long reply. I am currently finishing version 1.3.0 of VkFFT with major changes to the library design aimed at better maintainability and it already does not use global indexing with blockDim. I will upload it to the develop branch of the VkFFT repo next week. As for this pr, I can confirm that it makes code faster on MI200, however it breaks the global indexing if the axis are swapped (line 25958, for example). So I would rather leave the 1.2.33 version as it is and wait a bit for 1.3.0. Best regards, |
|
This has been implemented in v1.3.0. Thanks! |
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.
Using blockDim in hip kernels unfortunately incurs a large overhead, because this (dynamic) information is stored in the dispatch packet located in a host-coherent memory region. Since vkFFT always knows the work group size its going to use, just replace uses of blockDim with these values.
This means the load from non-cached memory is avoided, the dispatch pointer doesn't have to be loaded which frees up 2 SGPRs, and some indexing calculations might constant fold better.
The added option
useStaticWorkGroupSizehas three possible values:RDNA is disabled by default because this can actually decrease performance sometimes with the reason not fully known, details at 1
Performance of vkFFT sample 3, show 20-200% increase in bandwitdh particularly with smaller ffts. These line up with the numbers obtained from the HIP port of Gromacs where this optimization was originally applied.
Benchmarks with MI210 and V620 (equivalent to RX6900xt)
vkfft_hip_embed_blockdim.ods
Co-authored-by: [email protected]