Integrate Triton up to [6ea516a6e4](https://github.com/triton-lang/tr…

third_party/triton/common/mma_limit_pred.patch

@@ -19,5 +19,4 @@ fusions reported.
 +
    return (op->hasTrait<OpTrait::Elementwise>() && isMemoryEffectFree(op)) ||
           isView(op) ||
-          isa<Fp4ToFpOp, LoadOp, DescriptorLoadOp, BroadcastOp, ConvertLayoutOp>(
-
+          isa<Fp4ToFpOp, LoadOp, DescriptorLoadLikeOpInterface, BroadcastOp,

third_party/triton/common/no_accelerate_through_broadcast.patch

@@ -2,7 +2,7 @@ Fix for b/405045790.
 
 --- a/lib/Dialect/TritonGPU/Transforms/AccelerateMatmul.cpp
 +++ b/lib/Dialect/TritonGPU/Transforms/AccelerateMatmul.cpp
-@@ -272,10 +272,13 @@ static bool bwdFilter(Operation *op) {
+@@ -272,9 +272,12 @@ static bool bwdFilter(Operation *op) {
        return false;
    }
 
@@ -11,11 +11,8 @@ Fix for b/405045790.
 +  // small. This is just a heuristic to avoid a regression.
    return (op->hasTrait<OpTrait::Elementwise>() && isMemoryEffectFree(op)) ||
           isView(op) ||
--         isa<Fp4ToFpOp, LoadOp, DescriptorLoadOp, BroadcastOp, ConvertLayoutOp>(
--             op);
-+         isa<Fp4ToFpOp, LoadOp, DescriptorLoadOp, /*BroadcastOp,*/
-+             ConvertLayoutOp>(op);
+-         isa<Fp4ToFpOp, LoadOp, DescriptorLoadLikeOpInterface, BroadcastOp,
++         isa<Fp4ToFpOp, LoadOp, DescriptorLoadLikeOpInterface, /*BroadcastOp,*/
+              ConvertLayoutOp>(op);
  }
 
- // Finds the bitwidth with which the value x is loaded
-

third_party/triton/common/series.bzl

@@ -28,10 +28,7 @@ common_patch_list = [
     "//third_party/triton:common/avoid-0fc-mid-ptwas-128.patch",
     "//third_party/triton:common/wgmma_pipeline_fix.patch",
     "//third_party/triton:common/nvdisasm_bin_path.patch",
-    "//third_party/triton:common/llvm_cl887809531.patch",
-    "//third_party/triton:common/llvm_cl893899241.patch",
     "//third_party/triton:common/stage_and_cluster_map.patch",
-    "//third_party/triton:common/llvm_cl895542516.patch",
     "//third_party/triton:common/assert_fail.patch",
     "//third_party/triton:common/llvm_cl900404532.patch",
     # Add new patches just above this line

third_party/triton/workspace.bzl

@@ -7,8 +7,8 @@ load("//third_party/triton:oss_only/series.bzl", "oss_only_patch_list")
 def repo():
     """Imports Triton."""
 
-    TRITON_COMMIT = "6a0b5461eb439bdb3ef0c721837c64dc2c345fcb"
-    TRITON_SHA256 = "becf73e91a770d424ba98b49169ebd1cc5aa93c8077953021856d989305a5fd2"
+    TRITON_COMMIT = "6ea516a6e4cb1878fb6969d222e23294712b3955"
+    TRITON_SHA256 = "ae490e452ad73ed6d0ba8ee54fc96f72c265b627704ad8887d961a46a7888fac"
     tf_http_archive(
         name = "triton",
         sha256 = TRITON_SHA256,

xla/backends/gpu/codegen/triton/transforms/triton_xla_squeeze_dims_pass.cc

@@ -132,7 +132,8 @@ RankedTensorType SqueezeTensorType(RankedTensorType type,
         cast<DialectInferLayoutInterface>(&encoding.getDialect());
     [[maybe_unused]] LogicalResult result =
         inferLayoutInterface->inferReshapeOpEncoding(
-            type.getShape(), encoding, shape, encoding, /*loc=*/std::nullopt);
+            type.getShape(), encoding, shape, encoding,
+            /*allowReorder=*/false, /*loc=*/std::nullopt);
     CHECK(succeeded(result));
   }
   return RankedTensorType::get(shape, type.getElementType(), encoding);