[ROCm][CK][Inductor] Enable addmm for CK backend to gemm max autotune #130576

tenpercent · 2024-07-11T22:14:30Z

Add functional support for torch.addmm with CK backend. See also #125453

Implementation details

It turns out we can use the same template between addmm and matmul; essentially, matmul is addmm with empty bias
The Python generator in CK was updated to generate the shared cpp template. The pip package can be installed from pip install git+https://github.com/rocm/composable_kernel@add-addmm and will be merged into develop branch after this PR lands to avoid breaking the current matmul

Testing

pytest test/inductor/test_ck_backend.py -k addmm

cc @jeffdaily @sunway513 @jithunnair-amd @pruthvistony @ROCmSupport @dllehr-amd @jataylo @hongxiayang @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @peterbell10 @ipiszy @yf225 @chenyang78 @kadeng @muchulee8 @ColinPeppler @amjames @desertfire @chauhang @aakhundov @zjing14 @asroy @nmacchioni @aska-0096

pytorch-bot · 2024-07-11T22:14:33Z

🔗 Helpful Links

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

📄 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! (2 Unrelated Failures)

As of commit 7139093 with merge base e0514a5 ():

FLAKY - The following jobs failed but were likely due to flakiness present on trunk:

rocm / linux-focal-rocm6.1-py3.8 / test (default, 1, 6, linux.rocm.gpu.2) (gh) (similar failure)
test_transformers.py::TestSDPACudaOnlyCUDA::test_flash_attention_vs_math_ref_grads_batch_size_1_seq_len_q_1024_seq_len_k_1024_head_dim_128_is_causal_False_dropout_p_0_0_bfloat16_scale0_enable_gqa_True_n_heads0_cuda_bfloat16
rocm / linux-focal-rocm6.1-py3.8 / test (default, 6, 6, linux.rocm.gpu.2) (gh) (similar failure)
export/test_export.py::TestOneOffModelExportResult::test_scaled_dot_product_attention_cuda

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

test/inductor/test_ck_backend.py

torch/_inductor/codegen/rocm/ck_universal_gemm_template.py

torch/_inductor/kernel/mm.py

torch/_inductor/codegen/rocm/rocm_kernel.py

torch/_inductor/codegen/rocm/ck_universal_gemm_template.py

aakhundov · 2024-07-24T01:20:51Z

@tenpercent the linter seems failing, please take a look.

zoranzhao

Left some comments

torch/_inductor/kernel/mm.py

torch/_inductor/codegen/rocm/ck_addmm_template.py

tenpercent · 2024-07-25T18:28:07Z

@pytorchbot rebase -s

pytorchmergebot · 2024-07-25T18:29:35Z

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

pytorchmergebot · 2024-07-25T18:29:39Z

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

… ck-addmm

torch/_inductor/utils.py

torch/_inductor/codegen/rocm/ck_universal_gemm_template.py

torch/_inductor/codegen/rocm/rocm_kernel.py

torch/_inductor/config.py

chenyang78

Please fix the lint issues. Also, seems there are some rocm ci failures. Thanks!

tenpercent · 2024-08-02T23:43:26Z

@pytorchbot rebase -s

tenpercent · 2024-08-02T23:44:07Z

@pytorchbot label ciflow/trunk

pytorchmergebot · 2024-08-02T23:44:52Z

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

pytorchmergebot · 2024-08-02T23:44:53Z

Rebase failed due to Command git -C /home/runner/work/pytorch/pytorch rebase refs/remotes/origin/viable/strict pull/130576/head returned non-zero exit code 1

Rebasing (1/37)
Rebasing (2/37)
Rebasing (3/37)
Auto-merging torch/_inductor/config.py
CONFLICT (content): Merge conflict in torch/_inductor/config.py
Auto-merging torch/_inductor/utils.py
error: could not apply 3ba23799f2... rename supported_arch to ck_supported_arch in rocm config
hint: Resolve all conflicts manually, mark them as resolved with
hint: "git add/rm <conflicted_files>", then run "git rebase --continue".
hint: You can instead skip this commit: run "git rebase --skip".
hint: To abort and get back to the state before "git rebase", run "git rebase --abort".
hint: Disable this message with "git config advice.mergeConflict false"
Could not apply 3ba23799f2... rename supported_arch to ck_supported_arch in rocm config

Raised by https://github.com/pytorch/pytorch/actions/runs/10223394805

tenpercent · 2024-08-05T17:40:40Z

@pytorchbot merge

pytorchmergebot · 2024-08-05T17:43:33Z

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

pytorch-bot bot added module: inductor module: rocm AMD GPU support for Pytorch labels Jul 11, 2024

pytorchbot added the open source label Jul 11, 2024

tenpercent force-pushed the ck-addmm branch 2 times, most recently from f9b428e to 1b6db4c Compare July 15, 2024 19:45

tenpercent marked this pull request as ready for review July 15, 2024 21:36

colesbury requested a review from jeffdaily July 16, 2024 00:31

colesbury added the triaged This issue has been looked at a team member, and triaged and prioritized into an appropriate module label Jul 16, 2024

aakhundov added the topic: not user facing topic category label Jul 24, 2024

aakhundov reviewed Jul 24, 2024

View reviewed changes

zoranzhao reviewed Jul 25, 2024

View reviewed changes

torch/_inductor/kernel/mm.py Outdated Show resolved Hide resolved

torch/_inductor/codegen/rocm/ck_addmm_template.py Outdated Show resolved Hide resolved

pytorchmergebot force-pushed the ck-addmm branch from eb5da5f to 19f186e Compare July 25, 2024 18:29

pytorch-bot bot added ciflow/inductor ciflow/rocm Trigger "default" config CI on ROCm labels Jul 25, 2024

tenpercent added 10 commits July 25, 2024 19:47

try adding addmm

46c6937

merge addmm with matmul

c90d266

rename supported_arch to ck_supported_arch in rocm config

3ba2379

run lintrunner

c04335a

disable ck templplate for dynamic shapes

08ba57f

maybe fix type lint

f986090

cleanup tuple/sequence rendering

caaaf3b

add UT for autotuning matmul in same process

ccc19f6

refactor bias is not none check

3767c42

remove unused imports

481f354

tenpercent force-pushed the ck-addmm branch from e428fed to 481f354 Compare July 25, 2024 19:47

tenpercent added 7 commits July 30, 2024 22:59

Merge branch 'ck-addmm' of https://github.com/tenpercent/pytorch into…

cdfb09a

… ck-addmm

add a comment about numerics

9e28435

fix bias accumulation dtype

3d7820f

lint

00a84f8

fix bias broadcasted along the last dim

76f3775

lint

10f7677

tweak k in addmm unit test

8471d5e

tenpercent requested a review from chenyang78 August 1, 2024 20:45

update comment

beea4b3

chenyang78 reviewed Aug 2, 2024

View reviewed changes

tenpercent added 3 commits August 2, 2024 18:13

add comments

1b1ef1a

rename and simplify ck strides to leading dims

8cdf92e

update config

e2d6cba

chenyang78 approved these changes Aug 2, 2024

View reviewed changes

chenyang78 reviewed Aug 2, 2024

View reviewed changes

lint

d6fc230

pytorch-bot bot added the ciflow/trunk Trigger trunk jobs on your pull request label Aug 2, 2024

Merge remote-tracking branch 'upstream/viable/strict' into ck-addmm

7139093

pytorchmergebot added the merging label Aug 5, 2024

tenpercent mentioned this pull request Aug 5, 2024

CK companion PR to torch #130576 ROCm/composable_kernel#1441

Merged

pytorchmergebot added the Merged label Aug 5, 2024

pytorchmergebot closed this in 1c7dc33 Aug 5, 2024

pytorchmergebot removed the merging label Aug 5, 2024

[ROCm][CK][Inductor] Enable addmm for CK backend to gemm max autotune #130576

[ROCm][CK][Inductor] Enable addmm for CK backend to gemm max autotune #130576

Uh oh!

Conversation

tenpercent commented Jul 11, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Implementation details

Testing

Uh oh!

pytorch-bot bot commented Jul 11, 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/130576

✅ You can merge normally! (2 Unrelated Failures)

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

aakhundov commented Jul 24, 2024

Uh oh!

zoranzhao left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

tenpercent commented Jul 25, 2024

Uh oh!

pytorchmergebot commented Jul 25, 2024

Uh oh!

pytorchmergebot commented Jul 25, 2024

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

chenyang78 left a comment

Choose a reason for hiding this comment

Uh oh!

tenpercent commented Aug 2, 2024

Uh oh!

tenpercent commented Aug 2, 2024

Uh oh!

pytorchmergebot commented Aug 2, 2024

Uh oh!

pytorchmergebot commented Aug 2, 2024

Uh oh!

tenpercent commented Aug 5, 2024

Uh oh!

pytorchmergebot commented Aug 5, 2024

Merge started

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants

tenpercent commented Jul 11, 2024 •

edited

Loading

pytorch-bot bot commented Jul 11, 2024 •

edited

Loading