[Inductor][ROCm][CK] add CK grouped conv2d fwd kernels to ROCm codegen #137947

tenpercent · 2024-10-15T00:49:22Z

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

pytorch-bot · 2024-10-15T00:49:25Z

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/137947

📄 Preview Python docs built from this PR
📄 Preview C++ docs built from this PR
❓ Need help or want to give feedback on the CI? Visit the bot commands wiki or our office hours

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:

inductor-periodic / cuda12.1-py3.10-gcc9-sm80 / test (inductor_torchbench_smoketest_perf, 1, 1, linux.gcp.a100) (gh) (detected as infra flaky with no log or failing log classifier)

This comment was automatically generated by Dr. CI and updates every 15 minutes.

tenpercent · 2024-10-18T17:29:32Z

@pytorchbot label "topic: not user facing"

chenyang78 · 2024-10-19T23:08:02Z

Thanks, Max! Could we add some tests?

tenpercent · 2024-10-20T02:30:48Z

Thanks, Max! Could we add some tests?

Could we add the tests later, along with lowering?

tenpercent · 2024-10-20T02:42:43Z

@pytorchbot rebase -s

pytorchmergebot · 2024-10-20T02:44:03Z

@pytorchbot started a rebase job onto refs/remotes/origin/viable/strict. Check the current status here

pytorchmergebot · 2024-10-20T02:44:07Z

Successfully rebased ck-conv onto refs/remotes/origin/viable/strict, please pull locally before adding more changes (for example, via git checkout ck-conv && git pull --rebase)

tenpercent · 2024-10-21T16:59:47Z

Link #125453

ColinPeppler · 2024-10-21T21:10:12Z

torch/_inductor/kernel/mm.py

+    if (
+        is_nonzero
+        and use_ck_gemm_template(layout)
+        and V.graph.sizevars.size_hint(m * n * k, fallback=-1) > 0


was thinking whether it made sense to pull this into use_ck_gemm_template

ColinPeppler · 2024-10-21T21:13:28Z

torch/_inductor/codegen/rocm/ck_conv_template.py

+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:


we can use statically_known_equals(stride[-1], 1) here for dynamic shape support

https://github.com/pytorch/pytorch/blob/134f6cda7e192046921f44b9565a3b56cd28158d/torch/_inductor/sizevars.py#L340C9-L340C32

ColinPeppler · 2024-10-21T21:16:42Z

torch/_inductor/codegen/rocm/ck_conv_template.py

+        )
+        # 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)


also see this in ck_universal_gemm_template, might be nice to reference the same seed here!

pytorch/torch/_inductor/codegen/rocm/ck_universal_gemm_template.py

Line 421 in 364340c

random.seed(-11)

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

tenpercent · 2024-10-21T23:48:10Z

@pytorchbot rebase -s

pytorchmergebot · 2024-10-21T23:49:39Z

@pytorchbot started a rebase job onto refs/remotes/origin/viable/strict. Check the current status here

pytorchmergebot · 2024-10-21T23:49:43Z

Successfully rebased ck-conv onto refs/remotes/origin/viable/strict, please pull locally before adding more changes (for example, via git checkout ck-conv && git pull --rebase)

pytorchmergebot · 2024-10-22T15:09:41Z

Successfully rebased ck-conv onto refs/remotes/origin/viable/strict, please pull locally before adding more changes (for example, via git checkout ck-conv && git pull --rebase)

tenpercent · 2024-10-22T16:14:27Z

@pytorchbot merge

pytorchmergebot · 2024-10-22T16:16:44Z

Merge started

Your 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

Advanced Debugging

Check the merge workflow status
here

#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

pytorch-bot bot added ciflow/inductor module: inductor labels Oct 15, 2024

pytorchbot added the open source label Oct 15, 2024

tenpercent mentioned this pull request Oct 18, 2024

[PT Inductor] Add parsing grouped conv fwd instances ROCm/composable_kernel#1585

Merged

pytorch-bot bot added the topic: not user facing topic category label Oct 18, 2024

tenpercent changed the title ~~[Draft][Inductor][CK] add conv2d kernels~~ [Draft][Inductor][CK] add CK grouped conv2d fwd kernels to ROCm codegen Oct 18, 2024

pytorch-bot bot added ciflow/rocm Trigger "default" config CI on ROCm module: rocm AMD GPU support for Pytorch labels Oct 18, 2024

tenpercent marked this pull request as ready for review October 18, 2024 17:51

tenpercent requested review from ColinPeppler and chenyang78 October 18, 2024 17:51

tenpercent changed the title ~~[Draft][Inductor][CK] add CK grouped conv2d fwd kernels to ROCm codegen~~ [Inductor][ROCm][CK] add CK grouped conv2d fwd kernels to ROCm codegen Oct 18, 2024

pytorchmergebot force-pushed the ck-conv branch from 0d829fc to 75e6711 Compare October 20, 2024 02:44

ColinPeppler reviewed Oct 21, 2024

View reviewed changes

ColinPeppler approved these changes Oct 21, 2024

View reviewed changes

ColinPeppler reviewed Oct 21, 2024

View reviewed changes

pytorchmergebot force-pushed the ck-conv branch from 1089cb3 to 483ce35 Compare October 21, 2024 23:49

chenyang78 approved these changes Oct 22, 2024

View reviewed changes

tenpercent added 20 commits October 22, 2024 15:09

move null checks past workspace size reset

57aa931

finish debugging a segfault

db6131a

enable channels-last kernels

d834af4

test both nchw and channels-last inputs

f70baa4

refactor layout check

1a94af8

some renamings

952afcf

refactor filter

cc7de72

lints

8a51bb1

move out instance parsing

c10a384

cleanup

d01a862

roll back changes which plug in lowering

ead6a6e

roll back test

16bc0cd

run lintrunner

2e80f01

pacify mypy

3be5544

pacify ruff

01f6719

try fixing lint

ce67996

fix lint one more time

2be30a9

moar lints

718cb7d

statically_known_equals for stride check

ac79150

refactor check whether to use ck for gemm max autotune

2472fbb

pytorchmergebot force-pushed the ck-conv branch from 483ce35 to 2472fbb Compare October 22, 2024 15:09

pytorchmergebot added the merging label Oct 22, 2024

pytorchmergebot added the Merged label Oct 22, 2024

pytorchmergebot closed this in 2dab4cc Oct 22, 2024

pytorchmergebot removed the merging label Oct 22, 2024

[Inductor][ROCm][CK] add CK grouped conv2d fwd kernels to ROCm codegen #137947

[Inductor][ROCm][CK] add CK grouped conv2d fwd kernels to ROCm codegen #137947

Uh oh!

Conversation

tenpercent commented Oct 15, 2024 • edited by pytorch-bot bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pytorch-bot bot commented Oct 15, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/137947

✅ You can merge normally! (1 Unrelated Failure)

Uh oh!

tenpercent commented Oct 18, 2024

Uh oh!

chenyang78 commented Oct 19, 2024

Uh oh!

tenpercent commented Oct 20, 2024

Uh oh!

tenpercent commented Oct 20, 2024

Uh oh!

pytorchmergebot commented Oct 20, 2024

Uh oh!

pytorchmergebot commented Oct 20, 2024

Uh oh!

tenpercent commented Oct 21, 2024

Uh oh!

ColinPeppler Oct 21, 2024

Choose a reason for hiding this comment

Uh oh!

ColinPeppler Oct 21, 2024

Choose a reason for hiding this comment

Uh oh!

ColinPeppler Oct 21, 2024

Choose a reason for hiding this comment

Uh oh!

tenpercent Oct 21, 2024

Choose a reason for hiding this comment

Uh oh!

tenpercent commented Oct 21, 2024

Uh oh!

pytorchmergebot commented Oct 21, 2024

Uh oh!

pytorchmergebot commented Oct 21, 2024

Uh oh!

pytorchmergebot commented Oct 22, 2024

Uh oh!

tenpercent commented Oct 22, 2024

Uh oh!

pytorchmergebot commented Oct 22, 2024

Merge started

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

tenpercent commented Oct 15, 2024 •

edited by pytorch-bot bot

Loading

pytorch-bot bot commented Oct 15, 2024 •

edited

Loading