Implement fresh_triton_cache fixture through compilation.always_compile
#4363
Conversation
This file contains hidden or 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
…mpile' Signed-off-by: Anatoly Myachev <[email protected]>
Signed-off-by: Anatoly Myachev <[email protected]>
This was
linked to
issues
May 29, 2025
pbchekin
approved these changes
May 29, 2025
Contributor
Author
|
BTW: the idea came to my mind after talking to you, thanks @pbchekin! |
david-hls
pushed a commit
to david-hls/intel-xpu-backend-for-triton
that referenced
this pull request
Jun 18, 2025
The core Triton is a small number of people, and we receive many PRs (thank you!). To help us review your code more quickly, **if you are a new contributor (less than 3 PRs merged) we ask that you complete the following tasks and include the filled-out checklist in your PR description.** -------- This PR aims to address triton-lang/triton#4248, to correctly `device_print` the value if it is an signed integer. The signness info is lost when lowering TTGIR to LLIR (e.g. `i32` is always **signless** in MLIR), but the **lowered** data type is currently being used for constructing the format specifier in the `PrintOpToLLVM` implementation (triton-lang/triton#4248 (comment)), so a negative value is printed out as an unsigned int, thus confusing users. A minimal reproducer is ```python import torch import triton import triton.language as tl @triton.jit def print_kernel(ptr): value = tl.load(ptr) tl.device_print("value in kernel from device_print", value) print_kernel[(1,)](torch.tensor(10, dtype=torch.int32).cuda()) print_kernel[(1,)](torch.tensor(-10, dtype=torch.int32).cuda()) print_kernel[(1,)](torch.tensor((1 << 31) + 1000, dtype=torch.uint32).cuda()) ``` Currently, it prints ``` pid (0, 0, 0) idx () value in kernel from device_print: 10 ... pid (0, 0, 0) idx () value in kernel from device_print: 4294967286 ... pid (0, 0, 0) idx () value in kernel from device_print: 2147484648 ``` (always as unsigned int) This PR adds extra `isSigned` attribute in the `PrintOp` to indicate if each operand in the `PrintOp` should be printed as signed or not. With this, the program above now prints correctly ``` pid (0, 0, 0) idx () value in kernel from device_print: 10 ... pid (0, 0, 0) idx () value in kernel from device_print: -10 ... pid (0, 0, 0) idx () value in kernel from device_print: 2147484648 ``` Extra LIT tests and python unit tests are added as well; also manually verified that they failed without the fix and passing now by running ``` $ pytest python/test/unit/language/test_subprocess.py $ cd python/build/cmake.linux-x86_64-cpython-3.10; lit test ``` **Alternative considered**: adds `uint32` in the triton MLIR data type definition and then rely on the triton op data type to determine the format specifier, to retain the original signness info; as in commit triton-lang/triton@f7a7407. However, as PR reviewer pointed out, that means adding a new data type in the Triton IR just for this purpose, which is overkill and introduces unnecessary maintenance overhead and thus less ideal. -------- Complete the following tasks before sending your PR, and replace `[ ]` with `[x]` to indicate you have done them. - [X] I am not making a trivial change, such as fixing a typo in a comment. - [X] I have written a PR description following these [rules](https://cbea.ms/git-commit/#why-not-how). - [X] I have run `pre-commit run --from-ref origin/main --to-ref HEAD`. - Select one of the following. - [X] I have added tests. - `/test` for `lit` tests - `/unittest` for C++ tests - `/python/test` for end-to-end tests - [ ] This PR does not need a test because `FILL THIS IN`. - Select one of the following. - [ ] I have not added any `lit` tests. - [X] The `lit` tests I have added follow these [best practices](https://mlir.llvm.org/getting_started/TestingGuide/#filecheck-best-practices), including the "tests should be minimal" section. (Usually running Python code and using the instructions it generates is not minimal.)
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.
The main idea is that for tests where we need the cache to be generated each time, we can set this behavior via a special environment variable (
TRITON_ALWAYS_COMPILEor corresponding knobs), instead of creating a new temporary folder and deleting it each time, which is problematic on Windows. It seems like this solution is more lightweight, maybe it will be possible to upstream it.BMG+WIN CI: