[Frontend] Disallow access to globals that are not tl.constexpr. #3762
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
jlebar
force-pushed
the
dev-jlebar/disallow-non-constexpr-global
branch
from
3192153
to
21336e5
Compare
We've found that users often accidentally access global vars from inside a Triton kernel that they don't mean to. Moreover, they sometimes *change* these global vars after they first run the kernel, but Triton does not notice this (the initial value of the global is baked into the compiled kernel). As a result, the following runs have a stale value for the global. We hope that making users annotate their globals with tl.constexpr will remind them that they are kernel arguments that they should not change. In a separate PR, we will additionally explore the possibility of raising an error when the value of one of these globals changes. This is a breaking change, so we allow you to opt out of it by setting TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1. But we do not promise to support this envvar forever. GPC: disallow-non-constexpr-global
jlebar
force-pushed
the
dev-jlebar/disallow-non-constexpr-global
branch
from
21336e5
to
22e975b
Compare
ThomasRaoux
reviewed
Apr 25, 2024
ptillet
approved these changes
Apr 25, 2024
amjames
added a commit
to amjames/pytorch
that referenced
this pull request
May 16, 2024
…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
pytorchmergebot
pushed a commit
to pytorch/pytorch
that referenced
this pull request
Jun 6, 2024
… globals correctly. (#126195) [Triton #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: #126195 Approved by: https://github.com/alexbaden, https://github.com/bertmaher
TharinduRusira
pushed a commit
to TharinduRusira/pytorch
that referenced
this pull request
Jun 14, 2024
… 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
[Frontend] Disallow access to globals that are not tl.constexpr.
We've found that users often accidentally access global vars from inside a
Triton kernel that they don't mean to. Moreover, they sometimes change these
global vars after they first run the kernel, but Triton does not notice this
(the initial value of the global is baked into the compiled kernel). As a
result, the following runs have a stale value for the global.
We hope that making users annotate their globals with tl.constexpr will remind
them that they are kernel arguments that they should not change.
In a separate PR, we will additionally explore the possibility of raising an
error when the value of one of these globals changes.
This is a breaking change, so we allow you to opt out of it by setting
TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1. But we do not promise to support this
envvar forever.
PR chain