[inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. #126195

amjames · 2024-05-14T18:27:03Z

Stack from ghstack (oldest at bottom):

Triton #3762
disallows access to globals which are not tl.constexpr

Triton has always treated captured globals this way, but they now
require it be explicit in user code.

Updated codegen to make sure these variables are defined before writing
the kernel source when compiling a user defined triton kernel.

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

[ghstack-poisoned]

pytorch-bot · 2024-05-14T18:27:07Z

🔗 Helpful Links

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

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

❌ 1 New Failure, 6 Unrelated Failures

As of commit 15e5551 with merge base ade0754 ():

NEW FAILURE - The following job has failed:

trunk / macos-13-py3-arm64 / build (gh)
/Users/ec2-user/runner/_work/pytorch/pytorch/c10/util/StringUtil.cpp:45:8: error: 'wstring_convert<std::codecvt_utf8_utf16<wchar_t>>' is deprecated [-Werror,-Wdeprecated-declarations]

FLAKY - The following job failed but was likely due to flakiness present on trunk:

linux-binary-manywheel / manywheel-py3_8-cuda11_8-test / test (gh) (similar failure)
ImportError: libcudnn.so.9: cannot open shared object file: No such file or directory

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

inductor / cuda12.1-py3.10-gcc9-sm86 / test (dynamic_inductor_timm, 2, 2, linux.g5.4xlarge.nvidia.gpu, unstable) (gh) (#127438)
sebotnet33ts_256
inductor / linux-jammy-cpu-py3.8-gcc11-inductor / test (inductor_torchbench_cpu_smoketest_perf, 1, 1, linux.24xl.spr-metal, unstable) (gh) (#126993)
Process completed with exit code 1.
linux-binary-manywheel / manywheel-py3_8-cuda12_1-test / test (gh) (#127288)
ImportError: libcudnn.so.9: cannot open shared object file: No such file or directory
linux-binary-manywheel / manywheel-py3_8-cuda12_4-test / test (gh) (#127289)
ImportError: libcudnn.so.9: cannot open shared object file: No such file or directory
pull / linux-focal-cuda12.4-py3.10-gcc9-sm86 / test (default, 5, 5, linux.g5.4xlarge.nvidia.gpu, unstable) (gh) ()
inductor/test_efficient_conv_bn_eval.py::EfficientConvBNEvalCudaTests::test_basic_cuda

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

amjames · 2024-05-14T18:32:34Z

This is needed after moving the pin in #126098 I split it out since it does not seem to break anything with the current pin.

[ghstack-poisoned]

torch/_inductor/codegen/wrapper.py

…ser defined kernel source [Triton pytorch#3762](triton-lang/triton#3762) disallows access to globals which are not `tl.constexpr` Triton has always treated captured globals this way, but they now require it be explicit in user code. Updated codegen to make sure these variables are defined before writing the kernel source when compiling a user defined triton kernel. ghstack-source-id: 5fc2a44f6d39a699d704ca1f311eb84770af8647 Pull Request resolved: pytorch#126195

[ghstack-poisoned]

alexbaden

I tested this latest version locally with the CUDA backend and XPU backends and both now pass the inductor test_triton_kernels/test_triton_kernel_constants test with a recent Triton/TritonXPU commit.

test/inductor/test_triton_kernels.py

torch/_inductor/codegen/wrapper.py

[ghstack-poisoned]

torch/_inductor/codegen/wrapper.py

[ghstack-poisoned]

bertmaher · 2024-06-04T19:53:20Z

@pytorchbot merge

pytorchmergebot · 2024-06-04T19:55:12Z

Merge failed

Reason: This PR needs a release notes: label
If your changes are user facing and intended to be a part of release notes, please use a label starting with release notes:.

If not, please add the topic: not user facing label.

To add a label, you can comment to pytorchbot, for example
@pytorchbot label "topic: not user facing"

For more information, see
https://github.com/pytorch/pytorch/wiki/PyTorch-AutoLabel-Bot#why-categorize-for-release-notes-and-how-does-it-work.

Details for Dev Infra team

Raised by workflow job

bertmaher · 2024-06-04T20:03:15Z

@pytorchbot merge

pytorchmergebot · 2024-06-04T20:05:31Z

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

pytorchmergebot · 2024-06-04T20:05:53Z

Merge failed

Reason: 1 jobs have failed, first few of them are: trunk / macos-13-py3-arm64 / build

Details for Dev Infra team

Raised by workflow job

bertmaher · 2024-06-06T20:48:06Z

@pytorchbot merge -f "failure looks unrelated, I don't think triton can be involved with the macOS build"

pytorchmergebot · 2024-06-06T20:49:59Z

Merge started

Your change will be merged immediately since you used the force (-f) flag, bypassing any CI checks (ETA: 1-5 minutes). Please use -f as last resort and instead consider -i/--ignore-current to continue the merge ignoring current failures. This will allow currently pending tests to finish and report signal before the merge.

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

… globals correctly. (pytorch#126195) [Triton pytorch#3762](triton-lang/triton#3762) disallows access to globals which are not `tl.constexpr` Triton has always treated captured globals this way, but they now require it be explicit in user code. Updated codegen to make sure these variables are defined before writing the kernel source when compiling a user defined triton kernel. Pull Request resolved: pytorch#126195 Approved by: https://github.com/alexbaden, https://github.com/bertmaher

Update

1162c27

[ghstack-poisoned]

amjames mentioned this pull request May 14, 2024

Update triton pin #126098

Closed

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

pytorchbot added the open source label May 14, 2024

Update

092de78

[ghstack-poisoned]

alexbaden mentioned this pull request May 15, 2024

[PyTorch Upstream] Triton regression on Inductor test case:NameError("Cannot access global variable STRING_CONSTANT_C from within @jit'ed function) intel/intel-xpu-backend-for-triton#1126

Closed

alexbaden reviewed May 15, 2024

View reviewed changes

torch/_inductor/codegen/wrapper.py Outdated Show resolved Hide resolved

amjames mentioned this pull request May 15, 2024

Inconsistency between constants as arguments and captured globals triton-lang/triton#3924

Open

Update

48f4000

[ghstack-poisoned]

amjames mentioned this pull request May 16, 2024

[inductor] Don't inherit __future__ flags from the calling scope when compile -ing generated modules #126454

Closed

Update

617b298

[ghstack-poisoned]

amjames changed the title ~~[inductor][codegen] Define constexpr globals with annotation before user defined kernel source~~ [inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. May 16, 2024

Update

e1e4f1c

[ghstack-poisoned]

amjames requested a review from peterbell10 May 17, 2024 12:46

alexbaden approved these changes May 17, 2024

View reviewed changes

test/inductor/test_triton_kernels.py Outdated Show resolved Hide resolved

peterbell10 reviewed May 17, 2024

View reviewed changes

test/inductor/test_triton_kernels.py Outdated Show resolved Hide resolved

torch/_inductor/codegen/wrapper.py Outdated Show resolved Hide resolved

amjames added 4 commits May 21, 2024 13:03

Update

80599b5

[ghstack-poisoned]

Update

deb582f

[ghstack-poisoned]

Update

f179ae6

[ghstack-poisoned]

Update

bc0c488

[ghstack-poisoned]

peterbell10 reviewed May 29, 2024

View reviewed changes

torch/_inductor/codegen/wrapper.py Outdated Show resolved Hide resolved

Update

15e5551

[ghstack-poisoned]

amjames mentioned this pull request May 29, 2024

Update triton pin to 513f38c4 #127429

Closed

bertmaher approved these changes Jun 4, 2024

View reviewed changes

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

pytorchmergebot added the merging label Jun 4, 2024

pytorchmergebot removed the merging label Jun 4, 2024

bertmaher added the topic: not user facing topic category label Jun 4, 2024

pytorchmergebot added the merging label Jun 4, 2024

pytorchmergebot removed the merging label Jun 4, 2024

pytorchmergebot added the merging label Jun 6, 2024

pytorchmergebot closed this in 9554300 Jun 6, 2024

pytorchmergebot added Merged and removed merging labels Jun 6, 2024

github-actions bot deleted the gh/amjames/22/head branch July 7, 2024 01:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. #126195

[inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. #126195

amjames commented May 14, 2024 •

edited

Loading

pytorch-bot bot commented May 14, 2024 •

edited

Loading

amjames commented May 14, 2024

alexbaden left a comment

bertmaher commented Jun 4, 2024

pytorchmergebot commented Jun 4, 2024

bertmaher commented Jun 4, 2024

pytorchmergebot commented Jun 4, 2024

pytorchmergebot commented Jun 4, 2024

bertmaher commented Jun 6, 2024

pytorchmergebot commented Jun 6, 2024

[inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. #126195

[inductor][codegen] Codegen constexpr globals and constexpr annotated globals correctly. #126195

Conversation

amjames commented May 14, 2024 • edited Loading

pytorch-bot bot commented May 14, 2024 • edited Loading

🔗 Helpful Links

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

❌ 1 New Failure, 6 Unrelated Failures

amjames commented May 14, 2024

alexbaden left a comment

Choose a reason for hiding this comment

bertmaher commented Jun 4, 2024

pytorchmergebot commented Jun 4, 2024

Merge failed

bertmaher commented Jun 4, 2024

pytorchmergebot commented Jun 4, 2024

Merge started

pytorchmergebot commented Jun 4, 2024

Merge failed

bertmaher commented Jun 6, 2024

pytorchmergebot commented Jun 6, 2024

Merge started

amjames commented May 14, 2024 •

edited

Loading

pytorch-bot bot commented May 14, 2024 •

edited

Loading