From 1fe1a8725504da09cdab8d2b15642a10874f1765 Mon Sep 17 00:00:00 2001 From: ikushare <2697533527@qq.com> Date: Tue, 11 Nov 2025 12:31:37 +0000 Subject: [PATCH 1/7] Add a pass management interface for Ascend --- CMakeLists.txt | 4 ++++ third_party/ascend/backend/compiler.py | 3 +++ third_party/ascend/triton_ascend.cpp | 11 ++++++++++- 3 files changed, 17 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e2008ebfa..1dd46e0fa 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -135,6 +135,10 @@ function(add_triton_object name) ) if (FLAGTREE_BACKEND STREQUAL "ascend") + set(ASCENDNPU_IR_SRC_DIR ${PROJECT_SOURCE_DIR}/third_party/ascendnpu-ir) + set(ASCENDNPU_IR_BINARY_DIR ${PROJECT_BINARY_DIR}/third_party/ascendnpu-ir) + include_directories(${ASCENDNPU_IR_SRC_DIR}/bishengir/include) + include_directories(${ASCENDNPU_IR_BINARY_DIR}/bishengir/include) set(patched_depends "") foreach(dep ${ARG_DEPENDS}) list(FIND PATCHED_TRITON_DEPENDS "${dep}" index) diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index 865bee249..029fc268b 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -70,6 +70,9 @@ def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): pm.enable_debug() # Add pass here. ascend.passes.convert.add_triton_to_linalg_pipeline(pm) + ascend.passes.convert.add_triton_to_llvm(pm) + ascend.passes.convert.add_triton_to_hfusion(pm) + ascend.passes.convert.add_triton_to_hivm(pm) pm.run(mod) return str(mod) ''' diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index d323f95af..7260ade61 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -4,6 +4,9 @@ #include "mlir/Pass/PassManager.h" #include "passes.h" #include "triton-shared/Conversion/TritonToLinalgExperimental/TritonToLinalgExperimental.h" +#include "triton-shared/TritonToHFusion/TritonToHFusion.h" +#include "triton-shared/TritonToHIVM/TritonToHIVM.h" +#include "triton-shared/TritonToLLVM/TritonToLLVM.h" #define PY_SSIZE_T_CLEAN #include @@ -12,10 +15,16 @@ namespace py = pybind11; void init_triton_ascend_passes_convert(py::module &&m) { ADD_PASS_WRAPPER_0("add_triton_to_linalg_pipeline", mlir::triton::createTritonToLinalgExperimentalPass); + ADD_PASS_WRAPPER_0("add_triton_to_llvm", + mlir::triton::createTritonToLLVMPass); + ADD_PASS_WRAPPER_0("add_triton_to_hfusion", + mlir::triton::createTritonToHFusionPass); + ADD_PASS_WRAPPER_0("add_triton_to_hivm", + mlir::triton::createTritonToHIVMPass); } // register ascend passes to triton void init_triton_ascend(py::module &&m) { auto passes = m.def_submodule("passes"); init_triton_ascend_passes_convert(passes.def_submodule("convert")); -} \ No newline at end of file +} From 93632dbe8abfd766252148543d1d6caa1c1006c0 Mon Sep 17 00:00:00 2001 From: ph0375 Date: Wed, 19 Nov 2025 03:10:48 +0000 Subject: [PATCH 2/7] TritonLinearize_interface --- third_party/ascend/backend/compiler.py | 1 + third_party/ascend/triton_ascend.cpp | 3 +++ 2 files changed, 4 insertions(+) diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index 029fc268b..d81f6d450 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -73,6 +73,7 @@ def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): ascend.passes.convert.add_triton_to_llvm(pm) ascend.passes.convert.add_triton_to_hfusion(pm) ascend.passes.convert.add_triton_to_hivm(pm) + ascend.passes.convert.add_triton_linearize(pm) pm.run(mod) return str(mod) ''' diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 7260ade61..7d4bc4f72 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -7,6 +7,7 @@ #include "triton-shared/TritonToHFusion/TritonToHFusion.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" +#include "triton-shared/TritonLinearize/TritonLinearize.h" #define PY_SSIZE_T_CLEAN #include @@ -21,6 +22,8 @@ void init_triton_ascend_passes_convert(py::module &&m) { mlir::triton::createTritonToHFusionPass); ADD_PASS_WRAPPER_0("add_triton_to_hivm", mlir::triton::createTritonToHIVMPass); + ADD_PASS_WRAPPER_0("add_triton_linearize", + mlir::triton::createTritonLinearizePass); } // register ascend passes to triton From 1a70e1c535976757559b9b2c1c60d17540e44628 Mon Sep 17 00:00:00 2001 From: ph0375 Date: Wed, 19 Nov 2025 09:14:14 +0000 Subject: [PATCH 3/7] TritonToLinalgIncubated_interface --- third_party/ascend/CMakeLists.txt | 6 ++- third_party/ascend/backend/compiler.py | 52 ++++++++++++++++++++++---- third_party/ascend/triton_ascend.cpp | 29 ++++++++++---- 3 files changed, 72 insertions(+), 15 deletions(-) diff --git a/third_party/ascend/CMakeLists.txt b/third_party/ascend/CMakeLists.txt index 8b77a38ed..bf367f47c 100644 --- a/third_party/ascend/CMakeLists.txt +++ b/third_party/ascend/CMakeLists.txt @@ -2,6 +2,10 @@ #add_subdirectory(test) add_triton_plugin(TritonAscend ${CMAKE_CURRENT_SOURCE_DIR}/triton_ascend.cpp) -target_include_directories(TritonAscend PRIVATE ${CMAKE_SOURCE_DIR}/third_party/flir/include) +target_include_directories(TritonAscend PRIVATE + ${CMAKE_SOURCE_DIR}/third_party/flir/include + ${CMAKE_BINARY_DIR}/third_party/flir/include) + +add_dependencies(TritonAscend TritonToLinalgIncubatedPassIncGen) add_triton_library(Registrar Registrar.cc) diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index d81f6d450..61113e144 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -68,18 +68,56 @@ def make_ttir(mod, metadata, opt): def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): pm = ir.pass_manager(mod.context) pm.enable_debug() + enable_nd2nz_on_vector = metadata["enable_nd2nz_on_vector"] # Add pass here. - ascend.passes.convert.add_triton_to_linalg_pipeline(pm) - ascend.passes.convert.add_triton_to_llvm(pm) - ascend.passes.convert.add_triton_to_hfusion(pm) - ascend.passes.convert.add_triton_to_hivm(pm) + #ascend.passes.convert.add_triton_to_linalg_pipeline(pm) ascend.passes.convert.add_triton_linearize(pm) + ascend.passes.convert.add_triton_to_hivm(pm) + ascend.passes.convert.add_triton_to_hfusion(pm) + ascend.passes.convert.add_triton_to_llvm(pm) + ascend.passes.convert.add_triton_to_linalg_incubated( + pm, + global_kernel=False, + named_ops=True, + enable_nd2nz_on_vector=enable_nd2nz_on_vector + ) pm.run(mod) return str(mod) + # use triton_adapter to lower Triton-MLIR to linalg + # Get Triton-MLIR as string ''' - with open('/home/zhengyang/FlagTree/triton-op-ir/ops/01_vector_add/cache_ascend/add_kernel.ttadapter', 'r', encoding='utf-8') as f: - content = f.read() - return content + ttir_code = str(mod) + with tempfile.TemporaryDirectory() as tmpdir: + src_path = os.path.join(tmpdir, "kernel.ttir.mlir") + dst_path = os.path.join(tmpdir, "kernel.ttadapter.mlir") + Path(src_path).write_text(ttir_code) + triton_adapter_opt_path = _get_triton_adapter_opt_path() + + enable_nd2nz_on_vector = metadata["enable_nd2nz_on_vector"] + cmd_list = [ + triton_adapter_opt_path, + src_path, + "--discrete-mask-access-conversion", + "--triton-to-annotation", + "--triton-to-unstructure", + "--triton-to-hivm", + "--bubble-up-operation", + f"--triton-to-linalg=global-kernel=false named-ops={named_ops} "\ + f"enable-nd2nz-on-vector={enable_nd2nz_on_vector}", + "-o", + dst_path, + ] + if _is_ascend_sanitizer_enabled() or not _is_debug_line_info_disabled(): + cmd_list += ["--mlir-print-debuginfo"] # pass debug info + + ret = subprocess.run(cmd_list, capture_output=True, check=True) + if opt.debug: + dump_manager = get_dump_manager(metadata["hash"]) + dump_manager.put( + Path(dst_path).read_text(), "kernel.ttadapter.mlir", binary=False + ) + + return Path(dst_path).read_text() ''' diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 7d4bc4f72..0e8c6231a 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -5,9 +5,10 @@ #include "passes.h" #include "triton-shared/Conversion/TritonToLinalgExperimental/TritonToLinalgExperimental.h" #include "triton-shared/TritonToHFusion/TritonToHFusion.h" +#include "triton-shared/TritonLinearize/TritonLinearize.h" +#include "triton-shared/TritonToLinalgIncubated/TritonToLinalgPass.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" -#include "triton-shared/TritonLinearize/TritonLinearize.h" #define PY_SSIZE_T_CLEAN #include @@ -16,16 +17,30 @@ namespace py = pybind11; void init_triton_ascend_passes_convert(py::module &&m) { ADD_PASS_WRAPPER_0("add_triton_to_linalg_pipeline", mlir::triton::createTritonToLinalgExperimentalPass); - ADD_PASS_WRAPPER_0("add_triton_to_llvm", - mlir::triton::createTritonToLLVMPass); - ADD_PASS_WRAPPER_0("add_triton_to_hfusion", - mlir::triton::createTritonToHFusionPass); - ADD_PASS_WRAPPER_0("add_triton_to_hivm", - mlir::triton::createTritonToHIVMPass); ADD_PASS_WRAPPER_0("add_triton_linearize", mlir::triton::createTritonLinearizePass); + ADD_PASS_WRAPPER_0("add_triton_to_hivm", + mlir::triton::createTritonToHIVMPass); + ADD_PASS_WRAPPER_0("add_triton_to_hfusion", + mlir::triton::createTritonToHFusionPass); + ADD_PASS_WRAPPER_0("add_triton_to_llvm", + mlir::triton::createTritonToLLVMPass); + m.def( + "add_triton_to_linalg_incubated", + [](mlir::PassManager &pm, + bool global_kernel, + bool named_ops, + bool enable_nd2nz_on_vector) { + pm.addPass(mlir::triton::conv::createTritonToLinalgIncubatedPass( + global_kernel, named_ops, enable_nd2nz_on_vector)); + }, + py::arg("pm"), + py::arg("global_kernel"), + py::arg("named_ops"), + py::arg("enable_nd2nz_on_vector")); } + // register ascend passes to triton void init_triton_ascend(py::module &&m) { auto passes = m.def_submodule("passes"); From e63acab54e87b206a232e72401764d36b8c56f7e Mon Sep 17 00:00:00 2001 From: ph0375 Date: Fri, 21 Nov 2025 16:52:55 +0800 Subject: [PATCH 4/7] Rename --- third_party/ascend/triton_ascend.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 0e8c6231a..6c360225c 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -6,7 +6,7 @@ #include "triton-shared/Conversion/TritonToLinalgExperimental/TritonToLinalgExperimental.h" #include "triton-shared/TritonToHFusion/TritonToHFusion.h" #include "triton-shared/TritonLinearize/TritonLinearize.h" -#include "triton-shared/TritonToLinalgIncubated/TritonToLinalgPass.h" +#include "triton-shared/TritonToLinalgIncubated/TritonToLinalgIncubatedPass.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" @@ -31,7 +31,7 @@ void init_triton_ascend_passes_convert(py::module &&m) { bool global_kernel, bool named_ops, bool enable_nd2nz_on_vector) { - pm.addPass(mlir::triton::conv::createTritonToLinalgIncubatedPass( + pm.addPass(mlir::triton::Incubated::createTritonToLinalgIncubatedPass( global_kernel, named_ops, enable_nd2nz_on_vector)); }, py::arg("pm"), From 7a0da951c98c8642c7b988a4add701b65a4d5fb9 Mon Sep 17 00:00:00 2001 From: ikushare <2697533527@qq.com> Date: Tue, 25 Nov 2025 03:08:54 +0000 Subject: [PATCH 5/7] TritonToUnstructureIncubated --- CMakeLists.txt | 1 + third_party/ascend/backend/compiler.py | 1 + third_party/ascend/triton_ascend.cpp | 3 +++ 3 files changed, 5 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index 1dd46e0fa..86ca37ad8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -34,6 +34,7 @@ elseif(FLAGTREE_BACKEND STREQUAL "ascend") set(CMAKE_C_COMPILER clang) set(CMAKE_CXX_COMPILER clang++) add_compile_options("-Wno-deprecated-declarations") + add_compile_options("-Wno-dangling-assignment-gsl") add_compile_options("-Wno-error=deprecated-declarations") endif() set(FLAGTREE_PLUGIN "$ENV{FLAGTREE_PLUGIN}") diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index 61113e144..c5111b688 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -72,6 +72,7 @@ def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): # Add pass here. #ascend.passes.convert.add_triton_to_linalg_pipeline(pm) ascend.passes.convert.add_triton_linearize(pm) + ascend.passes.convert.add_triton_to_unstructure(pm) ascend.passes.convert.add_triton_to_hivm(pm) ascend.passes.convert.add_triton_to_hfusion(pm) ascend.passes.convert.add_triton_to_llvm(pm) diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 6c360225c..3103bc74f 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -9,6 +9,7 @@ #include "triton-shared/TritonToLinalgIncubated/TritonToLinalgIncubatedPass.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" +#include "triton-shared/TritonToUnstructureIncubated/UnstructureConversionPass.h" #define PY_SSIZE_T_CLEAN #include @@ -19,6 +20,8 @@ void init_triton_ascend_passes_convert(py::module &&m) { mlir::triton::createTritonToLinalgExperimentalPass); ADD_PASS_WRAPPER_0("add_triton_linearize", mlir::triton::createTritonLinearizePass); + ADD_PASS_WRAPPER_0("add_triton_to_unstructure", + mlir::triton::createTritonToUnstructureIncubatedPass); ADD_PASS_WRAPPER_0("add_triton_to_hivm", mlir::triton::createTritonToHIVMPass); ADD_PASS_WRAPPER_0("add_triton_to_hfusion", From 95bf484b03be381b9b7253d1cdfd2b85da96e739 Mon Sep 17 00:00:00 2001 From: ph0375 Date: Wed, 26 Nov 2025 10:30:29 +0000 Subject: [PATCH 6/7] DiscreteMaskAccessConversion --- third_party/ascend/backend/compiler.py | 2 +- third_party/ascend/triton_ascend.cpp | 7 ++++--- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index c5111b688..2dac61c1e 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -71,8 +71,8 @@ def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): enable_nd2nz_on_vector = metadata["enable_nd2nz_on_vector"] # Add pass here. #ascend.passes.convert.add_triton_to_linalg_pipeline(pm) + ascend.passes.convert.add_triton_discretemaskaccessconversion(pm) ascend.passes.convert.add_triton_linearize(pm) - ascend.passes.convert.add_triton_to_unstructure(pm) ascend.passes.convert.add_triton_to_hivm(pm) ascend.passes.convert.add_triton_to_hfusion(pm) ascend.passes.convert.add_triton_to_llvm(pm) diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 3103bc74f..5a316344b 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -3,25 +3,26 @@ */ #include "mlir/Pass/PassManager.h" #include "passes.h" +#include "triton-shared/DiscreteMaskAccessConversion/DiscreteMaskAccessConversionPass.h" #include "triton-shared/Conversion/TritonToLinalgExperimental/TritonToLinalgExperimental.h" #include "triton-shared/TritonToHFusion/TritonToHFusion.h" #include "triton-shared/TritonLinearize/TritonLinearize.h" #include "triton-shared/TritonToLinalgIncubated/TritonToLinalgIncubatedPass.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" -#include "triton-shared/TritonToUnstructureIncubated/UnstructureConversionPass.h" #define PY_SSIZE_T_CLEAN #include namespace py = pybind11; void init_triton_ascend_passes_convert(py::module &&m) { + + ADD_PASS_WRAPPER_0("add_triton_discretemaskaccessconversion", + mlir::triton::createDiscreteMaskAccessConversionPass); ADD_PASS_WRAPPER_0("add_triton_to_linalg_pipeline", mlir::triton::createTritonToLinalgExperimentalPass); ADD_PASS_WRAPPER_0("add_triton_linearize", mlir::triton::createTritonLinearizePass); - ADD_PASS_WRAPPER_0("add_triton_to_unstructure", - mlir::triton::createTritonToUnstructureIncubatedPass); ADD_PASS_WRAPPER_0("add_triton_to_hivm", mlir::triton::createTritonToHIVMPass); ADD_PASS_WRAPPER_0("add_triton_to_hfusion", From 881cd812696faa83a1a5d3fa97f07957d9cf9c9c Mon Sep 17 00:00:00 2001 From: ikushare <2697533527@qq.com> Date: Thu, 27 Nov 2025 02:19:57 +0000 Subject: [PATCH 7/7] TritonToAnnotation --- third_party/ascend/backend/compiler.py | 2 ++ third_party/ascend/triton_ascend.cpp | 6 ++++++ 2 files changed, 8 insertions(+) diff --git a/third_party/ascend/backend/compiler.py b/third_party/ascend/backend/compiler.py index 2dac61c1e..516d7ca33 100644 --- a/third_party/ascend/backend/compiler.py +++ b/third_party/ascend/backend/compiler.py @@ -73,6 +73,8 @@ def ttir_to_linalg(mod, metadata, opt, *, named_ops=False): #ascend.passes.convert.add_triton_to_linalg_pipeline(pm) ascend.passes.convert.add_triton_discretemaskaccessconversion(pm) ascend.passes.convert.add_triton_linearize(pm) + ascend.passes.convert.add_triton_to_annotation(pm) + ascend.passes.convert.add_triton_to_unstructure(pm) ascend.passes.convert.add_triton_to_hivm(pm) ascend.passes.convert.add_triton_to_hfusion(pm) ascend.passes.convert.add_triton_to_llvm(pm) diff --git a/third_party/ascend/triton_ascend.cpp b/third_party/ascend/triton_ascend.cpp index 5a316344b..00f86d87f 100644 --- a/third_party/ascend/triton_ascend.cpp +++ b/third_party/ascend/triton_ascend.cpp @@ -10,6 +10,8 @@ #include "triton-shared/TritonToLinalgIncubated/TritonToLinalgIncubatedPass.h" #include "triton-shared/TritonToHIVM/TritonToHIVM.h" #include "triton-shared/TritonToLLVM/TritonToLLVM.h" +#include "triton-shared/TritonToUnstructureIncubated/UnstructureConversionPass.h" +#include "triton-shared/TritonToAnnotation/TritonToAnnotation.h" #define PY_SSIZE_T_CLEAN #include @@ -23,6 +25,10 @@ void init_triton_ascend_passes_convert(py::module &&m) { mlir::triton::createTritonToLinalgExperimentalPass); ADD_PASS_WRAPPER_0("add_triton_linearize", mlir::triton::createTritonLinearizePass); + ADD_PASS_WRAPPER_0("add_triton_to_annotation", + mlir::triton::createTritonToAnnotationPass); + ADD_PASS_WRAPPER_0("add_triton_to_unstructure", + mlir::triton::createTritonToUnstructureIncubatedPass); ADD_PASS_WRAPPER_0("add_triton_to_hivm", mlir::triton::createTritonToHIVMPass); ADD_PASS_WRAPPER_0("add_triton_to_hfusion",