[BACKEND] Update ascend backend to 29d243e 20260129

.github/workflows/ascend-build-and-test.yml

@@ -12,8 +12,15 @@ concurrency:
 
 jobs:
   ascend-build-and-test:
-    runs-on: ascend
+    runs-on: flagtree-ascend
+    if: ${{ github.repository == 'FlagTree/flagtree' || github.repository == 'flagos-ai/flagtree' }}
     steps:
+      - name: Setup environment
+        shell: bash
+        run: |
+          source ~/env.sh
+          env | grep -E '^(http_proxy|https_proxy|all_proxy|no_proxy)=' >> $GITHUB_ENV || true
+
       - name: Checkout code (attempt 1)
         id: checkout1
         uses: actions/checkout@v6
@@ -60,11 +67,10 @@ jobs:
         shell: bash
         run: |
           set -x
-          pip uninstall -y triton
           export FLAGTREE_BACKEND=ascend
           source ~/env.sh
           cd python
-          MAX_JOBS=32 python3 -m pip install . --no-build-isolation
+          MAX_JOBS=32 python3 -m pip install . --no-build-isolation -vvv
 
       - name: FlagTree Test on Ascend
         if: steps.check_files.outputs.only_docs_changed != 'true'

.github/workflows/nv-build-and-test.yml

@@ -111,7 +111,10 @@ jobs:
         shell: bash
         run: |
           set -x
-          python3.11 -m pytest -s python/test/unit
+          python3.11 -m pytest -s python/test/unit \
+              --ignore=python/test/unit/test_debug.py \
+              --ignore=python/test/unit/test_debug_dump.py \
+              --ignore=python/test/unit/tools/test_disasm.py
           if [ -d "python/test/operators" ]; then
             python3.11 -m pytest -s python/test/operators
           fi

CMakeLists.txt

@@ -410,7 +410,7 @@ if(TRITON_BUILD_PYTHON_MODULE)
   elseif(FLAGTREE_BACKEND STREQUAL "ascend")
     set(PYTHON_ROOT_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/python/src)
     include_directories(${PYTHON_ROOT_SRC_PATH})
-    add_library(triton SHARED ${PYTHON_ROOT_SRC_PATH}/main.cc
+    add_library(triton SHARED ${PYTHON_SRC_PATH}/main.cc
                 ${PYTHON_SRC_PATH}/ir.cc
                 ${PYTHON_ROOT_SRC_PATH}/passes.cc
                 ${PYTHON_ROOT_SRC_PATH}/interpreter.cc

include/triton/Dialect/Triton/IR/Dialect.h

@@ -13,16 +13,14 @@
 #include "mlir/Interfaces/FunctionInterfaces.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "triton/Dialect/Triton/IR/Dialect.h.inc"
-#include "triton/Dialect/Triton/IR/OpsEnums.h.inc"
-#include "triton/Dialect/Triton/IR/Traits.h"
-#include "triton/Dialect/Triton/IR/Types.h"
 
 #if __has_include("flagtree_spec.h")
 #include "flagtree_spec.h"
 #endif
-#if __has_include("triton/Dialect/Triton/IR/OpInterfaces.h")
-#include "triton/Dialect/Triton/IR/OpInterfaces.h"
-#endif
+
+#include "triton/Dialect/Triton/IR/OpsEnums.h.inc"
+#include "triton/Dialect/Triton/IR/Traits.h"
+#include "triton/Dialect/Triton/IR/Types.h"
 
 #define GET_OP_CLASSES
 #include "triton/Dialect/Triton/IR/Ops.h.inc"

lib/Dialect/Triton/IR/Ops.cpp

@@ -822,7 +822,6 @@ OpFoldResult AdvanceOp::fold(FoldAdaptor adaptor) {
 // https://github.com/llvm/llvm-project/blob/main/mlir/lib/Dialect/Func/IR/FuncOps.cpp
 // We could revert it back once MLIR has a better inliner interface.
 //-- FuncOp --
-#ifndef FLAGTREE_SPEC_Dialect_Triton_IR_Ops_build
 void FuncOp::build(OpBuilder &builder, OperationState &state, StringRef name,
                    FunctionType type, ArrayRef<NamedAttribute> attrs,
                    ArrayRef<DictionaryAttr> argAttrs) {
@@ -835,11 +834,14 @@ void FuncOp::build(OpBuilder &builder, OperationState &state, StringRef name,
   if (argAttrs.empty())
     return;
   assert(type.getNumInputs() == argAttrs.size());
+#if LLVM_VERSION_MAJOR < 21
   function_interface_impl::addArgAndResultAttrs(
+#else // triton_v3.3.x
+  call_interface_impl::addArgAndResultAttrs(
+#endif
       builder, state, argAttrs, /*resultAttrs=*/std::nullopt,
       getArgAttrsAttrName(state.name), getResAttrsAttrName(state.name));
 }
-#endif
 
 ParseResult FuncOp::parse(OpAsmParser &parser, OperationState &result) {
   auto buildFuncType =
@@ -918,7 +920,6 @@ LogicalResult ReturnOp::verify() {
 }
 
 // -- JoinOp --
-#ifndef FLAGTREE_SPEC_Dialect_Triton_IR_Ops_inferReturnTypes
 LogicalResult
 JoinOp::inferReturnTypes(MLIRContext *context, std::optional<Location> location,
                          ValueRange operands, DictionaryAttr attributes,
@@ -950,7 +951,6 @@ JoinOp::inferReturnTypes(MLIRContext *context, std::optional<Location> location,
       RankedTensorType::get(retShape, srcTy.getElementType(), retEnc));
   return success();
 }
-#endif
 
 // -- SplitOp --
 LogicalResult SplitOp::inferReturnTypes(

python/setup_tools/utils/ascend.py

@@ -4,7 +4,7 @@
 downloader = DownloadManager()
 
 submodules = (Module(name="ascendnpu-ir", url="https://gitcode.com/Ascend/AscendNPU-IR.git",
-                     commit_id="04045a06ec7c9592b17de659307d5debe7be590a",
+                     commit_id="0501294d3e",
                      dst_path=os.path.join(flagtree_configs.flagtree_submodule_dir, "ascendnpu-ir")), )
 
 

python/test/unit/test_debug.py

@@ -4,8 +4,6 @@
 import triton.language as tl
 import triton
 
-
-@pytest.mark.skip(reason="flagtree")
 @pytest.mark.parametrize('cond, opt_flag, env_var', [
     (cond, opt_flag, env_var) for cond in [True, False] \
                               for opt_flag in [True, False] \
@@ -30,7 +28,6 @@ def _kernel(COND: tl.constexpr):
     getattr(torch, device).synchronize()
 
 
-@pytest.mark.skip(reason="flagtree")
 @pytest.mark.parametrize("cond", [False, True])
 def test_static_assert(cond):
 
@@ -64,7 +61,6 @@ def _test_overflow(x, y, x_dtype, y_dtype, debug, should_overflow, tri_func, ref
 # integer overflow sanitization
 
 
-@pytest.mark.skip(reason="flagtree")
 @pytest.mark.parametrize("x, y, x_dtype, y_dtype, debug, should_overflow", [
     (-2**31, -1, 'int32', 'int32', False, False),
     (-2**31, -1, 'int32', 'int32', True, True),
@@ -89,7 +85,6 @@ def _kernel_add(X, Y, Z):
 # mul overflow
 
 
-@pytest.mark.skip(reason="flagtree")
 @pytest.mark.parametrize("x, y, x_dtype, y_dtype, debug, should_overflow", [
     (2**30, 4, 'int32', 'int32', False, False),
     (2**30, 4, 'int32', 'int32', True, True),
@@ -111,7 +106,6 @@ def _kernel_mul(X, Y, Z):
 # sub overflow
 
 
-@pytest.mark.skip(reason="flagtree")
 @pytest.mark.parametrize("x, y, x_dtype, y_dtype, debug, should_overflow", [
     (-2**31, 1, 'int32', 'int32', False, False),
     (-2**31, 1, 'int32', 'int32', True, True),

python/test/unit/test_debug_dump.py

@@ -16,8 +16,6 @@ def enable_dump_context(pass_name="1"):
 
 
 def test_fn_dump(capfd, device, fresh_triton_cache):
-    return  # TODO: flagtree
-
     N = 1024
     src = torch.zeros(N, device=device)
 

python/test/unit/tools/test_disasm.py

@@ -5,7 +5,6 @@
 import triton.language as tl
 
 
-@pytest.mark.skip(reason="flagtree")
 def test_disam_cubin():
     if not triton.runtime.driver.active.get_current_target().backend == "cuda":
         pytest.skip("Test requires CUDA.")

third_party/ascend/backend/spec/include/flagtree_spec.h

@@ -1,7 +1,7 @@
 #ifndef ASCEND_FLAGTREE_SPEC_H_
 #define ASCEND_FLAGTREE_SPEC_H_
 
-#include "triton/Dialect/Triton/IR/ascend_Ops.h"
 #include "triton/Dialect/Triton/IR/ascend_Traits.h"
+#include "triton/Dialect/Triton/IR/OpInterfaces.h"
 
 #endif // ASCEND_FLAGTREE_SPEC_H_

third_party/ascend/backend/spec/include/triton/Dialect/Triton/IR/OpInterfaces.h

@@ -1,5 +1,4 @@
-// TODO: When upgrading to Triton 3.4.0, remove this file and use the upstream
-// Triton file.
+//FIXME TODO: When upgrading to Triton 3.4.0, del this file
 #ifndef TRITON_IR_OP_INTERFACES_H_
 #define TRITON_IR_OP_INTERFACES_H_
 

third_party/ascend/backend/spec/include/triton/Dialect/Triton/IR/TritonOpInterfaces.td

@@ -1,4 +1,4 @@
-// TODO: When upgrading to Triton 3.4.0, remove this file and use the upstream Triton file.
+// FIXME TODO: When upgrading to Triton 3.4.0, remove this file and use the upstream Triton file.
 #ifndef TRITON_OP_INTERFACES
 #define TRITON_OP_INTERFACES