[NFI] Clean up useTileLoadLinearLayout (#5232)

This PR removes the `useTileLoadLinearLayout` configuration option and
related infrastructure from the Intel Triton GPU backend. The change
simplifies the codebase by removing a conditional feature flag that was
previously enabled by default.

Signed-off-by: Whitney Tsang <[email protected]>

Author: whitneywhtsang

python/triton/knobs.py

@@ -549,7 +549,6 @@ class intel_knobs(base_knobs):
 
     dump_shader_info: env_bool = env_bool("TRITON_INTEL_ENABLE_IGC_SHADER_DUMP", False)
     gen_native_code: env_bool = env_bool("TRITON_XPU_GEN_NATIVE_CODE", False)
-    tile_load_ll: env_bool = env_bool("TRITON_XPU_ENABLE_TILE_LOAD_LINEAR_LAYOUT", True)
     opt_reduction_locality: env_bool = env_bool("TRITON_INTEL_OPTIMIZE_REDUCTION_LOCALITY", False)
     disable_igc_opt: env_bool = env_bool("TRITON_INTEL_DISABLE_IGC_OPT", False)
 

third_party/intel/backend/compiler.py

@@ -40,7 +40,6 @@ class XPUOptions:
     backend_name: str = 'intel'
     sanitize_overflow: bool = False
     generate_native_code: bool = False
-    enable_tile_load_linear_layout: bool = True
     arch: str = None
     # FIXME: enable for XPU: https://github.com/intel/intel-xpu-backend-for-triton/issues/4954
     instrumentation_mode: str = ""
@@ -133,7 +132,6 @@ def parse_target(self, tgt_prop) -> dict:
     def parse_options(self, opts) -> Any:
         args = {k: opts[k] for k in XPUOptions.__dataclass_fields__.keys() if k in opts}
         args["allow_fp8e4nv"] = True
-        args["enable_tile_load_linear_layout"] = knobs.intel.tile_load_ll
         return XPUOptions(**args)
 
     def pack_metadata(self, metadata):
@@ -298,7 +296,7 @@ def make_llir(src, metadata, options):
         # instrumentation point here so we can override IRs above (e.g., ttir and ttgir)
         if XPUBackend.instrumentation:
             XPUBackend.instrumentation.patch("ttgpuir_to_llvmir", pm, mod.context)
-        intel.passes.ttgpuir.add_to_llvmir(pm, options.enable_tile_load_linear_layout)
+        intel.passes.ttgpuir.add_to_llvmir(pm)
         intel.passes.ttgpuir.add_gen_to_llvm(pm)
         passes.common.add_canonicalizer(pm)
         intel.passes.ttgpuir.add_rewrite_stack_ptr(pm)

third_party/intel/include/TritonIntelGPUToLLVM/Passes.td

@@ -21,11 +21,6 @@ def ConvertTritonIntelGPUToLLVM
                            "mlir::triton::TritonDialect",
                            "mlir::triton::gpu::TritonGPUDialect",
                            "mlir::triton::TritonGEN::TritonGENDialect"];
-  let options = [
-    Option<"useTileLoadLinearLayout", "use_tile_load_linear_layout",
-           "bool", /*default*/"true",
-           "Use linear layouts to generate the tile load sizes and offsets">
-  ];
 }
 
 #endif // TRITONINTELGPU_CONVERSION_PASSES

third_party/intel/lib/TritonIntelGPUToLLVM/LoadStoreOpToLLVM.cpp

@@ -1012,10 +1012,9 @@ struct LoadOpToBlockIOConversion
   LoadOpToBlockIOConversion(
       LLVMTypeConverter &converter, const triton::intel::TargetInfo &targetInfo,
       const triton::intel::ModuleAxisInfoAnalysis &axisAnalysisPass,
-      PatternBenefit benefit, bool useTileLoadLinearLayout)
+      PatternBenefit benefit)
       : ConvertTritonGPUOpToLLVMPattern<triton::LoadOp>(converter, benefit),
-        BlockIOConversionBase(targetInfo, axisAnalysisPass),
-        useTileLoadLinearLayout(useTileLoadLinearLayout) {}
+        BlockIOConversionBase(targetInfo, axisAnalysisPass) {}
 
   LogicalResult
   rewriteTensorPointerLoad(triton::LoadOp op, OpAdaptor adaptor,
@@ -1540,18 +1539,11 @@ struct LoadOpToBlockIOConversion
     // Disable building the load layout if we are not going to use it. Building
     // the layout manually can cause an error which would abort the pass
     // pipeline and block us from getting debug info.
-    if (useTileLoadLinearLayout) {
-      // add the bases to the map and replace the tile layout with the new
-      // layout
-      bases[kLoad] = newLoadBases;
-      tileLayout = LinearLayout(bases, outDims,
-                                /*requiredSurjective=*/false);
-    } else {
-      // when linear layouts are disabled generate a single load, so we can have
-      // some reference for linear layout output without generating a layout
-      // that could abort the pass pipeline
-      tileLayout *= LinearLayout::identity1D(1, kLoad, dimOuterStr);
-    }
+    // add the bases to the map and replace the tile layout with the new
+    // layout
+    bases[kLoad] = newLoadBases;
+    tileLayout = LinearLayout(bases, outDims,
+                              /*requiredSurjective=*/false);
 
     LLVM_DEBUG({
       llvm::dbgs() << "Block load tile layout after adding loads: "
@@ -1657,33 +1649,19 @@ struct LoadOpToBlockIOConversion
               llvm::dbgs() << "y offset: "
                            << outer * repOuterStride + rep * repStride << "\n";
             });
-            if (useTileLoadLinearLayout) {
-              offsetY = b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
-                              b.i32_val(layoutOffsetY));
-              offsetX = b.i32_val(layoutOffsetX);
-            } else {
-              offsetY =
-                  b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
-                        b.i32_val(outer * repOuterStride + rep * repStride));
-              offsetX = b.i32_val(k * repKStride);
-            }
+            offsetY = b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
+                            b.i32_val(layoutOffsetY));
+            offsetX = b.i32_val(layoutOffsetX);
           } break;
           case DpasEncodingAttr::OpIdx::OperandB: {
             LLVM_DEBUG({
               llvm::dbgs() << "x offset: "
                            << outer * repOuterStride + rep * repStride << "\n";
               llvm::dbgs() << "y offset: " << k * repKStride << "\n";
             });
-            if (useTileLoadLinearLayout) {
-              offsetX = b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
-                              b.i32_val(layoutOffsetX));
-              offsetY = b.i32_val(layoutOffsetY);
-            } else {
-              offsetX =
-                  b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
-                        b.i32_val(outer * repOuterStride + rep * repStride));
-              offsetY = b.i32_val(k * repKStride);
-            }
+            offsetX = b.add(b.mul(outerDimWarpId, b.i32_val(warpOuterStride)),
+                            b.i32_val(layoutOffsetX));
+            offsetY = b.i32_val(layoutOffsetY);
           } break;
           case DpasEncodingAttr::OpIdx::OperandC: {
             llvm_unreachable("unexpected OpIdx::OperandC");
@@ -3164,9 +3142,6 @@ struct LoadOpToBlockIOConversion
 
     return success();
   }
-
-private:
-  bool useTileLoadLinearLayout;
 };
 
 struct LoadOpConversion : public ConvertOpToLLVMPattern<triton::LoadOp>,
@@ -4258,14 +4233,11 @@ void mlir::triton::intel::populateLoadStoreOpToLLVMPatterns(
     LLVMTypeConverter &typeConverter, const TargetInfo &targetInfo,
     RewritePatternSet &patterns,
     const intel::ModuleAxisInfoAnalysis &axisInfoAnalysis,
-    PatternBenefit benefit, bool useTileLoadLinearLayout) {
+    PatternBenefit benefit) {
   patterns.add<AtomicCASOpConversion, AtomicRMWOpConversion, LoadOpConversion,
                StoreOpConversion, PrefetchOpConversion>(
       typeConverter, targetInfo, axisInfoAnalysis, benefit);
   // BlockIO is more efficient than gather load or scatter store.
-  patterns.add<LoadOpToBlockIOConversion>(
-      typeConverter, targetInfo, axisInfoAnalysis, benefit.getBenefit() + 2,
-      useTileLoadLinearLayout);
-  patterns.add<StoreOpToBlockIOConversion>(
+  patterns.add<LoadOpToBlockIOConversion, StoreOpToBlockIOConversion>(
       typeConverter, targetInfo, axisInfoAnalysis, benefit.getBenefit() + 2);
 }

third_party/intel/lib/TritonIntelGPUToLLVM/PatternTritonGPUOpToLLVM.h

@@ -61,7 +61,7 @@ void populateFp4ToFpToLLVMPatterns(LLVMTypeConverter &typeConverter,
 void populateLoadStoreOpToLLVMPatterns(
     LLVMTypeConverter &typeConverter, const TargetInfo &targetInfo,
     RewritePatternSet &patterns, const ModuleAxisInfoAnalysis &axisInfoAnalysis,
-    PatternBenefit benefit, bool useTileLoadLinearLayout);
+    PatternBenefit benefit);
 
 void populateTensorPtrOpsToLLVMPatterns(LLVMTypeConverter &typeConverter,
                                         RewritePatternSet &patterns,

third_party/intel/lib/TritonIntelGPUToLLVM/PipelineManager.h

@@ -184,9 +184,8 @@ struct FuncOpConversion : public ConvertOpToLLVMPattern<triton::FuncOp> {
 /// block pointers or not.
 class TritonGPUToLLVMPipelineManager {
 public:
-  TritonGPUToLLVMPipelineManager(ModuleOp &mod, MLIRContext *ctx,
-                                 bool useTileLoadLinearLayout)
-      : mod(mod), ctx(ctx), useTileLoadLinearLayout(useTileLoadLinearLayout) {}
+  TritonGPUToLLVMPipelineManager(ModuleOp &mod, MLIRContext *ctx)
+      : mod(mod), ctx(ctx) {}
 
   /// Populate the conversion pipeline for function operations.
   void populateFunctionConversionPatterns(
@@ -213,9 +212,8 @@ class TritonGPUToLLVMPipelineManager {
     intel::populateDotOpToLLVMPatterns(typeConverter, patterns, benefit);
     intel::populateElementwiseOpToLLVMPatterns(
         typeConverter, patterns, axisInfoAnalysis, targetInfo, benefit);
-    intel::populateLoadStoreOpToLLVMPatterns(typeConverter, targetInfo,
-                                             patterns, axisInfoAnalysis,
-                                             benefit, useTileLoadLinearLayout);
+    intel::populateLoadStoreOpToLLVMPatterns(
+        typeConverter, targetInfo, patterns, axisInfoAnalysis, benefit);
     intel::populateReduceOpToLLVMPatterns(typeConverter, patterns, targetInfo,
                                           benefit);
     mlir::triton::populateScanOpToLLVMPatterns(typeConverter, patterns,
@@ -259,8 +257,6 @@ class TritonGPUToLLVMPipelineManager {
 private:
   ModuleOp &mod;
   MLIRContext *ctx;
-
-  bool useTileLoadLinearLayout = true;
 };
 
 } // namespace mlir::triton::intel

third_party/intel/lib/TritonIntelGPUToLLVM/TritonGPUToLLVM.cpp

@@ -60,10 +60,6 @@ struct ConvertTritonGPUToLLVM
     : public triton::gpu::intel::impl::ConvertTritonIntelGPUToLLVMBase<
           ConvertTritonGPUToLLVM> {
   using ConvertTritonIntelGPUToLLVMBase::ConvertTritonIntelGPUToLLVMBase;
-  ConvertTritonGPUToLLVM() = default;
-  ConvertTritonGPUToLLVM(bool useTileLoadLinearLayout) {
-    this->useTileLoadLinearLayout = useTileLoadLinearLayout;
-  }
 
   void getDependentDialects(DialectRegistry &registry) const override {
     registry.insert<LLVM::LLVMDialect, TritonGEN::TritonGENDialect,
@@ -75,7 +71,7 @@ struct ConvertTritonGPUToLLVM
     ModuleOp mod = getOperation();
 
     mlir::triton::intel::TritonGPUToLLVMPipelineManager pipelineManager(
-        mod, context, useTileLoadLinearLayout);
+        mod, context);
     mlir::LowerToLLVMOptions option(context);
     auto targetInfo = mlir::triton::intel::createTargetInfo(mod);
     TritonIntelGPUToLLVMTypeConverter typeConverter(context, option,

third_party/intel/triton_xpu.cc

@@ -60,8 +60,8 @@ void init_triton_intel_passes_ttir(py::module &&m) {
 }
 
 void init_triton_intel_passes_ttgpuir(py::module &&m) {
-  ADD_PASS_OPTION_WRAPPER_1(
-      "add_to_llvmir", gpu::intel::createConvertTritonIntelGPUToLLVM, bool);
+  ADD_PASS_WRAPPER_0("add_to_llvmir",
+                     gpu::intel::createConvertTritonIntelGPUToLLVM);
   ADD_PASS_WRAPPER_0("add_gen_to_llvm", createConvertTritonGENToLLVM);
   ADD_PASS_WRAPPER_0("add_accelerate_matmul",
                      gpu::intel::createTritonIntelGPUAccelerateMatmul);