CARVIEW |
Navigation Menu
-
Notifications
You must be signed in to change notification settings - Fork 24.7k
[Inductor][ROCm][CK] add CK grouped conv2d fwd kernels to ROCm codegen #137947
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
Conversation
đź”— Helpful Linksđź§Ş See artifacts and rendered test results at hud.pytorch.org/pr/137947
Note: Links to docs will display an error until the docs builds have been completed. âś… You can merge normally! (1 Unrelated Failure)As of commit 2472fbb with merge base 60c1433 ( FLAKY - The following job failed but was likely due to flakiness present on trunk:
This comment was automatically generated by Dr. CI and updates every 15 minutes. |
@pytorchbot label "topic: not user facing" |
Thanks, Max! Could we add some tests? |
Could we add the tests later, along with lowering? |
@pytorchbot rebase -s |
@pytorchbot started a rebase job onto refs/remotes/origin/viable/strict. Check the current status here |
Successfully rebased |
0d829fc
to
75e6711
Compare
Link #125453 |
torch/_inductor/kernel/mm.py
Outdated
if ( | ||
is_nonzero | ||
and use_ck_gemm_template(layout) | ||
and V.graph.sizevars.size_hint(m * n * k, fallback=-1) > 0 |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
was thinking whether it made sense to pull this into use_ck_gemm_template
def torch_layout_to_ck_layouts(torch_layout): | ||
# logically, torch tensors are always NCHW, | ||
# and channels-last memory layout is visible in the strides | ||
if torch_layout.stride[-1] == 1: |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
we can use statically_known_equals(stride[-1], 1)
here for dynamic shape support
) | ||
# NB: when using a fixed list order, most likely we will pick the subset of instances | ||
# which are very similar to each other. Randomizing the choice seems to solve this. | ||
random.seed(-11) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
also see this in ck_universal_gemm_template, might be nice to reference the same seed here!
random.seed(-11) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This is a copy/paste artifact, I might change it if the instances picked for test do not serve well for testing purpose - that is, when we introduce the end-to-end test
@pytorchbot rebase -s |
@pytorchbot started a rebase job onto refs/remotes/origin/viable/strict. Check the current status here |
Successfully rebased |
1089cb3
to
483ce35
Compare
Successfully rebased |
483ce35
to
2472fbb
Compare
@pytorchbot merge |
Merge startedYour change will be merged once all checks pass (ETA 0-4 Hours). Learn more about merging in the wiki. Questions? Feedback? Please reach out to the PyTorch DevX Team |
#137947) Plug into lowering and end to end test in a later PR Instance parsing companion PR ROCm/composable_kernel#1585 Pull Request resolved: #137947 Approved by: https://github.com/ColinPeppler, https://github.com/chenyang78
Plug into lowering and end to end test in a later PR
Instance parsing companion PR ROCm/composable_kernel#1585
cc @jeffdaily @sunway513 @jithunnair-amd @pruthvistony @ROCmSupport @dllehr-amd @jataylo @hongxiayang @naromero77amd @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @yf225 @chenyang78 @kadeng @muchulee8 @ColinPeppler @amjames @desertfire @chauhang @aakhundov @zjing14