[inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion #126068

jgong5 · 2024-05-13T14:21:37Z

Stack from ghstack (oldest at bottom):

As part of #125683, this PR adds the initial bf16/fp16 gemm template support with micro-gemm implemented with fused type casting and fp32 computation. It doesn't provide epilogue fusion support yet which will be added in the next PR.

cc @voznesenskym @penguinwu @EikanWang @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @peterbell10 @ipiszy @yf225 @chenyang78 @kadeng @muchulee8 @ColinPeppler @amjames @desertfire @chauhang

Differential Revision: D58017580

[ghstack-poisoned]

pytorch-bot · 2024-05-13T14:21:41Z

🔗 Helpful Links

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

📄 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.

❌ 21 Cancelled Jobs, 70 Unrelated Failures

As of commit cd4decd with merge base 94fea82 ():

CANCELLED JOBS - The following jobs were cancelled. Please retry:

inductor / linux-jammy-cpu-py3.8-gcc11-inductor / test (cpu_inductor_torchbench_freezing, 1, 2, linux.12xlarge) (gh)
##[error]The operation was canceled.
inductor / linux-jammy-cpu-py3.8-gcc11-inductor / test (cpu_inductor_torchbench, 1, 2, linux.12xlarge) (gh)
##[error]The operation was canceled.
inductor / rocm6.1-py3.8-inductor / test (inductor, 1, 1, linux.rocm.gpu.2) (gh)
##[error]The operation was canceled.
linux-binary-libtorch-cxx11-abi / libtorch-cpu-shared-with-deps-cxx11-abi-test / test (gh)
undefined reference to c10::detail::torchInternalAssertFail(char const*, char const*, unsigned int, char const*, std::string const&)'`
linux-binary-manywheel / manywheel-py3_8-cuda11_8-build / build (gh)
##[error]The operation was canceled.
linux-binary-manywheel / manywheel-py3_8-cuda12_1-build / build (gh)
##[error]The operation was canceled.
linux-binary-manywheel / manywheel-py3_8-cuda12_4-build / build (gh)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9-bazel-test / build-and-test (default, 1, 1, linux.4xlarge.nvidia.gpu) (gh)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9-sm86 / test (default, 3, 5, linux.g5.4xlarge.nvidia.gpu) (gh)
##[error]The operation was canceled.
pull / linux-focal-cuda12.4-py3.10-gcc9-bazel-test / build-and-test (default, 1, 1, linux.4xlarge.nvidia.gpu) (gh)
##[error]The operation was canceled.
pull / linux-focal-py3_8-clang9-xla / test (xla, 1, 1, linux.12xlarge) (gh)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (default, 1, 3, linux.2xlarge) (gh)
##[error]The operation was canceled.
pull / linux-focal-py3.12-clang10 / test (default, 1, 3, linux.2xlarge) (gh)
##[error]The operation was canceled.
pull / linux-focal-py3.8-clang10 / test (default, 1, 3, linux.2xlarge) (gh)
##[error]The operation was canceled.
pull / linux-jammy-py3.8-gcc11 / test (default, 1, 3, linux.2xlarge) (gh)
##[error]The operation was canceled.
pull / linux-jammy-py3.8-gcc11 / test (default, 2, 3, linux.2xlarge) (gh)
##[error]The operation was canceled.
trunk / linux-focal-rocm6.1-py3.8 / test (default, 1, 2, linux.rocm.gpu) (gh)
##[error]The operation was canceled.
trunk / linux-focal-rocm6.1-py3.8 / test (default, 2, 2, linux.rocm.gpu) (gh)
##[error]The operation was canceled.
trunk / linux-focal-rocm6.1-py3.8 / test (distributed, 1, 1, linux.rocm.gpu) (gh)
##[error]The operation was canceled.
trunk / macos-py3-arm64 / test (default, 1, 3, macos-m1-stable) (gh)
##[error]The operation was canceled.
trunk / pytorch-linux-focal-py3-clang9-android-ndk-r21e-build / build (default, 1, 1, linux.2xlarge) (gh)
##[error]The operation was canceled.

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

pull / linux-focal-cuda12.1-py3.10-gcc9 / test (default, 1, 5, linux.4xlarge.nvidia.gpu) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9 / test (default, 2, 5, linux.4xlarge.nvidia.gpu) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9 / test (default, 3, 5, linux.4xlarge.nvidia.gpu) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9 / test (default, 4, 5, linux.4xlarge.nvidia.gpu) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-cuda12.1-py3.10-gcc9 / test (default, 5, 5, linux.4xlarge.nvidia.gpu) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (default, 2, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (default, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (dynamo, 1, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (dynamo, 2, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.11-clang10 / test (dynamo, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.12-clang10 / test (default, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.12-clang10 / test (dynamo, 1, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.12-clang10 / test (dynamo, 2, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.12-clang10 / test (dynamo, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.8-clang10 / test (default, 2, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.8-clang10 / test (default, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.8-clang10 / test (dynamo, 1, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-focal-py3.8-clang10 / test (dynamo, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-jammy-py3.10-clang15-asan / test (default, 2, 6, linux.4xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-jammy-py3.8-gcc11 / test (default, 3, 3, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-jammy-py3.8-gcc11 / test (distributed, 1, 2, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.
pull / linux-jammy-py3.8-gcc11 / test (distributed, 2, 2, linux.2xlarge) (gh) (similar failure)
##[error]The operation was canceled.

UNSTABLE - The following jobs failed but were likely due to flakiness present on trunk and has been marked as unstable:

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

ghstack-source-id: 548f560 Pull Request resolved: #126068

[ghstack-poisoned]

ghstack-source-id: ffd469e Pull Request resolved: #126068

[ghstack-poisoned]

… w/o epilogue fusion" This PR adds the initial bf16/fp16 gemm template support with micro-gemm implemented with fused type casting and fp32 computation. It doesn't provide epilogue fusion support yet which will be added in the next PR. cc voznesenskym penguinwu EikanWang Guobing-Chen XiaobingSuper zhuhaozhe blzheng wenzhe-nrv jiayisunx peterbell10 ipiszy yf225 chenyang78 kadeng muchulee8 ColinPeppler amjames desertfire chauhang [ghstack-poisoned]

ghstack-source-id: 6a954b0 Pull Request resolved: #126068

… w/o epilogue fusion" As part of #125683, this PR adds the initial bf16/fp16 gemm template support with micro-gemm implemented with fused type casting and fp32 computation. It doesn't provide epilogue fusion support yet which will be added in the next PR. cc voznesenskym penguinwu EikanWang Guobing-Chen XiaobingSuper zhuhaozhe blzheng wenzhe-nrv jiayisunx peterbell10 ipiszy yf225 chenyang78 kadeng muchulee8 ColinPeppler amjames desertfire chauhang Differential Revision: [D58017580](https://our.internmc.facebook.com/intern/diff/D58017580) [ghstack-poisoned]

[ghstack-poisoned]

desertfire · 2024-06-03T14:17:29Z

@pytorchbot revert -m "failing internal tests" -c "ghfirst"

The cpu_select_algorithm tests you added/changed seem to be failing internally

Example of 1 error: pastebin.com/BB8E1Ngx

Generally the succinct lines are in the form of

222/torch/include/ATen/cpu/vec/vec_half.h:19:10: error: '__builtin_ia32_vcvtps2ph' needs target feature f16c

616/torch/include/ATen/cpu/vec/vec256/vec256_bfloat16.h:107:7: error: always_inline function '_mm256_cvtph_ps' requires target feature 'f16c', but would be inlined into function 'cvtfp16_fp32' that is compiled without support for 'f16c'

34/torch/include/ATen/cpu/vec/vec_half.h:37:10: error: always_inline function '_cvtsh_ss' requires target feature 'f16c', but would be inlined into function 'half2float_scalar' that is compiled without support for 'f16c'

etc.

@PaliC , fbcode and github are out of sync on this one. This PR is reverted but D58015598 isn't. Please help to solve the problem.

desertfire · 2024-06-04T02:41:00Z

@desertfire has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator.

[ghstack-poisoned]

jgong5 · 2024-06-05T11:36:10Z

@pytorchbot merge

pytorchmergebot · 2024-06-05T11:37:59Z

Merge failed

Reason: This PR has internal changes and must be landed via Phabricator

Details for Dev Infra team

Raised by workflow job

…o epilogue fusion (pytorch#126068)" This reverts commit dae33a4. Reverted pytorch#126068 on behalf of https://github.com/PaliC due to failing internal tests ([comment](pytorch#126068 (comment)))

[ghstack-poisoned]

jgong5 · 2024-06-07T08:17:53Z

@pytorchbot merge

pytorchmergebot · 2024-06-07T08:20:33Z

Merge failed

Reason: This PR has internal changes and must be landed via Phabricator

Details for Dev Infra team

Raised by workflow job

[ghstack-poisoned]

jgong5 · 2024-06-08T02:19:08Z

@pytorchbot merge

pytorchmergebot · 2024-06-08T02:21:15Z

Merge failed

Reason: This PR has internal changes and must be landed via Phabricator

Details for Dev Infra team

Raised by workflow job

[ghstack-poisoned]

…128472) reland for #126068 Pull Request resolved: #128472 Approved by: https://github.com/desertfire

…ytorch#128472) reland for pytorch#126068 Pull Request resolved: pytorch#128472 Approved by: https://github.com/desertfire

Update

f74cec9

[ghstack-poisoned]

jgong5 mentioned this pull request May 13, 2024

[inductor][cpp] epilogue support for gemm template #126019

Closed

pytorch-bot bot added ciflow/inductor module: inductor labels May 13, 2024

jgong5 pushed a commit that referenced this pull request May 13, 2024

[inductor][cpp] bf16/fp16 gemm template computed with fp32

a2b041d

ghstack-source-id: 548f560 Pull Request resolved: #126068

jgong5 marked this pull request as draft May 13, 2024 14:22

pytorchbot added the open source label May 13, 2024

Update

0dc98ee

[ghstack-poisoned]

jgong5 changed the title ~~[inductor][cpp] bf16/fp16 gemm template computed with fp32~~ [inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion May 14, 2024

Update

5819c0d

[ghstack-poisoned]

jgong5 pushed a commit that referenced this pull request May 14, 2024

[inductor][cpp] bf16/fp16 gemm template computed with fp32

0497a6b

ghstack-source-id: ffd469e Pull Request resolved: #126068

Update

5a59d0e

[ghstack-poisoned]

jgong5 marked this pull request as ready for review May 14, 2024 09:16

jgong5 requested review from jansel and lezcano and removed request for jansel May 14, 2024 12:16

jgong5 pushed a commit that referenced this pull request May 14, 2024

[inductor][cpp] bf16/fp16 gemm template computed with fp32

24ccda6

ghstack-source-id: 6a954b0 Pull Request resolved: #126068

jgong5 mentioned this pull request Jun 1, 2024

[inductor][cpp] align dtype of local buffer with the global one #127687

Closed

Update

588b138

[ghstack-poisoned]

desertfire mentioned this pull request Jun 3, 2024

[inductor][cpp] bf16/fp16 gemm template computed with fp32 #127765

Closed

Update

e2a84cd

[ghstack-poisoned]

pytorchmergebot added the merging label Jun 5, 2024

pytorchmergebot removed the merging label Jun 5, 2024

Update

752cacb

[ghstack-poisoned]

pytorchmergebot added the merging label Jun 7, 2024

pytorchmergebot removed the merging label Jun 7, 2024

Update

91513fd

[ghstack-poisoned]

pytorchmergebot added the merging label Jun 8, 2024

pytorchmergebot removed the merging label Jun 8, 2024

Jiong Gong added 2 commits June 8, 2024 01:27

Update

2d21c2f

[ghstack-poisoned]

Update

cd4decd

[ghstack-poisoned]

jgong5 mentioned this pull request Jun 12, 2024

[RELAND][inductor][cpp] bf16/fp16 gemm template computed with fp32 #128472

Closed

jgong5 closed this Jun 12, 2024

github-actions bot deleted the gh/jgong5/46/head branch July 12, 2024 01:58

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion #126068

[inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion #126068

Uh oh!

jgong5 commented May 13, 2024 •

edited

Loading

Uh oh!

pytorch-bot bot commented May 13, 2024 •

edited

Loading

Uh oh!

desertfire commented Jun 3, 2024

Uh oh!

desertfire commented Jun 4, 2024

Uh oh!

jgong5 commented Jun 5, 2024

Uh oh!

pytorchmergebot commented Jun 5, 2024

Uh oh!

jgong5 commented Jun 7, 2024

Uh oh!

pytorchmergebot commented Jun 7, 2024

Uh oh!

jgong5 commented Jun 8, 2024

Uh oh!

pytorchmergebot commented Jun 8, 2024

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

[inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion #126068

[inductor][cpp] bf16/fp16 gemm template computed with fp32 w/o epilogue fusion #126068

Uh oh!

Conversation

jgong5 commented May 13, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pytorch-bot bot commented May 13, 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/126068

❌ 21 Cancelled Jobs, 70 Unrelated Failures

Uh oh!

desertfire commented Jun 3, 2024

Uh oh!

desertfire commented Jun 4, 2024

Uh oh!

jgong5 commented Jun 5, 2024

Uh oh!

pytorchmergebot commented Jun 5, 2024

Merge failed

Uh oh!

jgong5 commented Jun 7, 2024

Uh oh!

pytorchmergebot commented Jun 7, 2024

Merge failed

Uh oh!

jgong5 commented Jun 8, 2024

Uh oh!

pytorchmergebot commented Jun 8, 2024

Merge failed

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

jgong5 commented May 13, 2024 •

edited

Loading

pytorch-bot bot commented May 13, 2024 •

edited

Loading