Poc tle gpu tileir - #811
Open
macqueen09 wants to merge 188 commits into
Open
Conversation
Cherry Pick of triton-lang/triton#4138 to release/3.0.x branch Co-authored-by: Keren Zhou <kerenzhou@openai.com>
…patibility (#4160) cherry pick #4049 for release/3.0.x Co-authored-by: Alex Baden <alexbaden@users.noreply.github.com>
… it is not needed (#4161) These need to be in this order to apply cleanly. Cherry pick #3816 and #3790 for release/3.0.x. --------- Co-authored-by: Alexander Efimov <efimov.alexander@gmail.com> Co-authored-by: Lixun Zhang <Lixun.Zhang@amd.com>
These need to be in this order to apply cleanly Cherry pick #4000 and #4015 for release/3.0.x --------- Signed-off-by: Finlay Marno <finlay.marno@codeplay.com> Co-authored-by: Finlay <finlay.marno@codeplay.com> Co-authored-by: Manman Ren <mren@meta.com> Co-authored-by: Manman Ren <mren@fb.com>
Cherry pick #4086 for release/3.0.x Co-authored-by: Lixun Zhang <Lixun.Zhang@amd.com>
Cherry pick #4097 for release/3.0.x
Cherry picks for release/3.0.x General: - e8bc45d [BACKEND][AMD] Disable linear layout due to perf regression (#4126) - 9a0a7c2 [AMD] Add basic verification to MFMA encoding (#4117) for RDNA: - 100e2aa [AMD][WMMA] Support dot3d (#3674) - 4a1ea8e [AMD][gfx11] Fix BF16 wmma instr generation (#4135) Proton HIP PRs: - 328b86d [PROTON] Refactor GPU profilers (#4056) - 60613fb [PROTON] Roctracer: convert agent id to gpu id for gpu ops (#4090) - c1776fa [PROTON][AMD] Add Proton HIP GPU Utilization Metrics (#4119) --------- Co-authored-by: Lei Zhang <antiagainst@gmail.com> Co-authored-by: Alexander Efimov <efimov.alexander@gmail.com> Co-authored-by: Ilya V <152324710+joviliast@users.noreply.github.com> Co-authored-by: Keren Zhou <kerenzhou@openai.com> Co-authored-by: mwootton <michael.wootton@amd.com> Co-authored-by: Corbin Robeck <corbin.robeck@amd.com>
…1) (#4223) Cherry-picking triton-lang/triton#4061 for release/3.0.x
This commit change the main `README.md` only. I updated AMD GPUs to the officially support hardware. CC: @antiagainst
To incorporate latest changes and release 3.0.0-post1 Same as: triton-lang/triton@ea70d77
Cherry Pick of : triton-lang/triton#4330 triton-lang/triton#4335 to release/3.0.x branch --------- Co-authored-by: Nikita Shulga <2453524+malfet@users.noreply.github.com> Co-authored-by: Keren Zhou <kerenzhou@openai.com>
Cherry-pick of 4383d2835d8533bedacac95915b6b740c9fa40d8 and 5eb4ba518eb5b07084562ed584f336a926f6ace3 from 2.0.X to 3.0.X branch
…AMD gpus (#4569) * Cherry-pick 06e6799 to 3.0.x to fix build issue * Cherry-pick cf2ad02 to fix hang while running flex attention on AMD gpus * Fix build issue with cherry-pick of cf2ad02 * `enableRegionSimplification` was changed from boolen type to enum in mainline branch --------- Co-authored-by: Philippe Tillet <phil@openai.com> Co-authored-by: Giuseppe Rossini <giuseppe.rossini@amd.com>
…696) 1. Set build version to 3.1.0 2. Remove: triton-lang/triton#4354 change deployed to build ``wheel-1``
Cherry Pick of #4311 to release/3.1.x branch Co-authored-by: peterbell10 <peterbell10@live.co.uk>
Cherry Pick of #4247 to release/3.1.x branch Co-authored-by: Hongtao Yu <hoy@meta.com>
…_bench` (#4733) This PR cherry picks a device agnostic change to `do_bench` to use the `DeviceInterface` over the hard-coding of the device type onto the release/3.1.x branch: * triton-lang/triton@36789d2 This will allow for further device agnostic changes in PyTorch Inductor. --- 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.** 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). - [ ] I have run `pre-commit run --from-ref origin/main --to-ref HEAD`. - Select one of the following. - [ ] I have added tests. - `/test` for `lit` tests - `/unittest` for C++ tests - `/python/test` for end-to-end tests - [x] This PR does not need a test because it is a cherry-pick of a commit which was landed to `main`. - Select one of the following. - [x] I have not added any `lit` tests. - [ ] 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.) 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.** Co-authored-by: Artemiy Bulavin <43750406+abulavin@users.noreply.github.com>
Per @jataylo : * libamdhip64.so discovery: * triton-lang/triton@13edc45 * triton-lang/triton@e0613c6 * wmma: triton-lang/triton@4a1ea8e (already cherry-picked) * mfma: triton-lang/triton@0a66c1b --------- Co-authored-by: Xinya Zhang <Xinya.Zhang@amd.com> Co-authored-by: Lei Zhang <antiagainst@gmail.com> Co-authored-by: Lixun Zhang <Lixun.Zhang@amd.com>
The string representation allows PyTorch Inductor to serialize/derserialize the `AttrsDescriptor` to the `@triton.heuristics` block in the generated code.
…on (#5033) Reductions have special handling for side effectful "combine ops" (e.g. "add" for a sum reduction). In the presence of side effects, a predicate is computed to determine whether a thread should participate in the reduction, to ensure that invalid/uninitialized data is not operated on. See #4811 for more details. ~Previously, the predicate logic was incorrect for 2D reductions. This PR fixes the logic and adds a python test.~ Edit: after additional discussion with @peterbell10, we removed the lanePred logic. Here's our thinking on why this is valid: * lanePred info is computed based entirely on the blocked layout info and properties of the reduction * the blocked layout won't tell you which threads do or don't have uninitialized data Instead, it sounds like the motivation for #4811 is based on uninitialized values that can be indicated by the `pred` variable passed into `warpReduce()`.
…(#5075) This is a follow up to #5033 but for scan ops, and also improving the testing as it was clearly insufficient before.
support asserts with scalar condition and only emit barrier for assert of tensors. Thanks to @peterbell10 for the suggestion.
…5081) Context: in `CodeGenerator.__init__`, globals for a given triton function are modified to handle remapping the libdevice module to cuda or hip (from triton-lang/triton#4539). In particular, this logic: ```python for k, v in gscope.items(): # gscope is a dict of fn.__globals__ ... self.gscope[k] = getattr(module_map[module_name], k) ``` was failing if you do this in the global scope: `from triton.language.extras.libdevice import fast_dividef as my_fast_dividef`.
…le (#5079) ExternElementwise ops have a `pure` attribute that marks the op as pure. If an op is pure, it should also be speculatable. In the reduction/scan ttgir->llvm passes, checks for speculatability are failing for ExternElementwise ops, causing additional conditional handling to be added. This PR makes ExternElementwise ops implement ConditionallySpeculatable, and mark the op as speculatable if the op is marked as pure. This removes the conditional branches from the generated scan/reduction code.
…rgs (#5083) Currently, when there are hooks accessing args in the `@triton.autotuner` (e.g., `pre_hook` / `post_hook` or those created from the `restore_value` or `reset_to_zero`) and the kernel arguments are passed as kwargs, this breaks, because the hooks take positional `args` only. This PR changes the first parameter of the `pre_hook` and `post_hook` of the `Autotuner` from (partial) `args` to (full) `kwargs`. As a result, we now have access to all arguments, positional or keyword, passed to the kernel call. The call sites and docs are updated accordingly. N.B.: This change is BC-breaking! In the signatures of `pre_hook` and `post_hook`, the first parameter type has changed from list (of positional args) to dict (of all kwargs). Submitting, as agreed with @Jokeren and @peterbell10 below. Fixes #5082. See the failing example, error, and code pointers there.
* fix CANN enum name and numpy import * add solution for some issue --------- Co-authored-by: zhongsanming <sanming.zhong@compiler-dev.com>
* feat(test): add matmul raw * feat(test): add test and dump * feat(test): use one loop * feat(test): use prologue * feat(test): remove unused variables * fix(test): use mat_c_l0c * feat(test): use sub_view * fix(test): preload 2 buffer * fix(test): fix gcd * fix(test): add out-of-bounds check
* feat(test): add matmul_add_residual * feat(test): add matmul_add_residual_cv * fix(test): remove mat_c_l0c * feat(test): add matmul_prefetch * convert tile.subview to memref.subview (#16) Co-authored-by: zhongsanming <zhongsanming@compiler-dev.com> --------- Co-authored-by: zhongsanming <116532846+zhongsanming@users.noreply.github.com> Co-authored-by: zhongsanming <zhongsanming@compiler-dev.com>
…idual-add fix native matmul add residual
* feat(test): add native_fa * feat(test): add native_fa * feat(test): use fused-attention * feat(test): use tile_allc/copy
* fix subview lowering issue, now it runs but produces wrong result * use MACRO to simplify pattern rewrite * allow pass num cores from cli to kernel impl --------- Co-authored-by: zhongsanming <sanming.zhong@compiler-dev.com>
Co-authored-by: zhongsanming <sanming.zhong@compiler-dev.com>
Co-authored-by: zhongsanming <sanming.zhong@compiler-dev.com>
Co-authored-by: macqueen09 <macqueen09@users.noreply.github.com>
macqueen09
requested review from
Galaxy1458,
i3wanna2,
menchunlei,
sunnycase and
zhzhcookie
as code owners
July 17, 2026 09:48
|
|
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
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.
主要功能: