Fix backward regression in PR #5712#5825
Open
avbokovoy wants to merge 4 commits into
Open
Conversation
Contributor
|
could you help clarify what cases will use this hip_mixed_d kernel? Please also put it in the PR summary. From previous discussion with Bernard, this is my understanding: could you modify the above based on the new push? cc: @avbokovoy @liligwu |
Contributor
Author
|
@spcyppt The dispatch logic after this PR should look like this: Let me know if you have other questions cc: @liligwu |
Contributor
|
@avbokovoy could you rebase this? There's a conflict and I could not import the PR. I don't have permission to resolve conflicts. |
Summary: CTA kernel optimizations from D86108064 split out: - Enable subwarp shuffle for CTA kernel on ROCm - Adaptive work group sizing based on total_L/total_B ratio - Small-D template parameter override (max_D <= 128) with kFixedMaxVecsPerThread=1, kThreadGroupSize=32, kUseVecBlocking=false for tighter compiler codegen - PROCESS_BLOCK macro for grad accumulation loop unrolling in compute_grad_sum (shared by CTA and warp kernels) - subwarp_reduce_add for ROCm GROUP_REDUCE_ALL_SUM (shared by CTA and warp kernels) This diff should not impact any performance on NVIDIA GPUs. Differential Revision: D102946299 Pulled By: spcyppt
Summary: Warp kernel optimizations from D86108064 split out: - New hip_mixed_d_warp kernel for mixed-D and VBE backward on ROCm - Batch preloading of per-row metadata via warp shuffle - Momentum value preloading (split_precomputation_preload) to eliminate separate global memory reads - Overloaded table_update_kernel accepting pre-resolved placement/offset/optimizer state - AMD __builtin_amdgcn_readlane for efficient broadcasting - Small-D template override (max_D <= 128) for warp kernel This diff should not impact any performance on NVIDIA GPUs. Differential Revision: D102946325 Pulled By: spcyppt
avbokovoy
force-pushed
the
abokovoi/pr_5712_backward_regression
branch
from
1445a35 to
5ff55c7
Compare
Contributor
Author
@spcyppt Rebase is done. Apologies for the delay |
Contributor
|
@spcyppt has imported this pull request. If you are a Meta employee, you can view this in D107407381. |
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
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.
Guarded short segments cases to skip optimized kernel launch. New warp_per_row kernel dispatch logic (
FBGEMM_TBE_ROCM_HIP_BACKWARD_KERNEL=1):