升级到triton3.5 重构tt到linalg

.github/workflows/format.yaml

@@ -34,20 +34,33 @@ jobs:
   clang-format:
     needs: markdownlint
     runs-on: ubuntu-latest
+
     steps:
       - uses: actions/checkout@v4
-      - uses: cpp-linter/cpp-linter-action@v2
-        id: cpp-lint
         with:
-          style: file
-          ignore: |
-            third_party/*
-            .git
-          tidy-checks: '-*' # disable clang tidy at this stage
-          version: 17
-      - name: Fail test
-        if: steps.cpp-lint.outputs.checks-failed > 0
-        run: echo "Some files failed the linting checks!" && exit 1
+          fetch-depth: 0
+
+      - name: Install clang-format
+        run: |
+          sudo apt-get update
+          sudo apt-get install -y clang-format-17
+
+      - name: Check changed files
+        run: |
+          git diff --diff-filter=d --name-only \
+            origin/${{ github.base_ref }}...HEAD \
+            > all_changed_files.txt
+
+          awk '/\.(cc|cpp|cxx|h|hpp)$/ && $0 !~ /^third_party\// { print }' \
+            all_changed_files.txt \
+            > changed_files.txt
+
+          if [ ! -s changed_files.txt ]; then
+            echo "No C/C++ files changed"
+            exit 0
+          fi
+
+          xargs -a changed_files.txt clang-format-17 --dry-run -Werror
 
   python-black:
     needs: markdownlint
@@ -58,4 +71,4 @@ jobs:
           fetch-depth: 0  # 确保能获取到 main 分支的历史
       - name: Check formatting for PR files
         run: |
-          bash scripts/ci_python_black.sh
+          bash scripts/ci_python_black.sh

.github/workflows/main.yml

@@ -10,7 +10,7 @@ on:
       - main
 
 env:
-  CI_PATH: "${{ vars.CI_BASE_PATH }}/GitHub/${{ github.repository }}/${GITHUB_RUN_NUMBER}"
+  CI_PATH: "${{ vars.CI_BASE_PATH }}/GitHub/${{ github.repository }}/${{ github.run_number }}"
   THIRD_PARTY_PATH: "${{ vars.CI_BASE_PATH }}/data/DLCompiler/third_party"
 
 concurrency:
@@ -31,23 +31,21 @@ jobs:
       - name: Create custom directory
         run: |
           set -ex
-          echo ${{ env.CI_PATH }}
-          mkdir -p ${{ env.CI_PATH }}
+          echo "$CI_PATH"
+          mkdir -p "$CI_PATH"
 
       - name: Clean custom directory
         run: |
           set -ex
-          if [ -d "${{ env.CI_PATH }}" ]; then
-            rm -rf ${{ env.CI_PATH }}/*
-            rm -rf ${{ env.CI_PATH }}/.github
-          fi
+          rm -rf "$CI_PATH"
+          mkdir -p "$CI_PATH"
 
       - name: Move code to custom directory
         run: |
           set -ex
-          mv $GITHUB_WORKSPACE/* ${{ env.CI_PATH }}/
-          mv $GITHUB_WORKSPACE/.github ${{ env.CI_PATH }}/
-          mv $GITHUB_WORKSPACE/.git ${{ env.CI_PATH }}/
+          mv "$GITHUB_WORKSPACE"/* "$CI_PATH"/
+          mv "$GITHUB_WORKSPACE/.github" "$CI_PATH"/
+          mv "$GITHUB_WORKSPACE/.git" "$CI_PATH"/
 
       - name: Setup Git Safe Directory
         run: |
@@ -68,16 +66,16 @@ jobs:
             which conda
             echo "which conda? $(which conda)"
             conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
+            source /usr/local/Ascend/ascend-toolkit/set_env.sh
+            source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
             cd ${{ env.CI_PATH }}
-            export JSON_PATH34=${{ vars.CI_BASE_PATH }}/data/v34/include.zip
-            export GOOGLETEST_DIR34=${{ vars.CI_BASE_PATH }}/data/v34/googletest
-            export LLVM_TGZ_PATH34=${{ vars.CI_BASE_PATH }}/data/v34/llvm-064f02da-ubuntu-arm64.tar.gz
+            export JSON_PATH35=${{ vars.CI_BASE_PATH }}/data/v35/include.zip
+            export GOOGLETEST_DIR35=${{ vars.CI_BASE_PATH }}/data/v35/googletest
+            export LLVM_TGZ_PATH35=${{ vars.CI_BASE_PATH }}/data/v35/llvm-7d5de303-ubuntu-arm64.tar.gz
 
             rm -rf ./third_party/*
             git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/ascendnpu-ir ./third_party/ascendnpu-ir
             git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/json ./third_party/json
-            git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/triton_shared ./third_party/triton_shared
             git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/triton ./third_party/triton
             echo "whoami before compile: $(whoami)"
             echo "which python? $(which python)"
@@ -89,37 +87,39 @@ jobs:
             bash compile_shared.sh apply_patch=true
           '
 
-      - name: Build and install tilelang-dlc
-        run: |
-          set -ex
-          # 切换到 root 用户执行
-          sudo -E bash -c '
-            source /home/dlc_ci/.bashrc
-            conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
-            cd ${{ env.CI_PATH }}
-            export TILELANG_DLC_PATH=${{ vars.CI_BASE_PATH }}/data/tilelang-dlc
-            export DLCOMPILER_SOURCE=${{ env.CI_PATH }}
-            export TILELANG_USE_DLCOMPILER=1
-            echo "whoami? $(whoami)"
-            echo "which python? $(which python)"
-            bash scripts/install_tilelang-dlc.sh
-          '
-
-      - name: Run tilelang-dlc tests on ascend
-        run: |
-          set -ex
-          # 切换到 root 用户执行
-          sudo -E bash -c '
-            source /home/dlc_ci/.bashrc
-            conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
-            cd ${{env.CI_PATH }}
-            export PATH=${{ vars.CI_BASE_PATH }}/data/bishengir_latest/:$PATH
-            export ASCEND_RT_VISIBLE_DEVICES=7
-            export TILELANG_USE_DLCOMPILER=1
-            bash test/commonir/run_tests.sh
-          '
+      # - name: Build and install tilelang-dlc
+      #   run: |
+      #     set -ex
+      #     # 切换到 root 用户执行
+      #     sudo -E bash -c '
+      #       source /home/dlc_ci/.bashrc
+      #       conda activate dlcompiler
+      #       source /usr/local/Ascend/ascend-toolkit/set_env.sh
+      #       source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
+      #       cd ${{ env.CI_PATH }}
+      #       export TILELANG_DLC_PATH=${{ vars.CI_BASE_PATH }}/data/tilelang-dlc
+      #       export DLCOMPILER_SOURCE=${{ env.CI_PATH }}
+      #       export TILELANG_USE_DLCOMPILER=1
+      #       echo "whoami? $(whoami)"
+      #       echo "which python? $(which python)"
+      #       bash scripts/install_tilelang-dlc.sh
+      #     '
+
+      # - name: Run tilelang-dlc tests on ascend
+      #   run: |
+      #     set -ex
+      #     # 切换到 root 用户执行
+      #     sudo -E bash -c '
+      #       source /home/dlc_ci/.bashrc
+      #       conda activate dlcompiler
+      #       source /usr/local/Ascend/ascend-toolkit/set_env.sh
+              # source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
+      #       cd ${{env.CI_PATH }}
+      #       export PATH=${{ vars.CI_BASE_PATH }}/data/bishengir_latest/:$PATH
+      #       export ASCEND_RT_VISIBLE_DEVICES=7
+      #       export TILELANG_USE_DLCOMPILER=1
+      #       bash test/commonir/run_tests.sh
+      #     '
 
       - name: Run triton tests on ascend
         run: |
@@ -128,11 +128,11 @@ jobs:
           sudo -E bash -c '
             source /home/dlc_ci/.bashrc
             conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
+            source /usr/local/Ascend/ascend-toolkit/set_env.sh
+            source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
             cd ${{env.CI_PATH }}
             echo "whoami? $(whoami)"
             echo "which python? $(which python)"
-            export PATH=${{ vars.CI_BASE_PATH }}/data/bishengir_latest/:$PATH
             export ASCEND_RT_VISIBLE_DEVICES=7
             bash test/ascend/run_tests.sh
           '
@@ -144,11 +144,11 @@ jobs:
           sudo -E bash -c '
             source /home/dlc_ci/.bashrc
             conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
+            source /usr/local/Ascend/ascend-toolkit/set_env.sh
+            source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
             cd ${{env.CI_PATH }}
             echo "whoami? $(whoami)"
             echo "which python? $(which python)"
-            export PATH=${{ vars.CI_BASE_PATH }}/data/bishengir_latest/:$PATH
             export ASCEND_RT_VISIBLE_DEVICES=7
             bash test/ascend/test_mlir.sh
           '
@@ -160,11 +160,11 @@ jobs:
           sudo -E bash -c '
             source /home/dlc_ci/.bashrc
             conda activate dlcompiler
-            source /usr/local/Ascend/cann-8.5.0/set_env.sh
+            source /usr/local/Ascend/ascend-toolkit/set_env.sh
+            source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
             cd ${{env.CI_PATH }}
             echo "whoami? $(whoami)"
             echo "which python? $(which python)"
-            export PATH=${{ vars.CI_BASE_PATH }}/data/bishengir_latest/:$PATH
             export ASCEND_RT_VISIBLE_DEVICES=7
             bash test/dsl/run_tests.sh
           '

.github/workflows/release.yml

@@ -6,7 +6,7 @@ on:
       - "v*.*.*"  # 只对形如 v1.2.3 的 tag 触发
 
 env:
-  RELEASE_PATH: "${{ vars.CI_BASE_PATH }}/GitHub/Release/${{ github.repository }}/${GITHUB_RUN_NUMBER}"
+  RELEASE_PATH: "${{ vars.CI_BASE_PATH }}/GitHub/Release/${{ github.repository }}/${{ github.run_number }}"
   THIRD_PARTY_PATH: "${{ vars.CI_BASE_PATH }}/data/DLCompiler/third_party"
 
 concurrency:
@@ -27,23 +27,21 @@ jobs:
       - name: Create custom directory
         run: |
           set -ex
-          echo ${{ env.RELEASE_PATH }}
-          mkdir -p ${{ env.RELEASE_PATH }}
+          echo "$RELEASE_PATH"
+          mkdir -p "$RELEASE_PATH"
 
       - name: Clean custom directory
         run: |
           set -ex
-          if [ -d "${{ env.RELEASE_PATH }}" ]; then
-            rm -rf ${{ env.RELEASE_PATH }}/*
-            rm -rf ${{ env.RELEASE_PATH }}/.github
-          fi
+          rm -rf "$RELEASE_PATH"
+          mkdir -p "$RELEASE_PATH"
 
       - name: Move code to custom directory
         run: |
           set -ex
-          mv $GITHUB_WORKSPACE/* ${{ env.RELEASE_PATH }}/
-          mv $GITHUB_WORKSPACE/.github ${{ env.RELEASE_PATH }}/
-          mv $GITHUB_WORKSPACE/.git ${{ env.RELEASE_PATH }}/
+          mv "$GITHUB_WORKSPACE"/* "$RELEASE_PATH"/
+          mv "$GITHUB_WORKSPACE/.github" "$RELEASE_PATH"/
+          mv "$GITHUB_WORKSPACE/.git" "$RELEASE_PATH"/
 
       - name: Setup Git Safe Directory
         run: |
@@ -60,15 +58,15 @@ jobs:
           source /home/dlc_ci/.bashrc
           source activate dlcompiler
           source /usr/local/Ascend/ascend-toolkit/set_env.sh
+          source /usr/local/Ascend/cann-9.1.0-beta.1/share/info/ascendnpu-ir/bin/set_env.sh
           cd ${{ env.RELEASE_PATH }}
-          export JSON_PATH34=${{ vars.CI_BASE_PATH }}/data/v34/include.zip
-          export GOOGLETEST_DIR34=${{ vars.CI_BASE_PATH }}/data/v34/googletest
-          export LLVM_TGZ_PATH34=${{ vars.CI_BASE_PATH }}/data/v34/llvm-064f02da-ubuntu-arm64.tar.gz
+          export JSON_PATH35=${{ vars.CI_BASE_PATH }}/data/v35/include.zip
+          export GOOGLETEST_DIR35=${{ vars.CI_BASE_PATH }}/data/v35/googletest
+          export LLVM_TGZ_PATH35=${{ vars.CI_BASE_PATH }}/data/v35/llvm-7d5de303-ubuntu-arm64.tar.gz
 
           rm -rf ./third_party/*
           git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/ascendnpu-ir ./third_party/ascendnpu-ir
           git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/json ./third_party/json
-          git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/triton_shared ./third_party/triton_shared
           git clone --no-hardlinks ${{ env.THIRD_PARTY_PATH }}/triton ./third_party/triton
           echo "whoami? $(whoami)"
           echo "which python? $(which python)"

.gitignore

@@ -99,6 +99,4 @@ fusion_result.json
 launcher_cxx11abi*
 
 # package
-backend/triton-shared-opt-v3*
-backend/dicp_opt
-third_party/triton-shared-opt
+backend/dicp_opt

.gitmodules

@@ -1,14 +1,8 @@
-[submodule "third_party/triton_linalg"]
-    path = third_party/triton_linalg
-    url = git@github.com:Cambricon/triton-linalg.git
 [submodule "third_party/triton"]
     path = third_party/triton
     url = git@github.com:triton-lang/triton.git
-[submodule "third_party/triton_shared"]
-    path = third_party/triton_shared
-    url = git@github.com:microsoft/triton-shared.git
-    depth = 1
+    branch = v3.5.0
 [submodule "third_party/ascendnpu-ir"]
     path = third_party/ascendnpu-ir
     url = https://gitcode.com/Ascend/AscendNPU-IR.git
-    depth = 1
+    branch = ef9139b323e25e8dae0a812ac585f3b47ab5d955

CMakeLists.txt

@@ -1,15 +1,17 @@
 set(DC_TRITON_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
 set(DC_TRITON_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
-set(TRTION_SHARED_NPU_SPECIFIC_SOURCES  "${DC_TRITON_SOURCE_DIR}/compiler/lib/Conversion/TritonToLinalgNPU")
 
 set(DC_TRITON_INCLUDE_DIR "")
 set(DC_TRITON_LINK_DIR "")
 
 include_directories(${CMAKE_CURRENT_BINARY_DIR}/compiler/include)
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/compiler/include)
-include_directories(${CMAKE_CURRENT_BINARY_DIR}/third_party/triton_shared/include)
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/json/include)
-include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/triton_shared/include)
+
+# ASCENDNPU_IR编译选项
+set(LLVM_MAJOR_VERSION_22_COMPATIBLE ON CACHE BOOL "NPUIR build with LLVM 22" FORCE)
+add_definitions(-D__LLVM_MAJOR_VERSION_22_COMPATIBLE__)
+add_compile_options(-Wno-switch)
 
 set(ASCENDNPU_IR_SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/third_party/ascendnpu-ir)
 set(ASCENDNPU_IR_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/third_party/ascendnpu-ir)
@@ -31,35 +33,45 @@ ${dialect_libs}
 ${conversion_libs}
 )
 
-list(APPEND CMAKE_MODULE_PATH 
+list(APPEND CMAKE_MODULE_PATH
   ${CMAKE_CURRENT_LIST_DIR}/cmake/
 )
 
 include(utils)
 add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/tools)
 add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/compiler)
-add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/third_party/triton_shared)
 
 if (TRITON_BUILD_PYTHON_MODULE)
-    add_triton_plugin(tritonDicpTriton ${CMAKE_CURRENT_SOURCE_DIR}/triton_dicp_triton.cc 
-    LINK_LIBS 
+    add_triton_plugin(tritonDicpTriton ${CMAKE_CURRENT_SOURCE_DIR}/triton_dicp_triton.cc
+    LINK_LIBS
 
     MLIRAffineToStandard
     MLIRIR
     MLIRPass
     MLIRTransforms
     MLIRSupport
     MLIRBytecodeWriter
-
-    TritonToLinalgNPUCoversion
-
-    LinalgExtTransforms
-    TritonExtTransforms
+    MLIRFuncInlinerExtension
+    MLIRBufferizationDialect
+    MLIRMemRefDialect
 
-    LinalgToLinked
-    LinkedToHIVM
-    DiscreteMaskAccessConversion
+    # DICP NPU pass libraries
+    TritonToLinalg
+    TritonToStructured
     TritonToUnstructure
+    TritonToHIVM
+    TritonToAnnotation
+    TritonToHFusion
+    TritonToLLVM
+    TritonToGraph
+    AutoBlockify
+    AscendLegalize
+    CommonIRTransforms
+    DiscreteMaskAccessConversion
+    MLIRTritonNPUUtils
+    TritonAffinityOpt
+    TritonDicpIR
+    TritonStructuredIR
     )
     target_link_libraries(tritonDicpTriton PRIVATE Python3::Module pybind11::headers)
-endif()
+endif()