Skip to content

Poc tle gpu tileir - #811

Open
macqueen09 wants to merge 188 commits into
flagos-ai:mainfrom
KernelLLM:poc-tle-gpu-tileir
Open

Poc tle gpu tileir#811
macqueen09 wants to merge 188 commits into
flagos-ai:mainfrom
KernelLLM:poc-tle-gpu-tileir

Conversation

@macqueen09

Copy link
Copy Markdown

主要功能:

  1. 新增 tle.gpu DSL,支持 alloc / copy / to_tensor / subview / store_tensor / to_buffer。
  2. tle.gpu 原语可以下降到 CommonIR TileIR:tile.alloc / tile.copy / tile.to_tensor / tile.subview / tile.store_tensor。
  3. GPU 侧 TileIR buffer memory space 增加 global / shared / local / register。
  4. FLIR 侧新增 tile-infer-layout pass,补默认 TileIR layout。
  5. NVIDIA 后端增加 TileIR -> TTGIR lowering,matmul 路径中 tile.copy/to_tensor 可继续降到 triton_gpu.local_alloc / async_copy_global_to_local / local_load。
  6. native_matmul_gpu.py 新增 tle.gpu matmul 对照测试。
  7. native_fa.py 已替换 Q/K/V staging 路径为 tle.gpu.alloc/copy/to_tensor。

atalman and others added 30 commits June 18, 2024 15:43
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 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>
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.
kateyijian and others added 15 commits July 3, 2026 14:52
* 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>
@CLAassistant

Copy link
Copy Markdown

CLA assistant check
Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
16 out of 19 committers have signed the CLA.

✅ Galaxy1458
✅ git-flyer
✅ small-cat
✅ lindylin1817
✅ Kylin1207
✅ Stardep-lmc
✅ Quamaly
✅ wangziy1
✅ yanminghui123
✅ zhongsanming
✅ ph0375
✅ toudefu
✅ CHAIYD-2025
✅ zhzhcookie
✅ SabYic
✅ huanghaoXcore
❌ macqueen09
❌ 103yiran
❌ kateyijian
You have signed the CLA already but the status is still pending? Let us recheck it.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.