[BACKEND] Add backend tsingmicro

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

@@ -0,0 +1,60 @@
+name: Tsingmicro-Build-And-Test
+
+on:
+  push:
+    branches: [ "triton_v3.3.x" ]
+  pull_request:
+    branches: [ "triton_v3.3.x" ]
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: true
+
+jobs:
+  tsingmicro-build-and-test:
+    runs-on: tsingmicro
+    steps:
+      - name: Checkout code (attempt 1)
+        id: checkout1
+        uses: actions/checkout@v4
+        continue-on-error: true
+
+      - name: Sleep before checkout2
+        if: steps.checkout1.outcome == 'failure'
+        run: |
+          echo "First checkout attempt failed. Sleeping for 120 seconds before retry..."
+          sleep 120
+
+      - name: Checkout code (attempt 2)
+        id: checkout2
+        if: steps.checkout1.outcome == 'failure'
+        uses: actions/checkout@v4
+        continue-on-error: true
+
+      - name: Sleep before final checkout
+        if: steps.checkout1.outcome == 'failure' && steps.checkout2.outcome == 'failure'
+        run: |
+          echo "Second checkout attempt failed. Sleeping for 180 seconds before final retry..."
+          sleep 180
+
+      - name: Checkout code (final attempt)
+        if: steps.checkout1.outcome == 'failure' && steps.checkout2.outcome == 'failure'
+        uses: actions/checkout@v4
+
+      - name: Verify checkout success
+        if: success()
+        run: echo "Checkout completed successfully"
+
+      - name: FlagTree Build on Tsingmicro
+        shell: bash
+        run: |
+          source ~/env.sh
+          export FLAGTREE_BACKEND=tsingmicro
+          cd python
+          python3.10 -m pip install . --no-build-isolation -v
+
+      - name: FlagTree Test on Tsingmicro
+        shell: bash
+        run: |
+          source ~/env.sh
+          python3.10 -c 'import triton; print(triton.__path__)'

.gitignore

@@ -26,6 +26,8 @@ third_party/cambricon/
 third_party/iluvatar/iluvatarTritonPlugin.so
 third_party/triton_shared/
 third_party/xpu/backend/xpu3
+third_party/tsingmicro/backend/lib
+third_party/tsingmicro/backend/bin
 
 # Proton
 python/triton/profiler

.pre-commit-config.yaml

@@ -39,25 +39,6 @@ repos:
     hooks:
       - id: clang-format
 
-  # Expand YAML anchors in files used by github workflows, because github can't
-  # do this itself.  This lets us use anchors, which avoids code duplication.
-  - repo: local
-    hooks:
-    - id: expand-yaml-anchors
-      name: Expand YAML anchors
-      language: golang
-      additional_dependencies: [github.com/mikefarah/yq/v4@latest]
-      entry: >
-        bash -c '
-          OUT=".github/workflows/integration-tests.yml"
-          IN="$OUT.in"
-          echo "# AUTOGENERATED by pre-commit, modify the .in file instead." > "$OUT" &&
-          echo >> "$OUT"
-          yq "explode(.)" "$IN" >> "$OUT"
-        '
-      files: ^.github/workflows/integration-tests.yml.*
-      pass_filenames: false
-
 exclude: |
   (?x)(
     ^include/triton/external/|

CMakeLists.txt

@@ -29,6 +29,9 @@ elseif(FLAGTREE_BACKEND STREQUAL "mthreads")
 elseif(FLAGTREE_BACKEND STREQUAL "aipu")
   add_definitions(-D__NVIDIA__)
   add_definitions(-D__AMD__)
+elseif(FLAGTREE_BACKEND STREQUAL "tsingmicro")
+  set(CMAKE_C_COMPILER clang)
+  set(CMAKE_CXX_COMPILER clang++)
 endif()
 set(FLAGTREE_PLUGIN "$ENV{FLAGTREE_PLUGIN}")
 if(FLAGTREE_PLUGIN)
@@ -204,7 +207,7 @@ include_directories(${PROJECT_SOURCE_DIR}/third_party)
 include_directories(${PROJECT_BINARY_DIR}/third_party) # Tablegen'd files
 
 # link_directories(${LLVM_LIBRARY_DIR})
-if (FLAGTREE_BACKEND MATCHES "^(cambricon|aipu)$")
+if (FLAGTREE_BACKEND MATCHES "^(cambricon|aipu|tsingmicro)$")
   include_directories(${PROJECT_SOURCE_DIR}/include)
   include_directories(${PROJECT_BINARY_DIR}/include) # Tablegen'd files
   add_subdirectory(include)
@@ -364,6 +367,12 @@ if(TRITON_BUILD_PYTHON_MODULE)
       LLVMXPUCodeGen
       LLVMXPUAsmParser
     )
+  elseif(FLAGTREE_BACKEND STREQUAL "tsingmicro")
+    list(APPEND TRITON_LIBRARIES
+	  # riscv
+      LLVMRISCVCodeGen
+      LLVMRISCVAsmParser
+    )
   endif()
 
   if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64" OR # Linux arm64
@@ -446,7 +455,7 @@ find_package(Threads REQUIRED)
 
 add_subdirectory(third_party/f2reduce)
 
-if(NOT FLAGTREE_BACKEND OR FLAGTREE_BACKEND STREQUAL "aipu")
+if(NOT FLAGTREE_BACKEND OR FLAGTREE_BACKEND MATCHES "^(aipu|tsingmicro)$")
   add_subdirectory(bin)
   add_subdirectory(test)
 endif()

README.md

@@ -51,6 +51,16 @@ cd ${YOUR_CODE_DIR}/flagtree/python
 export FLAGTREE_BACKEND=mthreads
 python3 -m pip install . --no-build-isolation -v
 ```
+```shell
+# tsingmicro
+# Recommended: Use Ubuntu 20.04
+mkdir -p ~/.flagtree/tsingmicro; cd ~/.flagtree/tsingmicro
+wget https://github.com/FlagTree/flagtree/releases/download/v0.2.0-build-deps/tsingmicro-llvm21-glibc2.35-glibcxx3.4.30-x64.tar.gz
+cd ${YOUR_CODE_DIR}/flagtree/
+git checkout -b triton_v3.3.x origin/triton_v3.3.x
+export FLAGTREE_BACKEND=tsingmicro
+python3 -m pip install . --no-build-isolation -v
+```
 
 To build with default backends (nvidia, amd, triton_shared):
 ```shell

README_cn.md

@@ -51,6 +51,16 @@ cd ${YOUR_CODE_DIR}/flagtree/python
 export FLAGTREE_BACKEND=mthreads
 python3 -m pip install . --no-build-isolation -v
 ```
+```shell
+# tsingmicro
+# 推荐使用镜像 Ubuntu 20.04
+mkdir -p ~/.flagtree/tsingmicro; cd ~/.flagtree/tsingmicro
+wget https://github.com/FlagTree/flagtree/releases/download/v0.2.0-build-deps/tsingmicro-llvm21-glibc2.35-glibcxx3.4.30-x64.tar.gz
+cd ${YOUR_CODE_DIR}/flagtree/
+git checkout -b triton_v3.3.x origin/triton_v3.3.x
+export FLAGTREE_BACKEND=tsingmicro
+python3 -m pip install . --no-build-isolation -v
+```
 
 使用默认的编译命令，可以编译安装 nvidia、amd、triton_shared 后端：
 ```shell

bin/RegisterTritonDialects.h

@@ -69,8 +69,8 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::triton::gpu::registerAllocateSharedMemoryPass();
   mlir::triton::gpu::registerTritonGPUAllocateWarpGroups();
   mlir::triton::gpu::registerTritonGPUGlobalScratchAllocationPass();
-  mlir::triton::registerConvertWarpSpecializeToLLVM();
 #ifdef __NVIDIA__
+  mlir::triton::registerConvertWarpSpecializeToLLVM();
   mlir::triton::registerConvertTritonGPUToLLVMPass();
   mlir::triton::registerConvertNVGPUToLLVMPass();
 #endif

python/setup.py

@@ -432,6 +432,7 @@ def build_extension(self, ext):
         thirdparty_cmake_args = get_thirdparty_packages([get_llvm_package_info()])
         thirdparty_cmake_args += self.get_pybind11_cmake_args()
         extdir = os.path.abspath(os.path.dirname(self.get_ext_fullpath(ext.path)))
+        ext_base_dir = os.path.abspath(os.path.dirname(self.get_ext_fullpath(ext.name)))
         # create build directories
         if not os.path.exists(self.build_temp):
             os.makedirs(self.build_temp)
@@ -471,6 +472,7 @@ def build_extension(self, ext):
                 "-DCMAKE_EXE_LINKER_FLAGS=-fuse-ld=lld",
                 "-DCMAKE_MODULE_LINKER_FLAGS=-fuse-ld=lld",
                 "-DCMAKE_SHARED_LINKER_FLAGS=-fuse-ld=lld",
+                f"-DCMAKE_INSTALL_PREFIX={ext_base_dir}",
             ]
 
         # Note that asan doesn't work with binaries that use the GPU, so this is
@@ -512,6 +514,7 @@ def build_extension(self, ext):
         subprocess.check_call(["cmake", self.base_dir] + cmake_args, cwd=cmake_dir, env=env)
         subprocess.check_call(["cmake", "--build", "."] + build_args, cwd=cmake_dir)
         subprocess.check_call(["cmake", "--build", ".", "--target", "mlir-doc"], cwd=cmake_dir)
+        subprocess.check_call(["cmake", "--install", "."], cwd=cmake_dir)
 
 
 nvidia_version_path = os.path.join(get_base_dir(), "cmake", "nvidia-toolchain-version.json")
@@ -597,7 +600,7 @@ def build_extension(self, ext):
 )
 
 if helper.flagtree_backend:
-    if helper.flagtree_backend == "aipu":
+    if helper.flagtree_backend in ("aipu", "tsingmicro"):
         backends = [
             *BackendInstaller.copy(helper.default_backends + helper.extend_backends),
             *BackendInstaller.copy_externals(),

python/setup_helper.py

@@ -10,13 +10,13 @@
 import hashlib
 from dataclasses import dataclass
 
+flagtree_backend = os.getenv("FLAGTREE_BACKEND", "").lower()
+flagtree_plugin = os.getenv("FLAGTREE_PLUGIN", "").lower()
 use_triton_shared = False
-necessary_third_party = ["flir"]
+necessary_third_party = ["" if flagtree_backend == "tsingmicro" else "flir"]
 default_backends = ["nvidia", "amd"]
 extend_backends = []
 ext_sourcedir = "triton/_C/"
-flagtree_backend = os.getenv("FLAGTREE_BACKEND", "").lower()
-flagtree_plugin = os.getenv("FLAGTREE_PLUGIN", "").lower()
 
 
 @dataclass
@@ -39,7 +39,6 @@ class FlagTreeBackend:
 }
 
 set_llvm_env = lambda path: set_env({
-    'LLVM_BUILD_DIR': path,
     'LLVM_INCLUDE_DIRS': Path(path) / "include",
     'LLVM_LIBRARY_DIR': Path(path) / "lib",
     'LLVM_SYSPATH': path,
@@ -239,7 +238,7 @@ def skip_package_dir(package):
     @staticmethod
     def get_package_dir(packages):
         package_dict = {}
-        if flagtree_backend and flagtree_backend not in ("cambricon", "aipu"):
+        if flagtree_backend and flagtree_backend not in ("cambricon", "aipu", "tsingmicro"):
             connection = []
             backend_triton_path = f"../third_party/{flagtree_backend}/python/"
             for package in packages:
@@ -284,7 +283,8 @@ def git_clone(lib, lib_path):
                       "so we couldn't compile triton_shared\n")
 
         third_partys = []
-        third_partys.append(flagtree_backend_info["flir"])
+        if flagtree_backend != "tsingmicro":
+            third_partys.append(flagtree_backend_info["flir"])
         if os.environ.get("USE_TRITON_SHARED", "ON") == "ON":
             third_partys.append(flagtree_backend_info["triton_shared"])
         else:
@@ -305,9 +305,10 @@ def handle_flagtree_backend():
     if flagtree_backend:
         print(f"flagtree_backend is {flagtree_backend}")
         extend_backends.append(flagtree_backend)
-        if "editable_wheel" in sys.argv and flagtree_backend != "aipu":
+        if "editable_wheel" in sys.argv and flagtree_backend not in ("aipu", "tsingmicro"):
             ext_sourcedir = os.path.abspath(f"../third_party/{flagtree_backend}/python/{ext_sourcedir}") + "/"
-    default_backends.append("flir")
+    if flagtree_backend != "tsingmicro":
+        default_backends.append("flir")
     if use_triton_shared:
         default_backends.append("triton_shared")
 
@@ -335,7 +336,7 @@ def check_env(env_val):
     file="iluvatar-llvm18-x86_64",
     condition=("iluvatar" == flagtree_backend),
     url="https://github.com/FlagTree/flagtree/releases/download/v0.1.0-build-deps/iluvatar-llvm18-x86_64.tar.gz",
-    pre_hock=lambda: check_env('LLVM_BUILD_DIR'),
+    pre_hock=lambda: check_env('LLVM_SYSPATH'),
     post_hock=set_llvm_env,
 )
 
@@ -344,7 +345,7 @@ def check_env(env_val):
     file="XTDK-llvm18-ubuntu2004_x86_64",
     condition=("xpu" == flagtree_backend),
     url="https://github.com/FlagTree/flagtree/releases/download/v0.1.0-build-deps/XTDK-llvm18-ubuntu2004_x86_64.tar",
-    pre_hock=lambda: check_env('LLVM_BUILD_DIR'),
+    pre_hock=lambda: check_env('LLVM_SYSPATH'),
     post_hock=set_llvm_env,
 )
 
@@ -355,10 +356,10 @@ def check_env(env_val):
 cache.store(
     files=("clang", "xpu-xxd", "xpu3-crt.xpu", "xpu-kernel.t", "ld.lld", "llvm-readelf", "llvm-objdump",
            "llvm-objcopy"), condition=("xpu" == flagtree_backend),
-    copy_src_path=f"{os.environ.get('LLVM_BUILD_DIR','')}/bin", copy_dst_path="third_party/xpu/backend/xpu3/bin")
+    copy_src_path=f"{os.environ.get('LLVM_SYSPATH','')}/bin", copy_dst_path="third_party/xpu/backend/xpu3/bin")
 
 cache.store(files=("libclang_rt.builtins-xpu3.a", "libclang_rt.builtins-xpu3s.a"),
-            condition=("xpu" == flagtree_backend), copy_src_path=f"{os.environ.get('LLVM_BUILD_DIR','')}/lib/linux",
+            condition=("xpu" == flagtree_backend), copy_src_path=f"{os.environ.get('LLVM_SYSPATH','')}/lib/linux",
             copy_dst_path="third_party/xpu/backend/xpu3/lib/linux")
 
 cache.store(files=("include", "so"), condition=("xpu" == flagtree_backend),
@@ -370,6 +371,16 @@ def check_env(env_val):
     condition=("mthreads" == flagtree_backend),
     url=
     "https://github.com/FlagTree/flagtree/releases/download/v0.1.0-build-deps/mthreads-llvm19-glibc2.34-glibcxx3.4.30-x64.tar.gz",
-    pre_hock=lambda: check_env('LLVM_BUILD_DIR'),
+    pre_hock=lambda: check_env('LLVM_SYSPATH'),
+    post_hock=set_llvm_env,
+)
+
+# tsingmicro
+cache.store(
+    file="tsingmicro-llvm21-glibc2.35-glibcxx3.4.30-x64",
+    condition=("tsingmicro" == flagtree_backend),
+    url=
+    "https://github.com/FlagTree/flagtree/releases/download/v0.2.0-build-deps/tsingmicro-llvm21-glibc2.35-glibcxx3.4.30-x64.tar.gz",
+    pre_hock=lambda: check_env('LLVM_SYSPATH'),
     post_hock=set_llvm_env,
 )

third_party/tsingmicro/CMakeLists.txt

@@ -0,0 +1,31 @@
+if(NOT DEFINED TX8_HOME)
+  if(DEFINED ENV{TX8_HOME})
+    set(TX8_HOME $ENV{TX8_HOME})
+  else()
+    message(FATAL_ERROR "TX8_HOME environment variable is not defined")
+  endif()
+endif()
+
+set(XUANTIE_NAME Xuantie-900-gcc-elf-newlib-x86_64-V2.10.2)
+
+include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include)
+include_directories(${CMAKE_CURRENT_BINARY_DIR}/include)
+include_directories(${CMAKE_CURRENT_SOURCE_DIR}/crt/include)
+include_directories(${TX8_HOME}/include)
+add_subdirectory(include)
+add_subdirectory(lib)
+add_subdirectory(bin)
+add_subdirectory(crt)
+if(TRITON_BUILD_PYTHON_MODULE)
+  # FIXME: Unify the libraries for TsingMicro into fewer ones
+  add_triton_plugin(TritonTsingMicro ${CMAKE_CURRENT_SOURCE_DIR}/python/triton_tsingmicro.cc
+                    LINK_LIBS ZTCAnalysis ZTCAnalysisStructured MagicKernelIR
+                    Tx81IR TritonTilingExtIR TritonStructuredIR TritonToCoreDialects
+                    TritonToLinalg TritonToStructured StructuredToMemref LinalgToMagicKernel
+                    TritonArithToLinalg CoreDialectsToMK Tx81ToLLVM Tx81MemrefToLLVM MKToTx81)
+  target_link_libraries(TritonTsingMicro PRIVATE Python3::Module pybind11::headers)
+endif()
+#if(TRITON_BUILD_UT)
+#  add_subdirectory(unittest)
+#endif()
+#add_subdirectory(test)