From a357ad2e3a2ca96ca90c1842cea9a27453e6ae8c Mon Sep 17 00:00:00 2001 From: Alina Sbirlea Date: Thu, 2 Apr 2026 23:21:38 -0700 Subject: [PATCH] Integrate LLVM at llvm/llvm-project@c329cc59d91b Updates LLVM usage to match [c329cc59d91b](https://github.com/llvm/llvm-project/commit/c329cc59d91b) PiperOrigin-RevId: 893905698 --- third_party/llvm/workspace.bzl | 4 +- third_party/shardy/temporary.patch | 795 +++++++++++++++++- third_party/shardy/workspace.bzl | 4 +- third_party/stablehlo/temporary.patch | 728 ++++++++++++++++ .../triton/common/llvm_cl893905698.patch | 146 ++++ third_party/triton/common/series.bzl | 1 + .../gpu/codegen/emitters/ir/xla_gpu_ops.cc | 2 +- xla/mlir_hlo/mhlo/IR/hlo_ops.cc | 176 ++-- 8 files changed, 1768 insertions(+), 88 deletions(-) create mode 100644 third_party/triton/common/llvm_cl893905698.patch diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl index 21ba653162cd6..6b8196b016e3b 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" - LLVM_COMMIT = "7ccd92e5e6e5c622b2b571d396fff9016241a8f1" - LLVM_SHA256 = "0365ff939e7ba4876437f45318ac4140e7b278fd8b2c5dc1e78c4dd04104c831" + LLVM_COMMIT = "c329cc59d91b3f9774383e5e933d812539f17d1e" + LLVM_SHA256 = "8c4e9bc4e6f99d922ecd2eb9946fe92975704bb5187d609eec33e8087fbd5d50" tf_http_archive( name = name, diff --git a/third_party/shardy/temporary.patch b/third_party/shardy/temporary.patch index 3b3f53353fb31..a352ec2f446db 100644 --- a/third_party/shardy/temporary.patch +++ b/third_party/shardy/temporary.patch @@ -1,15 +1,800 @@ +diff --git a/shardy/dialect/mpmd/ir/dialect.cc b/shardy/dialect/mpmd/ir/dialect.cc +index bfd82c0..0e00341 100644 +--- a/shardy/dialect/mpmd/ir/dialect.cc ++++ b/shardy/dialect/mpmd/ir/dialect.cc +@@ -1171,7 +1171,8 @@ LogicalResult UnassignOp::verify() { + + LogicalResult UnassignOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + UnassignOp::Adaptor adaptor(operands, attributes, properties, regions); + inferredReturnShapes.emplace_back( +diff --git a/shardy/dialect/sdy/ir/dialect.cc b/shardy/dialect/sdy/ir/dialect.cc +index 798e9c0..3c398a5 100644 +--- a/shardy/dialect/sdy/ir/dialect.cc ++++ b/shardy/dialect/sdy/ir/dialect.cc +@@ -1369,7 +1369,7 @@ bool ManualComputationOp::shouldKeepEdgeOwnerShardingsDivisible() { + LogicalResult ShardingGroupOp::inferReturnTypes(MLIRContext*, + std::optional, + ValueRange, DictionaryAttr, +- OpaqueProperties, RegionRange, ++ mlir::PropertyRef, RegionRange, + SmallVectorImpl&) { + return success(); + } +@@ -1384,8 +1384,8 @@ bool ConstantOp::isCompatibleReturnTypes(TypeRange l, TypeRange r) { + + LogicalResult ConstantOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl& inferredReturnTypes) { ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + ConstantOpAdaptor adaptor(operands, attributes, properties); + inferredReturnTypes.push_back(adaptor.getValue().getType()); + return success(); +@@ -1536,8 +1536,8 @@ Value NamedComputationOp::getEdgeOwnerFromSource(OpOperand& source) { + + LogicalResult NamedComputationOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl& inferredReturnTypes) { ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + NamedComputationOpAdaptor adaptor(operands, attributes, properties, regions); + llvm::copy(getBodyTerminatorOperands(adaptor).getTypes(), + std::back_inserter(inferredReturnTypes)); diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl -index 472beb1..21ba653 100644 +index 21ba653..6b8196b 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" -- LLVM_COMMIT = "9a0b003dde83d46f2ed6d95d85d1d9de3c1fe908" -- LLVM_SHA256 = "54ce29dc05966cb898dbb98a426b78e261d2830d2125977bc4a938ea53a7c05e" -+ LLVM_COMMIT = "7ccd92e5e6e5c622b2b571d396fff9016241a8f1" -+ LLVM_SHA256 = "0365ff939e7ba4876437f45318ac4140e7b278fd8b2c5dc1e78c4dd04104c831" +- LLVM_COMMIT = "7ccd92e5e6e5c622b2b571d396fff9016241a8f1" +- LLVM_SHA256 = "0365ff939e7ba4876437f45318ac4140e7b278fd8b2c5dc1e78c4dd04104c831" ++ LLVM_COMMIT = "c329cc59d91b3f9774383e5e933d812539f17d1e" ++ LLVM_SHA256 = "8c4e9bc4e6f99d922ecd2eb9946fe92975704bb5187d609eec33e8087fbd5d50" tf_http_archive( name = name, +diff --git a/third_party/stablehlo/temporary.patch b/third_party/stablehlo/temporary.patch +index f64aa01..8341fcc 100755 +--- a/third_party/stablehlo/temporary.patch ++++ b/third_party/stablehlo/temporary.patch +@@ -1,3 +1,731 @@ ++diff --ruN a/stablehlo/stablehlo/dialect/Base.h b/stablehlo/stablehlo/dialect/Base.h ++--- stablehlo/stablehlo/dialect/Base.h +++++ stablehlo/stablehlo/dialect/Base.h ++@@ -456,7 +456,7 @@ ++ static LogicalResult inferReturnTypes( ++ MLIRContext* /*context*/, std::optional location, ++ ValueRange operands, DictionaryAttr /*attributes*/, ++- OpaqueProperties /*properties*/, RegionRange /*regions*/, +++ PropertyRef /*properties*/, RegionRange /*regions*/, ++ SmallVectorImpl& inferredReturnTypes) { ++ // TODO(b/231358795): Review the use of InferTypeOpInterface for ops that ++ // support quantization or sparsity. ++@@ -478,7 +478,7 @@ ++ static LogicalResult inferReturnTypeComponentsFromOperands( ++ MLIRContext* context, std::optional location, ++ ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SmallVector inferredReturnTypes; ++ if (failed(inferReturnTypes(context, location, operands.getValues(), ++diff --ruN a/stablehlo/stablehlo/dialect/ChloOps.cpp b/stablehlo/stablehlo/dialect/ChloOps.cpp ++--- stablehlo/stablehlo/dialect/ChloOps.cpp +++++ stablehlo/stablehlo/dialect/ChloOps.cpp ++@@ -71,7 +71,7 @@ ++ LogicalResult Op::inferReturnTypeComponents( \ ++ MLIRContext* context, std::optional location, \ ++ ValueShapeRange operands, DictionaryAttr attributes, \ ++- OpaqueProperties properties, RegionRange regions, \ +++ PropertyRef properties, RegionRange regions, \ ++ SmallVectorImpl& inferredReturnShapes) { \ ++ return inferReturnTypeComponentsFromOperands( \ ++ context, location, operands, attributes, properties, regions, \ ++@@ -199,7 +199,7 @@ ++ ++ LogicalResult InferBroadcastBinaryOpReturnTypeComponents( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, +++ DictionaryAttr attributes, PropertyRef properties, ++ std::optional> broadcastDimensions, Type elementType, ++ SmallVectorImpl& inferredReturnShapes) { ++ // Handle unranked. ++@@ -252,8 +252,8 @@ ++ ++ LogicalResult BroadcastComplexOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ ShapedType lhsType = cast(operands[0].getType()); ++ Type elementType = ComplexType::get(lhsType.getElementType()); ++@@ -287,8 +287,8 @@ ++ ++ LogicalResult BroadcastCompareOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ Type elementType = IntegerType::get(context, 1); ++ Adaptor adaptor(operands, attributes, properties, regions); ++@@ -317,7 +317,7 @@ ++ ++ LogicalResult IsInfOp::inferReturnTypes( ++ MLIRContext* /*ctx*/, std::optional, ValueRange operands, ++- DictionaryAttr, OpaqueProperties, RegionRange, +++ DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); ++ return success(); ++@@ -329,7 +329,7 @@ ++ ++ LogicalResult IsNegInfOp::inferReturnTypes( ++ MLIRContext* /*ctx*/, std::optional, ValueRange operands, ++- DictionaryAttr, OpaqueProperties, RegionRange, +++ DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); ++ return success(); ++@@ -341,7 +341,7 @@ ++ ++ LogicalResult IsPosInfOp::inferReturnTypes( ++ MLIRContext* /*ctx*/, std::optional, ValueRange operands, ++- DictionaryAttr, OpaqueProperties, RegionRange, +++ DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); ++ return success(); ++@@ -355,7 +355,7 @@ ++ LogicalResult Op::inferReturnTypeComponents( \ ++ MLIRContext* context, std::optional location, \ ++ ValueShapeRange operands, DictionaryAttr attributes, \ ++- OpaqueProperties properties, RegionRange regions, \ +++ PropertyRef properties, RegionRange regions, \ ++ SmallVectorImpl& inferredReturnShapes) { \ ++ Adaptor adaptor(operands.getValues(), attributes, properties, regions); \ ++ return InferBroadcastBinaryOpReturnTypeComponents( \ ++@@ -400,8 +400,8 @@ ++ ++ LogicalResult ConstantLikeOp::inferReturnTypeComponents( ++ MLIRContext* /*context*/, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange /*regions*/, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange /*regions*/, ++ SmallVectorImpl& inferredReturnShapes) { ++ ConstantLikeOp::Adaptor op(operands, attributes, properties); ++ if (failed(op.verify(location.value()))) return failure(); ++@@ -437,7 +437,7 @@ ++ ++ LogicalResult BroadcastSelectOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr, OpaqueProperties, RegionRange, +++ DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnShapes) { ++ BroadcastSelectOp::Adaptor op(operands.getValues()); ++ auto onTrueType = cast(op.getOnTrue().getType()); ++@@ -757,8 +757,8 @@ ++ ++ LogicalResult TopKOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ TopKOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferTopKOp(location, adaptor.getOperand(), adaptor.getK(), ++@@ -773,7 +773,7 @@ ++ ++ LogicalResult ConstantOp::inferReturnTypes( ++ MLIRContext*, std::optional, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ Adaptor adaptor(operands, attributes, properties, regions); ++ Type type = cast(adaptor.getValueAttr()).getType(); ++@@ -787,7 +787,7 @@ ++ ++ LogicalResult ScanOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ ScanOp::Adaptor adaptor(operands, attributes, properties, regions); ++ if (regions.empty() || regions.front()->empty()) { ++diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp ++--- stablehlo/stablehlo/dialect/StablehloOps.cpp +++++ stablehlo/stablehlo/dialect/StablehloOps.cpp ++@@ -216,7 +216,7 @@ ++ LogicalResult Op::inferReturnTypeComponents( \ ++ MLIRContext* context, std::optional location, \ ++ ValueShapeRange operands, DictionaryAttr attributes, \ ++- OpaqueProperties properties, RegionRange regions, \ +++ PropertyRef properties, RegionRange regions, \ ++ SmallVectorImpl& inferredReturnShapes) { \ ++ return inferReturnTypeComponentsFromOperands( \ ++ context, location, operands, attributes, properties, regions, \ ++@@ -267,8 +267,8 @@ ++ ++ LogicalResult AddOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SmallVector inferredReturnTypes; ++ if (failed(inferReturnTypes(context, location, operands.getValues(), ++@@ -293,7 +293,7 @@ ++ ++ LogicalResult AfterAllOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ return hlo::inferAfterAllOp(getStablehloDialect(context), location, ++ inferredReturnTypes); ++@@ -357,7 +357,7 @@ ++ ++ LogicalResult ConstantOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ ConstantOpAdaptor adaptor(operands, attributes, properties); ++ return hlo::inferConstantOp(location, adaptor.getValue(), ++@@ -391,7 +391,7 @@ ++ ++ LogicalResult CreateTokenOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ return hlo::inferCreateTokenOp(getStablehloDialect(context), location, ++ inferredReturnTypes); ++@@ -729,7 +729,7 @@ ++ ++ LogicalResult CholeskyOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ CholeskyOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferCholeskyOp(location, adaptor.getA(), inferredReturnShapes); ++@@ -746,7 +746,7 @@ ++ ++ LogicalResult DotOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ DotOp::Adaptor adaptor(operands, attributes, properties, regions); ++ auto lhsType = mlir::cast(adaptor.getLhs().getType()); ++@@ -885,7 +885,7 @@ ++ ++ LogicalResult DotGeneralOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ DotGeneralOp::Adaptor adaptor(operands, attributes, properties, regions); ++ LLVM_DEBUG(llvm::dbgs() << "DotGeneralOp::inferReturnTypeComponents\n"); ++@@ -1183,7 +1183,7 @@ ++ ++ LogicalResult FftOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ FftOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferFftOp(location, adaptor.getOperand(), ++@@ -1304,8 +1304,8 @@ ++ ++ LogicalResult GatherOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ GatherOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferGatherOp( ++@@ -1345,8 +1345,8 @@ ++ ++ LogicalResult DynamicGatherOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ DynamicGatherOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferDynamicGatherOp( ++@@ -1365,7 +1365,7 @@ ++ ++ LogicalResult GetDimensionSizeOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ GetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1435,7 +1435,7 @@ ++ ++ LogicalResult DynamicUpdateSliceOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ DynamicUpdateSliceOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1450,7 +1450,7 @@ ++ ++ LogicalResult AbsOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ AbsOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferAbsOp(location, adaptor.getOperand(), inferredReturnTypes); ++@@ -1470,7 +1470,7 @@ ++ ++ LogicalResult CollectiveBroadcastOp::inferReturnTypes( ++ MLIRContext* /*context*/, std::optional location, ++- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, +++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ++ CollectiveBroadcastOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1480,8 +1480,8 @@ ++ ++ LogicalResult CollectiveBroadcastOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SmallVector inferredReturnTypes; ++ CollectiveBroadcastOp::Adaptor adaptor(operands, attributes, properties, ++@@ -1513,7 +1513,7 @@ ++ ++ LogicalResult CollectivePermuteOp::inferReturnTypes( ++ MLIRContext* /*context*/, std::optional location, ++- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, +++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ++ CollectivePermuteOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1523,8 +1523,8 @@ ++ ++ LogicalResult CollectivePermuteOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SmallVector inferredReturnTypes; ++ CollectivePermuteOp::Adaptor adaptor(operands, attributes, properties, ++@@ -1634,7 +1634,7 @@ ++ ++ LogicalResult AllToAllOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ AllToAllOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferAllToAllOp( ++@@ -1736,7 +1736,7 @@ ++ ++ LogicalResult AllReduceOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ AllReduceOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferAllReduceOp(location, adaptor.getOperands(), ++@@ -1752,7 +1752,7 @@ ++ // async_start references stablehlo.futures. ++ LogicalResult AsyncStartOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ if (regions.empty()) { ++ return emitOptionalError( ++@@ -1820,7 +1820,7 @@ ++ // TODO: Move to TypeInference.h and TypeInference.cpp. See TODO above. ++ LogicalResult AsyncDoneOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ auto futureType = dyn_cast(operands[0].getType()); ++ if (!futureType) { ++@@ -1845,8 +1845,8 @@ ++ ++ LogicalResult BatchNormGradOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ BatchNormGradOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferBatchNormGradOp( ++@@ -1861,8 +1861,8 @@ ++ ++ LogicalResult BatchNormTrainingOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ BatchNormTrainingOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1877,8 +1877,8 @@ ++ ++ LogicalResult BatchNormInferenceOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ BatchNormInferenceOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -1937,7 +1937,7 @@ ++ ++ LogicalResult BroadcastOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ BroadcastOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferBroadcastOp(location, adaptor.getOperand(), ++@@ -2117,7 +2117,7 @@ ++ ++ LogicalResult ClampOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ ClampOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferClampOp(location, adaptor.getMin(), adaptor.getOperand(), ++@@ -2138,7 +2138,7 @@ ++ ++ LogicalResult ComplexOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ ComplexOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferComplexOp(location, adaptor.getLhs(), inferredReturnTypes); ++@@ -2150,7 +2150,7 @@ ++ ++ LogicalResult ImagOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ ImagOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferImagOp(location, adaptor.getOperand(), inferredReturnTypes); ++@@ -2162,7 +2162,7 @@ ++ ++ LogicalResult IsFiniteOp::inferReturnTypes( ++ MLIRContext* ctx, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ IsFiniteOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferIsFiniteOp(ctx, location, adaptor.getX(), ++@@ -2175,7 +2175,7 @@ ++ ++ LogicalResult RealOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ RealOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferRealOp(location, adaptor.getOperand(), inferredReturnTypes); ++@@ -2187,7 +2187,7 @@ ++ ++ LogicalResult ConcatenateOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ ConcatenateOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferConcatenateOp(location, adaptor.getInputs().getTypes(), ++@@ -2283,7 +2283,7 @@ ++ ++ LogicalResult DynamicSliceOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ DynamicSliceOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferDynamicSliceOp(location, adaptor.getOperand().getType(), ++@@ -2367,7 +2367,7 @@ ++ ++ LogicalResult MapOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ MapOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferMapOp(location, adaptor.getInputs(), adaptor.getDimensions(), ++@@ -2399,7 +2399,7 @@ ++ ++ LogicalResult OutfeedOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ return hlo::inferOutfeedOp(getStablehloDialect(context), location, ++ inferredReturnTypes); ++@@ -2411,7 +2411,7 @@ ++ ++ LogicalResult SendOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ SendOp::Adaptor adaptor(operands, attributes, properties, regions); ++ bool isDeviceToDevice = adaptor.getChannelHandle().getType() == 1; ++@@ -2439,7 +2439,7 @@ ++ ++ LogicalResult ReduceWindowOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ ReduceWindowOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferReduceWindowOp( ++@@ -2535,7 +2535,7 @@ ++ ++ LogicalResult ReduceOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ ReduceOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferReduceOp(location, adaptor.getInputs().getTypes(), ++@@ -2555,8 +2555,8 @@ ++ Attribute encoding; ++ ReduceOp::Adaptor adaptor( ++ odsState.operands, ++- odsState.attributes.getDictionary(odsState.getContext()), {}, ++- odsState.regions); +++ odsState.attributes.getDictionary(odsState.getContext()), +++ ::mlir::PropertyRef{}, odsState.regions); ++ ++ auto inputArgTensorTypes = ++ llvm::map_to_vector(adaptor.getInputs().getTypes(), ++@@ -2626,7 +2626,7 @@ ++ //===----------------------------------------------------------------------===// ++ LogicalResult OptimizationBarrierOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ OptimizationBarrierOp::Adaptor adaptor(operands, attributes, properties); ++ return hlo::inferOptimizationBarrierOp(location, adaptor.getOperand(), ++@@ -2642,7 +2642,7 @@ ++ ++ LogicalResult ReverseOp::inferReturnTypes( ++ MLIRContext* /*context*/, std::optional location, ++- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, +++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ++ ReverseOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferReverseOp(location, adaptor.getOperand().getType(), ++@@ -2672,8 +2672,8 @@ ++ ++ LogicalResult RngOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ RngOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferRngOp( ++@@ -2697,7 +2697,7 @@ ++ ++ LogicalResult SelectOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange, ++ SmallVectorImpl& inferredReturnShapes) { ++ SelectOp::Adaptor op(operands, attributes); ++ return hlo::inferSelectOp(location, op.getPred(), op.getOnTrue(), ++@@ -2718,8 +2718,8 @@ ++ ++ LogicalResult SetDimensionSizeOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -2762,7 +2762,7 @@ ++ ++ LogicalResult PadOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ PadOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferPadOp(location, adaptor.getOperand().getType(), ++@@ -2917,7 +2917,7 @@ ++ ++ LogicalResult ReplicaIdOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr, OpaqueProperties, RegionRange, +++ DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ return hlo::inferReplicaIdOp(context, location, inferredReturnTypes); ++ } ++@@ -2928,7 +2928,7 @@ ++ ++ LogicalResult PartitionIdOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ++- ValueRange /*operands*/, DictionaryAttr, OpaqueProperties, RegionRange, +++ ValueRange /*operands*/, DictionaryAttr, PropertyRef, RegionRange, ++ SmallVectorImpl& inferredReturnTypes) { ++ return hlo::inferPartitionIdOp(context, location, inferredReturnTypes); ++ } ++@@ -2939,7 +2939,7 @@ ++ ++ LogicalResult IfOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ IfOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferIfOp(location, adaptor.getPred(), adaptor.getRegions(), ++@@ -2952,7 +2952,7 @@ ++ ++ LogicalResult CaseOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ CaseOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferCaseOp(location, adaptor.getIndex(), adaptor.getRegions(), ++@@ -2997,7 +2997,7 @@ ++ ++ LogicalResult SliceOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ SliceOpAdaptor adaptor(operands, attributes, properties); ++ return hlo::inferSliceOp(location, adaptor.getOperand().getType(), ++@@ -3022,7 +3022,7 @@ ++ ++ LogicalResult SortOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ SortOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferSortOp(location, adaptor.getInputs(), inferredReturnShapes); ++@@ -3074,7 +3074,7 @@ ++ ++ LogicalResult TransposeOp::inferReturnTypes( ++ MLIRContext*, std::optional loc, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ TransposeOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferTransposeOp(loc, adaptor.getOperand(), ++@@ -3100,7 +3100,7 @@ ++ ++ LogicalResult TriangularSolveOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ TriangularSolveOp::Adaptor adaptor(operands, attributes, properties, regions); ++ bool isTransposeAInvalid = ++@@ -3133,7 +3133,7 @@ ++ ++ LogicalResult GetTupleElementOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ GetTupleElementOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferGetTupleElementOp(location, adaptor.getOperand(), ++@@ -3146,7 +3146,7 @@ ++ ++ LogicalResult TupleOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ TupleOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferTupleOp(context, location, adaptor.getVal(), ++@@ -3171,8 +3171,8 @@ ++ ++ LogicalResult CompareOp::inferReturnTypeComponents( ++ MLIRContext* context, std::optional location, ++- ValueShapeRange operands, DictionaryAttr attributes, ++- OpaqueProperties properties, RegionRange regions, +++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, +++ RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ CompareOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferCompareOp(context, location, adaptor.getLhs(), ++@@ -3192,7 +3192,7 @@ ++ ++ LogicalResult SelectAndScatterOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ SelectAndScatterOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++@@ -3214,7 +3214,7 @@ ++ ++ LogicalResult ScatterOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ ScatterOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferScatterOp(location, adaptor.getInputs(), ++@@ -3255,7 +3255,7 @@ ++ ++ LogicalResult WhileOp::inferReturnTypes( ++ MLIRContext*, std::optional location, ValueRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnTypes) { ++ WhileOp::Adaptor adaptor(operands, attributes, properties, regions); ++ return hlo::inferWhileOp(location, adaptor.getOperand(), inferredReturnTypes); ++@@ -3279,7 +3279,7 @@ ++ ++ LogicalResult UniformDequantizeOp::inferReturnTypeComponents( ++ MLIRContext*, std::optional location, ValueShapeRange operands, ++- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, ++ SmallVectorImpl& inferredReturnShapes) { ++ UniformDequantizeOp::Adaptor adaptor(operands, attributes, properties, ++ regions); ++diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.td b/stablehlo/stablehlo/dialect/StablehloOps.td ++--- stablehlo/stablehlo/dialect/StablehloOps.td +++++ stablehlo/stablehlo/dialect/StablehloOps.td ++@@ -850,7 +850,7 @@ ++ static LogicalResult inferReturnTypes( ++ MLIRContext * /*context*/, std::optional location, ++ ValueRange operands, DictionaryAttr /*attributes*/, ++- OpaqueProperties /*properties*/, RegionRange /*regions*/, +++ PropertyRef /*properties*/, RegionRange /*regions*/, ++ SmallVectorImpl &inferredReturnTypes) { ++ if (operands.empty()) ++ return emitOptionalError( + diff --ruN a/stablehlo/stablehlo/integrations/c/StablehloTypes.cpp b/stablehlo/stablehlo/integrations/c/StablehloTypes.cpp + --- stablehlo/stablehlo/integrations/c/StablehloTypes.cpp + +++ stablehlo/stablehlo/integrations/c/StablehloTypes.cpp diff --git a/third_party/shardy/workspace.bzl b/third_party/shardy/workspace.bzl index d19e78c34aa1e..17612bbda4453 100644 --- a/third_party/shardy/workspace.bzl +++ b/third_party/shardy/workspace.bzl @@ -3,8 +3,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls") def repo(): - SHARDY_COMMIT = "7991023acf075463d524ee94491f634514c5aa32" - SHARDY_SHA256 = "c007b67fd21e484848d9ec3b334abea31a895b71d998ba11a3a0b6e60826dd2a" + SHARDY_COMMIT = "71aa6159555e8da42911a0d60cf9225d72d42980" + SHARDY_SHA256 = "055f595eff9eb09cf39bcac3696d1b14f287c1dd6262c221ecf4f1c7f1b5128c" tf_http_archive( name = "shardy", diff --git a/third_party/stablehlo/temporary.patch b/third_party/stablehlo/temporary.patch index f64aa01c51d9a..8341fcc6fab9c 100755 --- a/third_party/stablehlo/temporary.patch +++ b/third_party/stablehlo/temporary.patch @@ -1,3 +1,731 @@ +diff --ruN a/stablehlo/stablehlo/dialect/Base.h b/stablehlo/stablehlo/dialect/Base.h +--- stablehlo/stablehlo/dialect/Base.h ++++ stablehlo/stablehlo/dialect/Base.h +@@ -456,7 +456,7 @@ + static LogicalResult inferReturnTypes( + MLIRContext* /*context*/, std::optional location, + ValueRange operands, DictionaryAttr /*attributes*/, +- OpaqueProperties /*properties*/, RegionRange /*regions*/, ++ PropertyRef /*properties*/, RegionRange /*regions*/, + SmallVectorImpl& inferredReturnTypes) { + // TODO(b/231358795): Review the use of InferTypeOpInterface for ops that + // support quantization or sparsity. +@@ -478,7 +478,7 @@ + static LogicalResult inferReturnTypeComponentsFromOperands( + MLIRContext* context, std::optional location, + ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SmallVector inferredReturnTypes; + if (failed(inferReturnTypes(context, location, operands.getValues(), +diff --ruN a/stablehlo/stablehlo/dialect/ChloOps.cpp b/stablehlo/stablehlo/dialect/ChloOps.cpp +--- stablehlo/stablehlo/dialect/ChloOps.cpp ++++ stablehlo/stablehlo/dialect/ChloOps.cpp +@@ -71,7 +71,7 @@ + LogicalResult Op::inferReturnTypeComponents( \ + MLIRContext* context, std::optional location, \ + ValueShapeRange operands, DictionaryAttr attributes, \ +- OpaqueProperties properties, RegionRange regions, \ ++ PropertyRef properties, RegionRange regions, \ + SmallVectorImpl& inferredReturnShapes) { \ + return inferReturnTypeComponentsFromOperands( \ + context, location, operands, attributes, properties, regions, \ +@@ -199,7 +199,7 @@ + + LogicalResult InferBroadcastBinaryOpReturnTypeComponents( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, ++ DictionaryAttr attributes, PropertyRef properties, + std::optional> broadcastDimensions, Type elementType, + SmallVectorImpl& inferredReturnShapes) { + // Handle unranked. +@@ -252,8 +252,8 @@ + + LogicalResult BroadcastComplexOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + ShapedType lhsType = cast(operands[0].getType()); + Type elementType = ComplexType::get(lhsType.getElementType()); +@@ -287,8 +287,8 @@ + + LogicalResult BroadcastCompareOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + Type elementType = IntegerType::get(context, 1); + Adaptor adaptor(operands, attributes, properties, regions); +@@ -317,7 +317,7 @@ + + LogicalResult IsInfOp::inferReturnTypes( + MLIRContext* /*ctx*/, std::optional, ValueRange operands, +- DictionaryAttr, OpaqueProperties, RegionRange, ++ DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); + return success(); +@@ -329,7 +329,7 @@ + + LogicalResult IsNegInfOp::inferReturnTypes( + MLIRContext* /*ctx*/, std::optional, ValueRange operands, +- DictionaryAttr, OpaqueProperties, RegionRange, ++ DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); + return success(); +@@ -341,7 +341,7 @@ + + LogicalResult IsPosInfOp::inferReturnTypes( + MLIRContext* /*ctx*/, std::optional, ValueRange operands, +- DictionaryAttr, OpaqueProperties, RegionRange, ++ DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + inferredReturnTypes.push_back(getIsInfLikeReturnType(operands.front())); + return success(); +@@ -355,7 +355,7 @@ + LogicalResult Op::inferReturnTypeComponents( \ + MLIRContext* context, std::optional location, \ + ValueShapeRange operands, DictionaryAttr attributes, \ +- OpaqueProperties properties, RegionRange regions, \ ++ PropertyRef properties, RegionRange regions, \ + SmallVectorImpl& inferredReturnShapes) { \ + Adaptor adaptor(operands.getValues(), attributes, properties, regions); \ + return InferBroadcastBinaryOpReturnTypeComponents( \ +@@ -400,8 +400,8 @@ + + LogicalResult ConstantLikeOp::inferReturnTypeComponents( + MLIRContext* /*context*/, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange /*regions*/, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange /*regions*/, + SmallVectorImpl& inferredReturnShapes) { + ConstantLikeOp::Adaptor op(operands, attributes, properties); + if (failed(op.verify(location.value()))) return failure(); +@@ -437,7 +437,7 @@ + + LogicalResult BroadcastSelectOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr, OpaqueProperties, RegionRange, ++ DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnShapes) { + BroadcastSelectOp::Adaptor op(operands.getValues()); + auto onTrueType = cast(op.getOnTrue().getType()); +@@ -757,8 +757,8 @@ + + LogicalResult TopKOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + TopKOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferTopKOp(location, adaptor.getOperand(), adaptor.getK(), +@@ -773,7 +773,7 @@ + + LogicalResult ConstantOp::inferReturnTypes( + MLIRContext*, std::optional, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + Adaptor adaptor(operands, attributes, properties, regions); + Type type = cast(adaptor.getValueAttr()).getType(); +@@ -787,7 +787,7 @@ + + LogicalResult ScanOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + ScanOp::Adaptor adaptor(operands, attributes, properties, regions); + if (regions.empty() || regions.front()->empty()) { +diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp +--- stablehlo/stablehlo/dialect/StablehloOps.cpp ++++ stablehlo/stablehlo/dialect/StablehloOps.cpp +@@ -216,7 +216,7 @@ + LogicalResult Op::inferReturnTypeComponents( \ + MLIRContext* context, std::optional location, \ + ValueShapeRange operands, DictionaryAttr attributes, \ +- OpaqueProperties properties, RegionRange regions, \ ++ PropertyRef properties, RegionRange regions, \ + SmallVectorImpl& inferredReturnShapes) { \ + return inferReturnTypeComponentsFromOperands( \ + context, location, operands, attributes, properties, regions, \ +@@ -267,8 +267,8 @@ + + LogicalResult AddOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SmallVector inferredReturnTypes; + if (failed(inferReturnTypes(context, location, operands.getValues(), +@@ -293,7 +293,7 @@ + + LogicalResult AfterAllOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + return hlo::inferAfterAllOp(getStablehloDialect(context), location, + inferredReturnTypes); +@@ -357,7 +357,7 @@ + + LogicalResult ConstantOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + ConstantOpAdaptor adaptor(operands, attributes, properties); + return hlo::inferConstantOp(location, adaptor.getValue(), +@@ -391,7 +391,7 @@ + + LogicalResult CreateTokenOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + return hlo::inferCreateTokenOp(getStablehloDialect(context), location, + inferredReturnTypes); +@@ -729,7 +729,7 @@ + + LogicalResult CholeskyOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + CholeskyOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferCholeskyOp(location, adaptor.getA(), inferredReturnShapes); +@@ -746,7 +746,7 @@ + + LogicalResult DotOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + DotOp::Adaptor adaptor(operands, attributes, properties, regions); + auto lhsType = mlir::cast(adaptor.getLhs().getType()); +@@ -885,7 +885,7 @@ + + LogicalResult DotGeneralOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + DotGeneralOp::Adaptor adaptor(operands, attributes, properties, regions); + LLVM_DEBUG(llvm::dbgs() << "DotGeneralOp::inferReturnTypeComponents\n"); +@@ -1183,7 +1183,7 @@ + + LogicalResult FftOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + FftOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferFftOp(location, adaptor.getOperand(), +@@ -1304,8 +1304,8 @@ + + LogicalResult GatherOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + GatherOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferGatherOp( +@@ -1345,8 +1345,8 @@ + + LogicalResult DynamicGatherOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + DynamicGatherOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferDynamicGatherOp( +@@ -1365,7 +1365,7 @@ + + LogicalResult GetDimensionSizeOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + GetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1435,7 +1435,7 @@ + + LogicalResult DynamicUpdateSliceOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + DynamicUpdateSliceOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1450,7 +1450,7 @@ + + LogicalResult AbsOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + AbsOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferAbsOp(location, adaptor.getOperand(), inferredReturnTypes); +@@ -1470,7 +1470,7 @@ + + LogicalResult CollectiveBroadcastOp::inferReturnTypes( + MLIRContext* /*context*/, std::optional location, +- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, ++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + CollectiveBroadcastOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1480,8 +1480,8 @@ + + LogicalResult CollectiveBroadcastOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SmallVector inferredReturnTypes; + CollectiveBroadcastOp::Adaptor adaptor(operands, attributes, properties, +@@ -1513,7 +1513,7 @@ + + LogicalResult CollectivePermuteOp::inferReturnTypes( + MLIRContext* /*context*/, std::optional location, +- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, ++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + CollectivePermuteOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1523,8 +1523,8 @@ + + LogicalResult CollectivePermuteOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SmallVector inferredReturnTypes; + CollectivePermuteOp::Adaptor adaptor(operands, attributes, properties, +@@ -1634,7 +1634,7 @@ + + LogicalResult AllToAllOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + AllToAllOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferAllToAllOp( +@@ -1736,7 +1736,7 @@ + + LogicalResult AllReduceOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + AllReduceOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferAllReduceOp(location, adaptor.getOperands(), +@@ -1752,7 +1752,7 @@ + // async_start references stablehlo.futures. + LogicalResult AsyncStartOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + if (regions.empty()) { + return emitOptionalError( +@@ -1820,7 +1820,7 @@ + // TODO: Move to TypeInference.h and TypeInference.cpp. See TODO above. + LogicalResult AsyncDoneOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + auto futureType = dyn_cast(operands[0].getType()); + if (!futureType) { +@@ -1845,8 +1845,8 @@ + + LogicalResult BatchNormGradOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + BatchNormGradOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferBatchNormGradOp( +@@ -1861,8 +1861,8 @@ + + LogicalResult BatchNormTrainingOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + BatchNormTrainingOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1877,8 +1877,8 @@ + + LogicalResult BatchNormInferenceOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + BatchNormInferenceOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -1937,7 +1937,7 @@ + + LogicalResult BroadcastOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + BroadcastOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferBroadcastOp(location, adaptor.getOperand(), +@@ -2117,7 +2117,7 @@ + + LogicalResult ClampOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + ClampOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferClampOp(location, adaptor.getMin(), adaptor.getOperand(), +@@ -2138,7 +2138,7 @@ + + LogicalResult ComplexOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + ComplexOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferComplexOp(location, adaptor.getLhs(), inferredReturnTypes); +@@ -2150,7 +2150,7 @@ + + LogicalResult ImagOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + ImagOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferImagOp(location, adaptor.getOperand(), inferredReturnTypes); +@@ -2162,7 +2162,7 @@ + + LogicalResult IsFiniteOp::inferReturnTypes( + MLIRContext* ctx, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + IsFiniteOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferIsFiniteOp(ctx, location, adaptor.getX(), +@@ -2175,7 +2175,7 @@ + + LogicalResult RealOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + RealOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferRealOp(location, adaptor.getOperand(), inferredReturnTypes); +@@ -2187,7 +2187,7 @@ + + LogicalResult ConcatenateOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + ConcatenateOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferConcatenateOp(location, adaptor.getInputs().getTypes(), +@@ -2283,7 +2283,7 @@ + + LogicalResult DynamicSliceOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + DynamicSliceOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferDynamicSliceOp(location, adaptor.getOperand().getType(), +@@ -2367,7 +2367,7 @@ + + LogicalResult MapOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + MapOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferMapOp(location, adaptor.getInputs(), adaptor.getDimensions(), +@@ -2399,7 +2399,7 @@ + + LogicalResult OutfeedOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + return hlo::inferOutfeedOp(getStablehloDialect(context), location, + inferredReturnTypes); +@@ -2411,7 +2411,7 @@ + + LogicalResult SendOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + SendOp::Adaptor adaptor(operands, attributes, properties, regions); + bool isDeviceToDevice = adaptor.getChannelHandle().getType() == 1; +@@ -2439,7 +2439,7 @@ + + LogicalResult ReduceWindowOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + ReduceWindowOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferReduceWindowOp( +@@ -2535,7 +2535,7 @@ + + LogicalResult ReduceOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + ReduceOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferReduceOp(location, adaptor.getInputs().getTypes(), +@@ -2555,8 +2555,8 @@ + Attribute encoding; + ReduceOp::Adaptor adaptor( + odsState.operands, +- odsState.attributes.getDictionary(odsState.getContext()), {}, +- odsState.regions); ++ odsState.attributes.getDictionary(odsState.getContext()), ++ ::mlir::PropertyRef{}, odsState.regions); + + auto inputArgTensorTypes = + llvm::map_to_vector(adaptor.getInputs().getTypes(), +@@ -2626,7 +2626,7 @@ + //===----------------------------------------------------------------------===// + LogicalResult OptimizationBarrierOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + OptimizationBarrierOp::Adaptor adaptor(operands, attributes, properties); + return hlo::inferOptimizationBarrierOp(location, adaptor.getOperand(), +@@ -2642,7 +2642,7 @@ + + LogicalResult ReverseOp::inferReturnTypes( + MLIRContext* /*context*/, std::optional location, +- ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, ++ ValueRange operands, DictionaryAttr attributes, PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + ReverseOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferReverseOp(location, adaptor.getOperand().getType(), +@@ -2672,8 +2672,8 @@ + + LogicalResult RngOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + RngOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferRngOp( +@@ -2697,7 +2697,7 @@ + + LogicalResult SelectOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange, + SmallVectorImpl& inferredReturnShapes) { + SelectOp::Adaptor op(operands, attributes); + return hlo::inferSelectOp(location, op.getPred(), op.getOnTrue(), +@@ -2718,8 +2718,8 @@ + + LogicalResult SetDimensionSizeOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -2762,7 +2762,7 @@ + + LogicalResult PadOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + PadOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferPadOp(location, adaptor.getOperand().getType(), +@@ -2917,7 +2917,7 @@ + + LogicalResult ReplicaIdOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr, OpaqueProperties, RegionRange, ++ DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + return hlo::inferReplicaIdOp(context, location, inferredReturnTypes); + } +@@ -2928,7 +2928,7 @@ + + LogicalResult PartitionIdOp::inferReturnTypes( + MLIRContext* context, std::optional location, +- ValueRange /*operands*/, DictionaryAttr, OpaqueProperties, RegionRange, ++ ValueRange /*operands*/, DictionaryAttr, PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { + return hlo::inferPartitionIdOp(context, location, inferredReturnTypes); + } +@@ -2939,7 +2939,7 @@ + + LogicalResult IfOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + IfOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferIfOp(location, adaptor.getPred(), adaptor.getRegions(), +@@ -2952,7 +2952,7 @@ + + LogicalResult CaseOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + CaseOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferCaseOp(location, adaptor.getIndex(), adaptor.getRegions(), +@@ -2997,7 +2997,7 @@ + + LogicalResult SliceOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + SliceOpAdaptor adaptor(operands, attributes, properties); + return hlo::inferSliceOp(location, adaptor.getOperand().getType(), +@@ -3022,7 +3022,7 @@ + + LogicalResult SortOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + SortOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferSortOp(location, adaptor.getInputs(), inferredReturnShapes); +@@ -3074,7 +3074,7 @@ + + LogicalResult TransposeOp::inferReturnTypes( + MLIRContext*, std::optional loc, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + TransposeOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferTransposeOp(loc, adaptor.getOperand(), +@@ -3100,7 +3100,7 @@ + + LogicalResult TriangularSolveOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + TriangularSolveOp::Adaptor adaptor(operands, attributes, properties, regions); + bool isTransposeAInvalid = +@@ -3133,7 +3133,7 @@ + + LogicalResult GetTupleElementOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + GetTupleElementOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferGetTupleElementOp(location, adaptor.getOperand(), +@@ -3146,7 +3146,7 @@ + + LogicalResult TupleOp::inferReturnTypes( + MLIRContext* context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + TupleOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferTupleOp(context, location, adaptor.getVal(), +@@ -3171,8 +3171,8 @@ + + LogicalResult CompareOp::inferReturnTypeComponents( + MLIRContext* context, std::optional location, +- ValueShapeRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, ++ ValueShapeRange operands, DictionaryAttr attributes, PropertyRef properties, ++ RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + CompareOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferCompareOp(context, location, adaptor.getLhs(), +@@ -3192,7 +3192,7 @@ + + LogicalResult SelectAndScatterOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + SelectAndScatterOp::Adaptor adaptor(operands, attributes, properties, + regions); +@@ -3214,7 +3214,7 @@ + + LogicalResult ScatterOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + ScatterOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferScatterOp(location, adaptor.getInputs(), +@@ -3255,7 +3255,7 @@ + + LogicalResult WhileOp::inferReturnTypes( + MLIRContext*, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { + WhileOp::Adaptor adaptor(operands, attributes, properties, regions); + return hlo::inferWhileOp(location, adaptor.getOperand(), inferredReturnTypes); +@@ -3279,7 +3279,7 @@ + + LogicalResult UniformDequantizeOp::inferReturnTypeComponents( + MLIRContext*, std::optional location, ValueShapeRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, ++ DictionaryAttr attributes, PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnShapes) { + UniformDequantizeOp::Adaptor adaptor(operands, attributes, properties, + regions); +diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.td b/stablehlo/stablehlo/dialect/StablehloOps.td +--- stablehlo/stablehlo/dialect/StablehloOps.td ++++ stablehlo/stablehlo/dialect/StablehloOps.td +@@ -850,7 +850,7 @@ + static LogicalResult inferReturnTypes( + MLIRContext * /*context*/, std::optional location, + ValueRange operands, DictionaryAttr /*attributes*/, +- OpaqueProperties /*properties*/, RegionRange /*regions*/, ++ PropertyRef /*properties*/, RegionRange /*regions*/, + SmallVectorImpl &inferredReturnTypes) { + if (operands.empty()) + return emitOptionalError( diff --ruN a/stablehlo/stablehlo/integrations/c/StablehloTypes.cpp b/stablehlo/stablehlo/integrations/c/StablehloTypes.cpp --- stablehlo/stablehlo/integrations/c/StablehloTypes.cpp +++ stablehlo/stablehlo/integrations/c/StablehloTypes.cpp diff --git a/third_party/triton/common/llvm_cl893905698.patch b/third_party/triton/common/llvm_cl893905698.patch new file mode 100644 index 0000000000000..00bbe700605f8 --- /dev/null +++ b/third_party/triton/common/llvm_cl893905698.patch @@ -0,0 +1,146 @@ +--- a/include/triton/Dialect/TritonNvidiaGPU/IR/Dialect.h ++++ b/include/triton/Dialect/TritonNvidiaGPU/IR/Dialect.h +@@ -70,6 +70,7 @@ + + struct TensorMemory : public SideEffects::Resource::Base { + StringRef getName() const final { return ""; } ++ Resource* getParent() const override { return nullptr; } + }; + + struct TMemAllocation { + +--- a/lib/Dialect/Triton/IR/Ops.cpp ++++ b/lib/Dialect/Triton/IR/Ops.cpp +@@ -214,11 +214,10 @@ + } + + //-- DotOp -- +-LogicalResult +-DotOp::inferReturnTypes(MLIRContext *context, std::optional location, +- ValueRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++LogicalResult DotOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + // type is the same as the accumulator + auto accTy = cast(operands[2].getType()); + inferredReturnTypes.push_back(accTy); +@@ -481,11 +480,10 @@ + return success(); + } + +-LogicalResult +-ReduceOp::inferReturnTypes(MLIRContext *context, std::optional loc, +- ValueRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++LogicalResult ReduceOp::inferReturnTypes( ++ MLIRContext* context, std::optional loc, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + Properties *prop = properties.as(); + int axis = prop->axis.getInt(); + for (auto arg : operands) { +@@ -636,11 +634,10 @@ + ScanOp::build(builder, state, inferredReturnTypes, operands, axis, reverse); + } + +-LogicalResult +-ScanOp::inferReturnTypes(MLIRContext *context, std::optional location, +- ValueRange operands, DictionaryAttr attributes, +- OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++LogicalResult ScanOp::inferReturnTypes( ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + for (auto arg : operands) + inferredReturnTypes.push_back(arg.getType()); + return success(); +@@ -737,9 +734,9 @@ + } + + LogicalResult UnsplatOp::inferReturnTypes( +- MLIRContext *context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + auto dstTy = cast(operands[0].getType()).getElementType(); + inferredReturnTypes.push_back(dstTy); + return success(); +@@ -747,9 +744,9 @@ + + //-- ExpandDimsOp -- + LogicalResult ExpandDimsOp::inferReturnTypes( +- MLIRContext *context, std::optional loc, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional loc, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + // infer shape + auto arg = operands[0]; + auto argTy = cast(arg.getType()); +@@ -1344,9 +1341,9 @@ + } + + LogicalResult GatherOp::inferReturnTypes( +- MLIRContext *context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + GatherOpAdaptor adaptor(operands, attributes, properties, regions); + auto indicesType = cast(adaptor.getIndices().getType()); + auto srcType = cast(adaptor.getSrc().getType()); + +--- a/lib/Dialect/TritonNvidiaGPU/IR/Ops.cpp ++++ b/lib/Dialect/TritonNvidiaGPU/IR/Ops.cpp +@@ -49,9 +49,9 @@ + + // -- WarpGroupDotOp -- + LogicalResult WarpGroupDotOp::inferReturnTypes( +- MLIRContext *context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + // type is the same as the accumulator + auto accTy = cast(operands[2].getType()); + inferredReturnTypes.push_back(accTy); +@@ -156,9 +156,9 @@ + + // -- WarpGroupDotWaitOp -- + LogicalResult WarpGroupDotWaitOp::inferReturnTypes( +- MLIRContext *context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + for (Value operand : operands) + inferredReturnTypes.push_back(operand.getType()); + return success(); + +--- a/third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp ++++ b/third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp +@@ -851,9 +851,9 @@ + // prefetch. Values are the byte offsets added to the base pointer for each + // prefetch instruction. + LogicalResult TDMPrefetchOp::inferReturnTypes( +- MLIRContext *context, std::optional location, ValueRange operands, +- DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, +- SmallVectorImpl &inferredReturnTypes) { ++ MLIRContext* context, std::optional location, ValueRange operands, ++ DictionaryAttr attributes, mlir::PropertyRef properties, ++ RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + TDMPrefetchOp::Adaptor ad(operands, attributes, properties, regions); + + // If returnOffsets is not set the op will not return any results + diff --git a/third_party/triton/common/series.bzl b/third_party/triton/common/series.bzl index 7f05a219a8337..7d12d914c7902 100644 --- a/third_party/triton/common/series.bzl +++ b/third_party/triton/common/series.bzl @@ -29,5 +29,6 @@ common_patch_list = [ "//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_cl893905698.patch", # Add new patches just above this line ] diff --git a/xla/backends/gpu/codegen/emitters/ir/xla_gpu_ops.cc b/xla/backends/gpu/codegen/emitters/ir/xla_gpu_ops.cc index 1635c533b6196..352a4746d01c6 100644 --- a/xla/backends/gpu/codegen/emitters/ir/xla_gpu_ops.cc +++ b/xla/backends/gpu/codegen/emitters/ir/xla_gpu_ops.cc @@ -139,7 +139,7 @@ SmallVector inferReductionInitTypes(TypeRange input_types) { LogicalResult ReduceOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - mlir::DictionaryAttr attributes, mlir::OpaqueProperties properties, + mlir::DictionaryAttr attributes, mlir::PropertyRef properties, mlir::RegionRange regions, mlir::SmallVectorImpl& inferredReturnTypes) { ReduceOp::Adaptor adaptor(operands, attributes, properties, regions); diff --git a/xla/mlir_hlo/mhlo/IR/hlo_ops.cc b/xla/mlir_hlo/mhlo/IR/hlo_ops.cc index 423bb3ae2f875..0779fa58ce9fe 100644 --- a/xla/mlir_hlo/mhlo/IR/hlo_ops.cc +++ b/xla/mlir_hlo/mhlo/IR/hlo_ops.cc @@ -347,7 +347,7 @@ LogicalResult ReduceScatterOp::verify() { LogicalResult Op::inferReturnTypeComponents( \ MLIRContext* context, std::optional location, \ ValueShapeRange operands, DictionaryAttr attributes, \ - OpaqueProperties properties, RegionRange regions, \ + mlir::PropertyRef properties, RegionRange regions, \ SmallVectorImpl& inferredReturnShapes) { \ return inferReturnTypeComponentsFromOperands( \ context, location, operands, attributes, properties, regions, \ @@ -517,8 +517,8 @@ LogicalResult AsyncUpdateOp::verify() { LogicalResult AsyncUpdateOp::inferReturnTypes( MLIRContext*, std::optional, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { AsyncUpdateOp::Adaptor adaptor(operands, attributes, properties, regions); auto stateType = cast(adaptor.getBundle().getType()); inferredReturnTypes.push_back(stateType); @@ -546,8 +546,8 @@ LogicalResult AsyncDoneOp::verify() { LogicalResult AsyncDoneOp::inferReturnTypes( MLIRContext*, std::optional, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { AsyncDoneOp::Adaptor adaptor(operands, attributes, properties, regions); AsyncStartOp startOp = findAsyncChainStart(operands[0].getDefiningOp()); @@ -575,7 +575,7 @@ LogicalResult AsyncDoneOp::inferReturnTypes( LogicalResult AfterAllOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange, - DictionaryAttr, OpaqueProperties, RegionRange, + DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { return hlo::inferAfterAllOp(getMhloDialect(context), location, inferredReturnTypes); @@ -631,8 +631,8 @@ void ConstantOp::build(OpBuilder& /*builder*/, OperationState& result, LogicalResult ConstantOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ConstantOpAdaptor adaptor(operands, attributes, properties, regions); return hlo::inferConstantOp(location, adaptor.getValue(), inferredReturnTypes); @@ -721,7 +721,7 @@ LogicalResult FusionOp::verify() { return verifyOutputOperandAliasing(this); } LogicalResult CreateTokenOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange, - DictionaryAttr, OpaqueProperties, RegionRange, + DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { return hlo::inferCreateTokenOp(getMhloDialect(context), location, inferredReturnTypes); @@ -879,7 +879,8 @@ void CustomCallOp::getEffects( LogicalResult CholeskyOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { CholeskyOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferCholeskyOp(location, adaptor.getA(), inferredReturnShapes); @@ -1428,7 +1429,8 @@ static LogicalResult verify1dTensor(std::optional loc, LogicalResult FftOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { FftOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getFftLength(), "fft_length"))) @@ -1615,7 +1617,7 @@ LogicalResult GatherOp::reifyReturnTypeShapes( LogicalResult GatherOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { GatherOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getSliceSizes(), "slice_sizes"))) @@ -1678,7 +1680,7 @@ LogicalResult DynamicGatherOp::reifyReturnTypeShapes( LogicalResult DynamicGatherOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { DynamicGatherOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferDynamicGatherOp( @@ -1697,7 +1699,8 @@ LogicalResult DynamicGatherOp::inferReturnTypeComponents( LogicalResult GetDimensionSizeOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { GetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -1876,7 +1879,8 @@ LogicalResult DynamicIotaOp::reifyReturnTypeShapes( LogicalResult DynamicUpdateSliceOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { DynamicUpdateSliceOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -1917,8 +1921,8 @@ OpFoldResult DynamicUpdateSliceOp::fold(FoldAdaptor /*adaptor*/) { LogicalResult AbsOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { AbsOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferAbsOp(location, adaptor.getOperand(), inferredReturnTypes); } @@ -2385,7 +2389,8 @@ void TupleOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult AllToAllOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { AllToAllOp::Adaptor adaptor(operands, attributes, properties, regions); std::optional splitDimension = adaptor.getSplitDimension(); @@ -2507,7 +2512,8 @@ void AllReduceOp::build(OpBuilder& odsBuilder, OperationState& odsState, LogicalResult AllReduceOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { AllReduceOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -2540,7 +2546,7 @@ LogicalResult AllReduceOp::inferReturnTypeComponents( LogicalResult BatchNormGradOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { BatchNormGradOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferBatchNormGradOp( @@ -2556,7 +2562,7 @@ LogicalResult BatchNormGradOp::inferReturnTypeComponents( LogicalResult BatchNormTrainingOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { BatchNormTrainingOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -2572,7 +2578,7 @@ LogicalResult BatchNormTrainingOp::inferReturnTypeComponents( LogicalResult BatchNormInferenceOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { BatchNormInferenceOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -2676,7 +2682,8 @@ OpFoldResult BroadcastOp::fold(FoldAdaptor adaptor) { LogicalResult BroadcastOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { BroadcastOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getBroadcastSizes(), @@ -3016,8 +3023,8 @@ LogicalResult DynamicBroadcastInDimOp::reifyReturnTypeShapes( LogicalResult ComplexOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ComplexOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferComplexOp(location, adaptor.getLhs(), inferredReturnTypes); } @@ -3038,8 +3045,8 @@ OpFoldResult ComplexOp::fold(FoldAdaptor) { LogicalResult ImagOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ImagOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferImagOp(location, adaptor.getOperand(), inferredReturnTypes); } @@ -3058,8 +3065,8 @@ OpFoldResult ImagOp::fold(FoldAdaptor) { LogicalResult IsFiniteOp::inferReturnTypes( MLIRContext* ctx, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { IsFiniteOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferIsFiniteOp(ctx, location, adaptor.getX(), inferredReturnTypes); @@ -3071,8 +3078,8 @@ LogicalResult IsFiniteOp::inferReturnTypes( LogicalResult RealOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { RealOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferRealOp(location, adaptor.getOperand(), inferredReturnTypes); } @@ -3169,8 +3176,8 @@ class ConcatenateForwarding : public OpRewritePattern { LogicalResult ConcatenateOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ConcatenateOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferConcatenateOp(location, adaptor.getVal().getTypes(), adaptor.getDimension(), inferredReturnTypes); @@ -3508,7 +3515,8 @@ void DynamicSliceOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult DynamicSliceOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { DynamicSliceOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getSliceSizes(), "slice_sizes"))) @@ -3660,7 +3668,8 @@ LogicalResult InfeedOp::verify() { LogicalResult MapOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { MapOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getDimensions(), "dimensions"))) @@ -3699,7 +3708,7 @@ LogicalResult MapOp::reifyReturnTypeShapes( LogicalResult OutfeedOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange, - DictionaryAttr, OpaqueProperties, RegionRange, + DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { return hlo::inferOutfeedOp(getMhloDialect(context), location, inferredReturnTypes); @@ -3711,8 +3720,8 @@ LogicalResult OutfeedOp::inferReturnTypes( LogicalResult SendOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { SendOp::Adaptor adaptor(operands, attributes, properties, regions); bool isDeviceToDevice = adaptor.getChannelHandle().getType() == 1; bool isDeviceToHost = adaptor.getChannelHandle().getType() == 2; @@ -3739,7 +3748,8 @@ LogicalResult RecvOp::verify() { LogicalResult ReduceWindowOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { ReduceWindowOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferReduceWindowOp( @@ -4107,7 +4117,8 @@ ParseResult ReduceOp::parse(OpAsmParser& parser, OperationState& result) { LogicalResult ReduceOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { ReduceOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferReduceOp( @@ -4129,8 +4140,8 @@ void ReduceOp::build(OpBuilder&, OperationState& odsState, ValueRange inputs, Attribute encoding; ReduceOp::Adaptor adaptor( odsState.operands, - odsState.attributes.getDictionary(odsState.getContext()), {}, - odsState.regions); + odsState.attributes.getDictionary(odsState.getContext()), + mlir::PropertyRef{}, odsState.regions); SmallVector inputArgTensorTypes{ llvm::map_range(adaptor.getInputs().getTypes(), @@ -4309,8 +4320,8 @@ LogicalResult ReduceOp::reifyReturnTypeShapes( //===----------------------------------------------------------------------===// LogicalResult OptimizationBarrierOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { OptimizationBarrierOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferOptimizationBarrierOp(location, adaptor.getOperand(), @@ -4345,7 +4356,7 @@ LogicalResult RngBitGeneratorOp::verify() { LogicalResult RngOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { RngOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferRngOp( @@ -4384,7 +4395,8 @@ LogicalResult XlaRngGetAndUpdateStateOp::verify() { LogicalResult XlaRngGetAndUpdateStateOp::inferReturnTypes( MLIRContext* ctx, std::optional, ValueRange, DictionaryAttr, - OpaqueProperties, RegionRange, SmallVectorImpl& inferredReturnTypes) { + mlir::PropertyRef, RegionRange, + SmallVectorImpl& inferredReturnTypes) { inferredReturnTypes.push_back(mlir::RankedTensorType::get( {2}, mlir::IntegerType::get(ctx, 64, IntegerType::Unsigned))); return success(); @@ -4396,7 +4408,8 @@ LogicalResult XlaRngGetAndUpdateStateOp::inferReturnTypes( LogicalResult ScanOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { ScanOp::Adaptor adaptor(operands, attributes, properties, regions); if (regions.empty() || regions.front()->empty()) { @@ -4611,7 +4624,8 @@ void SelectOp::getCanonicalizationPatterns(RewritePatternSet& results, // the return type based on operand type. LogicalResult SelectOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { SelectOp::Adaptor op(operands, attributes, properties, regions); return hlo::inferSelectOp(location, op.getPred(), op.getOnTrue(), @@ -4657,7 +4671,7 @@ OpFoldResult SetDimensionSizeOp::fold(FoldAdaptor adaptor) { LogicalResult SetDimensionSizeOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { SetDimensionSizeOp::Adaptor adaptor(operands, attributes, properties, regions); @@ -4672,8 +4686,8 @@ LogicalResult SetDimensionSizeOp::inferReturnTypeComponents( LogicalResult PadOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { PadOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getEdgePaddingLow(), "edge_padding_low")) || @@ -5045,7 +5059,7 @@ void ReshapeOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult ReplicaIdOp::inferReturnTypes( MLIRContext* context, std::optional location, - ValueRange /*operands*/, DictionaryAttr, OpaqueProperties, RegionRange, + ValueRange /*operands*/, DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { return hlo::inferReplicaIdOp(context, location, inferredReturnTypes); } @@ -5056,7 +5070,7 @@ LogicalResult ReplicaIdOp::inferReturnTypes( LogicalResult PartitionIdOp::inferReturnTypes( MLIRContext* context, std::optional location, - ValueRange /*operands*/, DictionaryAttr, OpaqueProperties, RegionRange, + ValueRange /*operands*/, DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { return hlo::inferPartitionIdOp(context, location, inferredReturnTypes); } @@ -5067,7 +5081,7 @@ LogicalResult PartitionIdOp::inferReturnTypes( LogicalResult AddDependencyOp::inferReturnTypes( MLIRContext* context, std::optional, ValueRange operands, - DictionaryAttr, OpaqueProperties, RegionRange, + DictionaryAttr, mlir::PropertyRef, RegionRange, SmallVectorImpl& inferredReturnTypes) { inferredReturnTypes.push_back(operands.getTypes()[0]); return success(); @@ -5079,8 +5093,8 @@ LogicalResult AddDependencyOp::inferReturnTypes( LogicalResult IfOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { IfOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferIfOp(location, adaptor.getPred(), adaptor.getRegions(), inferredReturnTypes); @@ -5110,8 +5124,8 @@ void IfOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult CaseOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { CaseOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferCaseOp(location, adaptor.getIndex(), adaptor.getRegions(), inferredReturnTypes); @@ -5666,7 +5680,8 @@ OpFoldResult ClampOp::fold(FoldAdaptor adaptor) { LogicalResult ClampOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { ClampOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferClampOp(location, adaptor.getMin(), adaptor.getOperand(), @@ -5687,8 +5702,9 @@ LogicalResult ClampOp::reifyReturnTypeShapes( LogicalResult SliceOp::inferReturnTypes( MLIRContext* /*context*/, std::optional location, - ValueRange operands, DictionaryAttr attributes, OpaqueProperties properties, - RegionRange regions, SmallVectorImpl& inferredReturnTypes) { + ValueRange operands, DictionaryAttr attributes, + mlir::PropertyRef properties, RegionRange regions, + SmallVectorImpl& inferredReturnTypes) { SliceOpAdaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(location, adaptor.getStartIndices(), "start_indices")) || @@ -5919,7 +5935,8 @@ void SortOp::build(OpBuilder& builder, OperationState& state, LogicalResult SortOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { SortOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferSortOp(location, adaptor.getInputs(), inferredReturnShapes); @@ -6009,7 +6026,8 @@ void SortOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult TopKOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { TopKOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferTopKOp(location, adaptor.getOperand(), adaptor.getK(), @@ -6168,8 +6186,8 @@ LogicalResult TransposeOp::reifyReturnTypeShapes( LogicalResult TransposeOp::inferReturnTypes( MLIRContext*, std::optional loc, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { TransposeOp::Adaptor adaptor(operands, attributes, properties, regions); if (failed(verify1dTensor(loc, adaptor.getPermutation(), "permutation"))) return failure(); @@ -6185,7 +6203,8 @@ LogicalResult TransposeOp::inferReturnTypes( LogicalResult TriangularSolveOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { TriangularSolveOp::Adaptor adaptor(operands, attributes, properties, regions); bool isTransposeAInvalid = @@ -6209,8 +6228,8 @@ OpFoldResult GetTupleElementOp::fold(FoldAdaptor /*adaptor*/) { LogicalResult GetTupleElementOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { GetTupleElementOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferGetTupleElementOp(location, adaptor.getOperand(), adaptor.getIndex(), inferredReturnTypes); @@ -6222,8 +6241,8 @@ LogicalResult GetTupleElementOp::inferReturnTypes( LogicalResult TupleOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { TupleOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferTupleOp(context, location, adaptor.getVal(), inferredReturnTypes); @@ -6244,7 +6263,7 @@ void CompareOp::build(OpBuilder& builder, OperationState& result, Value lhs, LogicalResult CompareOp::inferReturnTypeComponents( MLIRContext* context, std::optional location, ValueShapeRange operands, DictionaryAttr attributes, - OpaqueProperties properties, RegionRange regions, + mlir::PropertyRef properties, RegionRange regions, SmallVectorImpl& inferredReturnShapes) { CompareOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferCompareOp(context, location, adaptor.getLhs(), @@ -6374,8 +6393,8 @@ OpFoldResult CompareOp::fold(FoldAdaptor adaptor) { LogicalResult SelectAndScatterOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { SelectAndScatterOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferSelectAndScatterOp(location, adaptor.getOperand(), @@ -6409,8 +6428,8 @@ LogicalResult SelectAndScatterOp::verify() { LogicalResult ScatterOp::inferReturnTypes( MLIRContext*, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { ScatterOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferScatterOp(location, adaptor.getInputs(), adaptor.getUpdateComputation(), @@ -6646,8 +6665,8 @@ void ScatterOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult WhileOp::inferReturnTypes( MLIRContext* context, std::optional location, ValueRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, - SmallVectorImpl& inferredReturnTypes) { + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnTypes) { WhileOp::Adaptor adaptor(operands, attributes, properties, regions); return hlo::inferWhileOp(location, adaptor.getOperand(), inferredReturnTypes); } @@ -6750,7 +6769,8 @@ void WhileOp::getCanonicalizationPatterns(RewritePatternSet& results, LogicalResult UniformDequantizeOp::inferReturnTypeComponents( MLIRContext*, std::optional location, ValueShapeRange operands, - DictionaryAttr attributes, OpaqueProperties properties, RegionRange regions, + DictionaryAttr attributes, mlir::PropertyRef properties, + RegionRange regions, SmallVectorImpl& inferredReturnShapes) { UniformDequantizeOp::Adaptor adaptor(operands, attributes, properties, regions);