[XPU] Sync internal 5a664566 cluster + SDNN v0.3.6.7.0 and conceal IR

CMakeLists.txt

@@ -113,9 +113,38 @@ include(AddMLIR)
 # Utilities
 function(add_triton_object name)
   cmake_parse_arguments(ARG "" "" "DEPENDS;LINK_LIBS" ${ARGN})
+
+  # FLAGTREE SPEC SOURCE OVERRIDE:
+  # For a per-backend whole-file replacement of a main-tree source, drop a copy
+  # (derived from the FlagTree source file + backend edits) at
+  #   third_party/${FLAGTREE_BACKEND}/backend/spec/<path-relative-to-project>/<file>
+  # This lets the backend compile its own version of a shared .cpp WITHOUT
+  # editing the pristine main-tree file (mirrors set_flagtree_backend_td for .td).
+  # EXISTS-gated: no effect for backends / files without an override file.
+  set(_flagtree_srcs "")
+  foreach(_src ${ARG_UNPARSED_ARGUMENTS})
+    set(_resolved "${_src}")
+    if(FLAGTREE_BACKEND)
+      if(IS_ABSOLUTE "${_src}")
+        set(_abs "${_src}")
+      else()
+        set(_abs "${CMAKE_CURRENT_SOURCE_DIR}/${_src}")
+      endif()
+      file(RELATIVE_PATH _rel "${PROJECT_SOURCE_DIR}" "${_abs}")
+      if(NOT _rel MATCHES "^\\.\\.")
+        set(_spec "${PROJECT_SOURCE_DIR}/third_party/${FLAGTREE_BACKEND}/backend/spec/${_rel}")
+        if(EXISTS "${_spec}")
+          set(_resolved "${_spec}")
+          message(STATUS "[FlagTree spec] source override: ${_rel} -> ${_spec}")
+        endif()
+      endif()
+    endif()
+    list(APPEND _flagtree_srcs "${_resolved}")
+  endforeach()
+
   add_library(${name} OBJECT)
   target_sources(${name}
-    PRIVATE ${ARG_UNPARSED_ARGUMENTS}
+    PRIVATE ${_flagtree_srcs}
     INTERFACE $<TARGET_OBJECTS:${name}>
   )
 

python/setup_tools/utils/xpu.py

@@ -392,8 +392,8 @@ def register_cache(cache, flagtree_backend, check_env, set_llvm_env):
                 copy_src_path=f"{cache.dir_path}/{flagtree_backend}/xpu-device-libs",
                 copy_dst_path=f"third_party/{flagtree_backend}/device")
     cache.store(file="xpu-sdnn-objects", condition=is_xpu,
-                url="https://klx-sdk-release-public.su.bcebos.com/XTriton/xpu-sdnn-objects_v0.3.6.6.0.tar.gz",
-                version="v0.3.6.6.0", post_hook=lambda path: install_sdnn_objects(path, cache.flagtree_dir))
+                url="https://klx-sdk-release-public.su.bcebos.com/XTriton/xpu-sdnn-objects_v0.3.6.7.1.tar.gz",
+                version="v0.3.6.7.1", post_hook=lambda path: install_sdnn_objects(path, cache.flagtree_dir))
     cache.store(
         files=("clang", "xpu-xxd", "xpu3-elfconv", "xpu3-elfconv-triton", "xpu-kernel.t", "ld.lld", "llvm-readelf",
                "llvm-objdump", "llvm-objcopy"), condition=is_xpu,

third_party/xpu/CMakeLists.txt

@@ -1,3 +1,14 @@
+# XPU: conceal all IR text (MLIR/LLVM/ASM) in non-Debug builds, aligning with
+# internal Triton 3.0 TRITON_CONCEAL_IR. XPU-only. Also cover
+# TritonRelBuildWithAsserts, which FlagTree uses as the default RelWithAsserts
+# type and does not inherit CMAKE_CXX_FLAGS_RELEASE.
+set(CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} -DTRITON_CONCEAL_IR=1")
+set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -DTRITON_CONCEAL_IR=1")
+set(CMAKE_C_FLAGS_TRITONRELBUILDWITHASSERTS "${CMAKE_C_FLAGS_TRITONRELBUILDWITHASSERTS} -DTRITON_CONCEAL_IR=1")
+set(CMAKE_CXX_FLAGS_TRITONRELBUILDWITHASSERTS "${CMAKE_CXX_FLAGS_TRITONRELBUILDWITHASSERTS} -DTRITON_CONCEAL_IR=1")
+add_compile_definitions($<$<NOT:$<CONFIG:Debug>>:TRITON_CONCEAL_IR=1>)
+message(STATUS "[XPU] TRITON_CONCEAL_IR=1 enabled for non-Debug XPU builds")
+
 # ============================================================================
 # XPU SDNN Object File Validation
 # ============================================================================
@@ -109,8 +120,17 @@ add_subdirectory(include)
 # compiles, without changing main-tree CMake files.
 if(TARGET TritonGPUIR)
   add_dependencies(TritonGPUIR TritonXPUAttrDefsIncGen)
+  target_include_directories(TritonGPUIR BEFORE PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}/include
+    ${CMAKE_CURRENT_BINARY_DIR}/include)
 endif()
 
+# ir.cc / llvm.cc are compiled as part of the main `triton` target, which is
+# created after this subdirectory. Defer so the conceal macro actually reaches
+# those sources.
+cmake_language(DEFER DIRECTORY "${CMAKE_SOURCE_DIR}" CALL
+  target_compile_definitions triton PRIVATE TRITON_CONCEAL_IR=1)
+
 add_subdirectory(lib)
 
 # ==================== FLAGTREE XPU SYNC MARK ====================

third_party/xpu/backend/compiler.py

@@ -107,6 +107,7 @@ class XPUOptions:
     buffer_size_limit: int = int(os.environ.get("TRITONXPU_BUFFER_SIZE", 512))
     groups_per_cluster: int = int(os.environ.get("TRITONXPU_GROUPS_PER_CLUSTER", 1))
     unroll_num: int = int(os.environ.get("TRITONXPU_UNROLL_NUM", 2))
+    vrf_budget: int = int(os.environ.get("TRITONXPU_VRF_BUDGET", 24))
     is_use_mask_zero: bool = int(os.environ.get("TRITONXPU_IS_USE_MASK_ZERO", 0))
     extern_libs: dict = None
     is_sdnn: bool = False
@@ -143,6 +144,7 @@ class XPUOptions:
 
     # use_int4_w4a8
     use_int4_w4a8: bool = False
+    load_tile_size: int = 131072
 
     def __post_init__(self):
         default_libdir = Path(__file__).parent / f"xpu{self.arch}"
@@ -243,6 +245,7 @@ def make_ttxir(mod, metadata, opt):
         elem_bytes = int(os.environ.get("TRITONXPU_ELEMBYTES", 0))
         groups_per_cluster = metadata["groups_per_cluster"]
         unroll_num = metadata["unroll_num"]
+        vrf_budget = metadata.get("vrf_budget", 24)
         XPUBackend.buffer_len = xpu.get_buffer_len(mod, max_buffer_size, elem_bytes)
         # print(f"XPUBackend.buffer_len = {XPUBackend.buffer_len}")
         core_num = metadata["core_num"]
@@ -278,31 +281,48 @@ def make_ttxir(mod, metadata, opt):
                     xpu.passes.ttsdnnir.add_triton_convert_type_pass(pm, opt.arch, 2)
                 else:
                     xpu.passes.ttsdnnir.add_triton_convert_type_pass(pm, opt.arch, TTSDNN_F_MATMUL_FAST_MODE)
-            xpu.passes.ttsdnnir.add_convert_triton_to_tritonsdnn_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_convert_triton_to_tritonsdnn_pass(pm, opt.arch, opt.load_tile_size)
             passes.ttir.add_loop_aware_cse(pm)
             xpu.passes.ttsdnnir.add_linalg_to_tritonsdnn_pass(pm, opt.arch)
             passes.ttir.add_loop_aware_cse(pm)
             xpu.passes.ttsdnnir.add_tritonsdnn_legalize_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_merge_extern_ew_pass(pm)
             xpu.passes.ttsdnnir.add_tritonsdnn_combine_before_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_resolve_layout_conflict_pass(pm)
+            xpu.passes.ttsdnnir.add_tritonsdnn_hoist_ds_pass(pm)
+            xpu.passes.ttsdnnir.add_tritonsdnn_transpose_mma_pass(pm)
+            xpu.passes.ttsdnnir.add_tritonsdnn_transpose_ew_pass(pm)
+            xpu.passes.ttsdnnir.add_tritonsdnn_combine_before_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_eliminate_mma_acc_zero_pass(pm, opt.arch)
+            if opt.arch == 4:
+                xpu.passes.ttsdnnir.add_tritonsdnn_fuse_mma_vector_bias_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_optimize_rc_layout_pass(pm, opt.arch)
+            if opt.arch == 4:
+                xpu.passes.ttsdnnir.add_tritonsdnn_ewlite_scheduling_pass(pm, opt.arch)
+                if TTSDNN_F_DMA_MODE:
+                    xpu.passes.ttsdnnir.add_tritonsdnn_remove_ds_op_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_dsa_copy_pass(pm)
             xpu.passes.ttsdnnir.add_tritonsdnn_bufferize_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_mx_scale_layout_pass(pm)
             xpu.passes.ttsdnnir.add_tritonsdnn_combine_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_fuse_relu_activation_pass(pm, opt.arch)
             if opt.exp_range != ":0":
                 res = parse_floating_range_string(opt.exp_range)
                 xpu.passes.ttsdnnir.add_tritonsdnn_ew_act_table_pass(pm, res)
             else:
                 xpu.passes.ttsdnnir.add_tritonsdnn_ew_act_table_pass(pm, None)
             xpu.passes.ttsdnnir.add_tritonsdnn_loop_grid_pass(pm)
-            if opt.arch == 4 and TTSDNN_F_DMA_MODE:
-                xpu.passes.ttsdnnir.add_tritonsdnn_remove_ds_op_pass(pm, opt.arch)
-            if not TTSDNN_F_SINGLE_CORE_MODE:
-                xpu.passes.ttsdnnir.add_tritonsdnn_pipeline_pass(pm)
+            xpu.passes.ttsdnnir.add_tritonsdnn_hoist_loop_invariant_dma_pass(pm, opt.arch)
+            xpu.passes.ttsdnnir.add_tritonsdnn_pipeline_pass(pm, opt.arch)
             if opt.arch == 4 and TTSDNN_F_KILL_EW_FILL_MODE:
                 xpu.passes.ttsdnnir.add_tritonsdnn_kloop_acc_elimination_pass(pm)
             xpu.passes.ttsdnnir.add_tritonsdnn_multi_buffer_pass(pm, opt.arch, opt.num_stages)
+            xpu.passes.ttsdnnir.add_tritonsdnn_lower_rc_subview_pass(pm, opt.arch)
             passes.common.add_symbol_dce(pm)
             passes.common.add_canonicalizer(pm)
             passes.common.add_cse(pm)
         else:
+            xpu.passes.ttxpuir.add_tritonxpu_legalize_extern_ew_pass(pm)
             xpu.passes.ttxpuir.add_convert_triton_to_tritonxpu_pass(pm, opt.arch, XPUBackend.buffer_len, core_num)
             xpu.passes.ttxpuir.add_tritonxpu_print_pass(pm)
             xpu.passes.ttxpuir.add_tritonxpu_gm2lm_pass(pm, opt.arch, TTXPU_O_ATOMIC_SIM, opt.isClusterOneCoreActOnly,
@@ -313,15 +333,20 @@ def make_ttxir(mod, metadata, opt):
             passes.common.add_canonicalizer(pm)
             if TTXPU_F_DTYPE_CONVERT:
                 xpu.passes.ttxpuir.add_tritonxpu_dtype_convert_pass(pm, opt.arch)
+            xpu.passes.ttxpuir.add_tritonxpu_vectorizability_analysis_pass(pm, True, True)
             if not metadata["isCloseCoreTiling"]:
                 xpu.passes.ttxpuir.add_tritonxpu_core_tiling_pass(
                     pm, 0, XPUBackend.buffer_len, core_num, groups_per_cluster,
                     metadata["isAutoCoreTiling"]) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
             # xpu.passes.ttxpuir.add_tritonxpu_lm_to_sm_pass(pm)
             passes.common.add_cse(pm)
             if not metadata["isCloseOffsetAnalysis"]:
+                xpu.passes.ttxpuir.add_tritonxpu_scalar_analysis_pass(
+                    pm, False) if not TTXPU_O_CLOSE_OPT else None
                 xpu.passes.ttxpuir.add_tritonxpu_offset_state_pass(
                     pm, 0, XPUBackend.buffer_len, is_use_mask_zero) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
+                xpu.passes.ttxpuir.add_tritonxpu_scalar_analysis_pass(
+                    pm, True) if not TTXPU_O_CLOSE_OPT else None
             passes.common.add_canonicalizer(pm)
             xpu.passes.ttxpuir.add_tritonxpu_legalize_pass(pm, XPUBackend.buffer_len, core_num, groups_per_cluster,
                                                            is_use_mask_zero)
@@ -337,17 +362,21 @@ def make_ttxir(mod, metadata, opt):
             passes.common.add_canonicalizer(pm)
             if not metadata["isCloseVectorization"]:
                 compareFusion = int(os.environ.get("TRITONXPU_COMPARE_FUSION", 0))
+                xpu.passes.ttxpuir.add_tritonxpu_normalize_pass(
+                    pm, 0, compareFusion) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
+                xpu.passes.ttxpuir.add_tritonxpu_vectorizability_analysis_pass(pm, True, False)
                 xpu.passes.ttxpuir.add_tritonxpu_vectorize_pass(
                     pm, 0, compareFusion) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
             passes.common.add_canonicalizer(pm)
             xpu.passes.ttxpuir.add_tritonxpu_alloca_pass(pm, XPUBackend.buffer_len, core_num)
             if not metadata["isCloseMemoryAsync"]:
-                xpu.passes.ttxpuir.add_tritonxpu_memory_async_pass(pm,
-                                                                   0) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
+                xpu.passes.ttxpuir.add_tritonxpu_async_load_schedule_pass(
+                    pm, 0) if not TTXPU_O_CLOSE_OPT else None  # dumpFlag=0
             if not metadata["isCloseUnrollControl"]:
-                xpu.passes.ttxpuir.add_tritonxpu_unroll_control_pass(pm, XPUBackend.buffer_len, core_num,
-                                                                     is_use_mask_zero,
-                                                                     unroll_num) if not TTXPU_O_CLOSE_OPT else None
+                xpu.passes.ttxpuir.add_tritonxpu_tile_analysis_pass(pm, vrf_budget)
+                xpu.passes.ttxpuir.add_tritonxpu_unroll_control_pass(
+                    pm, XPUBackend.buffer_len, core_num, is_use_mask_zero, unroll_num,
+                    vrf_budget, False, -1) if not TTXPU_O_CLOSE_OPT else None
             xpu.passes.ttxpuir.add_tritonxpu_store_control_pass(pm) if not TTXPU_O_CLOSE_OPT else None
             if not TTXPU_F_OHTER_VALUE_SIM:
                 xpu.passes.ttxpuir.add_tritonxpu_other_sim_pass(pm, XPUBackend.buffer_len, core_num)
@@ -356,6 +385,8 @@ def make_ttxir(mod, metadata, opt):
             if not metadata["isCloseClusterLoopGrid"]:
                 xpu.passes.ttxpuir.add_tritonxpu_cf_to_scf_pass(pm)
                 xpu.passes.ttxpuir.add_tritonxpu_loop_grid_pass(pm)
+                if int(os.environ.get("TRITONXPU_LOOP_INVARIANT_STAGING", 0)):
+                    xpu.passes.ttxpuir.add_tritonxpu_loop_invariant_staging_pass(pm)
             passes.common.add_cse(pm)
             passes.common.add_licm(pm)
             passes.common.add_symbol_dce(pm)

third_party/xpu/backend/spec/PROVENANCE.md

@@ -0,0 +1,31 @@
+# FlagTree XPU — main-tree source overrides (whole-file vendored replacement)
+
+> 机制：`add_triton_object`（top `CMakeLists.txt`）对每个源文件检查本目录下是否有同「项目相对路径」的覆盖文件；有则 XPU build 编本副本，非 XPU 编主树原版。EXISTS 门控，非 XPU / 无覆盖零影响。
+> 该机制用于保持共享主树源码不含 XPU 专属语义，同时允许后端维护完整的源文件替换。
+
+## 为什么用整文件替换而非就地改主树
+主树为所有后端共享，禁止就地改（Q0b）。这些副本**以 FlagTree 主树文件为底**（保留 `flagtree_hints` 等本地适配）+ 叠加 XPU 专属改动；**不以 internal 版为底**（实测 `Ops.cpp` 与 internal 有 162 行冲突 gap，如 `LoadOp::build` 的 `flagtree_hints` vs `offsetState/syncMode`）。
+
+## Provenance（派生基线，用于反 drift）
+- 派生自 FlagTree `main` 提交：`40a57023ddbd075c732eefb6886d6d3a152a181e`
+- 迁移目标 internal Triton：`triton_3.6` @ `d3c64dd65e401239608164d6be4d893b261b4869`
+- 生成时间：2026-07-10；2026-07-16 对最新 `main` 复核
+
+上述五个主树源文件在原始派生提交 `b26ec15c65bea71f4e011c7626f7de9f5146937e`
+与当前 `main` 基线之间内容一致；本次更新提交号用于准确反映 PR 的实际基线。
+
+## 覆盖文件清单（副本 = 上述 FlagTree 文件 + 下列改动；括号为相对 pristine 的变动行）
+| 覆盖路径 | 叠加的 XPU 改动 | 变动行 |
+|---|---|---|
+| `lib/Dialect/Triton/IR/Ops.cpp` | DotOp::verify 放宽 i8×i4(w4a8) · ReshapeOp::fold 去 `!getAllowReorder()` · BitcastOp::verify vector↔vector | 46 |
+| `lib/Dialect/Triton/IR/Traits.cpp` | verifyTensorSize：允许非 pow2 元素数；verifyTensorLayouts：SliceEncoding 递归解析到 triton_xpu parent | 38 |
+| `lib/Conversion/TritonGPUToLLVM/ViewOpToLLVM.cpp` | ArithConstantSplatOpConversion splat guard | 2 |
+| `lib/Dialect/TritonGPU/IR/Dialect.cpp` | ceil-offset：`getTotalElemsPerThread`/`getElemsPerThread(Attribute,shape)` 对 XPU 层（ClusterLayoutAttr / XPU-backed SliceEncoding）走 ceil-based 派发（新增 `isXPUBackedLayout` + `#include TritonXPU/IR/Dialect.h` + 文件内前置声明 `getElemsPerThread(Attribute,shape)`）；绕开泛型 LinearEncoding 的 pow2 断言 | 59 |
+| `include/triton/Dialect/Triton/IR/TritonTypes.td` | 将 `TT_Vector`/`TT_VectorTensor` 组成的 `TT_VectorLike` 加入 `TT_Type`，允许 XPU vectorize 后的 `tt.extern_elementwise` 等主 Triton op 接受 vector-like operand/result | 6 |
+
+> `Dialect.cpp` 关键点：internal 在主树 header `TritonGPU/IR/Dialect.h` 加了 `getElemsPerThread(Attribute,ArrayRef)` 声明；FlagTree 该 header **无**此声明（Q0b 不改共享 header），故 `getTotalElemsPerThread` 里 line~112 的 `getElemsPerThread(layout,shape)` 会误配到 `getElemsPerThread(Type)` 报 `Attribute→Type` 转换错。修法：在本 vendored 副本内 `namespace mlir::triton::gpu` 前置声明该 overload（不动主树 header）。已经 XTDK clang22 实测编译+链接通过。
+
+> 注：`Dialect.cpp` 副本保留了 pristine 顶部的 `flagtree_spec.h` 原生守卫（`#if __has_include("flagtree_spec.h")` / `#ifndef FLAGTREE_SPEC_Dialect_TritonGPU_IR_Dialect`）。XPU 未提供 `third_party/xpu/backend/spec/include/flagtree_spec.h`，故 `__has_include` 为假、宏未定义、整个 body 正常编译——守卫无副作用。新增的 `#include "triton/Dialect/TritonXPU/IR/Dialect.h"` 经 XPU 后端 include dir（`third_party/xpu/include`）解析，仅存在于本副本。
+
+## 维护须知（drift）
+FlagTree 主树每次升级上述任一文件，**必须**用新版主树文件重做副本（以新主树为底重叠 XPU 改动），并更新本文件的派生提交号。否则 XPU 编译的是过期主树逻辑。