[KMCompiler][TLE] Common IR triton35 by zhongsanming · Pull Request #974 · flagos-ai/FlagTree

zhongsanming · 2026-08-13T02:32:37Z

引入一种新的Common IR作为TTIR的补充，CommonIR的主要功能是对算子层暴露更低层次的抽象，让算子层能解锁更多优化的机会。这个PR在triton 3.5 分支上完成了昇腾设备的POC工作，打通了从tle.dsa到Common IR到linalg IR到npubin的全流程。

现有flaggems全量op的精度测试结果：飞书文档

如果飞书文档没有权限请先使用以下链接：
summary

环境配置：

CANN： 9.1.0 ，目前9.1.0以下的CANN版本无法通过编译
硬件：昇腾910B/910C

* fix(test): pass combine_batch * fix(test): replace PP * fix(test): pass BLOCK_M, BLOCK_N, DIM * fix(test): add tl.constexpr * fix(test): fix workspace type

* fix(test): remove vid and HALF_M * feat(test): add fa dummy * feat(test): comment code * feat(test): add linalg_ir dump * fix(third_party): comment add_triton_to_annotation

Co-authored-by: zsm <zsm@compiler-dev.com>

* 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>

* test(test): add matmul subview * test(test): add no sync fa * test(test): add gpu pass fa * test(test): add fa_serial * test(test): add missing sync * fix(test): recover sub function * refactor(test): remove used variables * fix(third_party): add tile::MemorySpace::Global * fix(CMakeLists.txt): add no-dangling-assignment-gsl

* fix(test): fix L0C to UB error * fix(test): fix L0C to UB error * test(test): add fa 4func * test(test): add fa serial * test(test): add matmul_low_dim

* fix(test): fix cube use ub error * refactor(test): use tl.load

* fix(test): fix cube use ub error * refactor(test): use tl.load * fix(test): remove l1 alloc and copy, fix sync error

Co-authored-by: zhongsanming <zhongsanming@compiler-dev.com>

* fix(test): fix cube use ub error * refactor(test): use tl.load * fix(test): fix fa compile error * fix(test): fix rebase error

CLAassistant · 2026-08-13T02:32:47Z

All committers have signed the CLA.

sunnycase

LGTM

kateyijian · 2026-08-14T08:49:19Z


  registry.insert<
      mlir::triton::TritonDialect, mlir::cf::ControlFlowDialect,
+      mlir::triton::tile::TileIRDialect, mlir::hivm::HIVMDialect,


使用宏进行隔离（目前只影响ascend 后端）

kateyijian · 2026-08-14T08:50:24Z

@@ -0,0 +1,184 @@
+# 2× L1 缓冲区在 Linalg 降级流水线中被意外消除


文档需要删除

kateyijian · 2026-08-14T08:53:00Z

-                    >();
+                    DLTIDialect,                   // flagtree tle raw
+                    mlir::triton::tle::TleDialect, // flagtree tle raw
+                    ::mlir::index::IndexDialect, func::FuncDialect>();


CI 所有厂商 triton 3.5.x: 燧原，昇腾

kateyijian · 2026-08-14T08:54:43Z



+@triton.jit
+def copy_buffer(x):


dsa用例放哪里合理

kateyijian · 2026-08-14T09:03:31Z

             return op->getResult(0);
           })
-      // Add an annotation
+      .def("create_tanh",


删除文件

kateyijian · 2026-08-14T09:17:47Z

 if(FLAGTREE_BACKEND STREQUAL "ascend")
  include_directories(${CMAKE_CURRENT_SOURCE_DIR}/dsa/dialect/include)
  include_directories(${CMAKE_CURRENT_BINARY_DIR}/dsa/dialect/include)
+  include_directories(${CMAKE_SOURCE_DIR}/third_party/flir/include)


删除一行

kateyijian · 2026-08-14T09:18:47Z

    MLIRTargetLLVMIRImport
+    TleIR
+    TileIRIR
+    BiShengIRHIVMDialect


删除这一行

kateyijian · 2026-08-14T09:20:25Z

 if(NOT MSVC)
  if(FLAGTREE_BACKEND)
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -Wno-dangling-assignment-gsl")


@macqueen09 删除这一行

kateyijian · 2026-08-14T09:20:34Z

    endif()
  else()
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -fvisibility=hidden")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -Wno-dangling-assignment-gsl -fvisibility=hidden")


@103yiran 删除这一行

kateyijian · 2026-08-14T09:24:08Z

@103yiran 加到flagtree CI用例

kateyijian and others added 30 commits August 13, 2026 10:18

Initialize Common IR

df2c9a9

Initialize Common IR

0ca982f

Add DSL extension

0f0f46f

Add build script

0230eab

fix compiler error

020b269

Add compiler README file

748e4c5

Add the conversion tileir to hivm ir

9c18fbc

CommonIR: emit TileIR buffer ops from TLE DSA frontend

41c39d4

CommonIR: add buffer subview frontend support

7bd6f99

tileir: convert copy(GM->X) to memref.copy

1269935

Rewrite FA Op

77981de

Fa op fix (#8)

1e8308a

* fix(test): pass combine_batch * fix(test): replace PP * fix(test): pass BLOCK_M, BLOCK_N, DIM * fix(test): add tl.constexpr * fix(test): fix workspace type

add Fa dummy kernel (#9)

4e2621c

* fix(test): remove vid and HALF_M * feat(test): add fa dummy * feat(test): comment code * feat(test): add linalg_ir dump * fix(third_party): comment add_triton_to_annotation

Remove redundant copy (#10)

d36c34c

Co-authored-by: zsm <zsm@compiler-dev.com>

Add replace mlir

328c794

fix(test): fix pv_part name

f879710

feat(test): use with

807475d

fix(test): fix attn_score name

f4faf58

fix(test): remove l0c alloc

36f21c4

fix(third_party): disable ubuf_saving

d46fab3

CommonIR: add scope marker and cleanup TileIR lowering

b873095

feat(test): use scope

9379436

Fix TileIR copy memory-space lowering

cb45e7a

Emit scope op for TLE scope

37a923d

Fix compile and runtime issue

0375e3a

Next step to fixx error

ec77cdd

Common ir poc fix enum import (#14)

3ab5bda

* fix CANN enum name and numpy import * add solution for some issue --------- Co-authored-by: zhongsanming <sanming.zhong@compiler-dev.com>

Enable compile flow and Success to run matmul

5f50b75

103yiran and others added 13 commits August 13, 2026 10:20

test(test): add matmul bench (#32)

e3c5b95

move case into success/fail folder

d655b0e

move matmul_double_buffer into success folder

db8e073

create performance case

e1b1ca2

[TEST] add matmul and fa simple case (#57)

87731f0

* fix(test): fix L0C to UB error * fix(test): fix L0C to UB error * test(test): add fa 4func * test(test): add fa serial * test(test): add matmul_low_dim

Fa 3task fix (#58)

b084696

* fix(test): fix cube use ub error * refactor(test): use tl.load

Fa l1 fix (#64)

1f6bc16

* fix(test): fix cube use ub error * refactor(test): use tl.load * fix(test): remove l1 alloc and copy, fix sync error

fix pre-commit issues (#66)

64cc4cf

Co-authored-by: zhongsanming <zhongsanming@compiler-dev.com>

Add mhc kernel implementation (#65)

e6a5510

[TEST] fix fa compile error (#68)

e7f6785

* fix(test): fix cube use ub error * refactor(test): use tl.load * fix(test): fix fa compile error * fix(test): fix rebase error

fix fa_4func.py (#61)

48e7898

reformat

e59ce83

github-actions Bot added DOC Improvements or additions to documentation tle ascend CORE triton_v3.5.x labels Aug 13, 2026

zhongsanming marked this pull request as ready for review August 13, 2026 06:00

zhongsanming requested review from Galaxy1458, sunnycase and zhzhcookie as code owners August 13, 2026 06:00

zhongsanming changed the title ~~Common ir poc triton35~~ [KMCompiler][TLE] Common IR triton35 Aug 13, 2026

zhongsanming marked this pull request as draft August 13, 2026 07:24

zhongsanming force-pushed the common-ir-poc-triton35 branch from 73e293d to e59ce83 Compare August 13, 2026 07:25

sunnycase previously approved these changes Aug 13, 2026

View reviewed changes

kateyijian reviewed Aug 14, 2026

View reviewed changes

pin flir commit before merge for CI checks

f4706de

zhongsanming dismissed sunnycase’s stale review via f4706de August 15, 2026 03:11

Navigation Menu

Uh oh!

[KMCompiler][TLE] Common IR triton35 - #974

[KMCompiler][TLE] Common IR triton35#974
zhongsanming wants to merge 85 commits into
flagos-ai:triton_v3.5.xfrom
KernelLLM:common-ir-poc-triton35

zhongsanming commented Aug 13, 2026 •

edited

Loading

Uh oh!

CLAassistant commented Aug 13, 2026 •

edited

Loading

Uh oh!

sunnycase left a comment

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

kateyijian Aug 14, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

		@@ -0,0 +1,184 @@
		# 2× L1 缓冲区在 Linalg 降级流水线中被意外消除



		@triton.jit
		def copy_buffer(x):

Uh oh!

Conversation

zhongsanming commented Aug 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

CLAassistant commented Aug 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sunnycase left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

zhongsanming commented Aug 13, 2026 •

edited

Loading

CLAassistant commented Aug 13, 2026 •

edited

Loading