Integrate LLVM at llvm/llvm-project@7189c4bb83d9 (#2945)
diff --git a/build_tools/llvm_version.txt b/build_tools/llvm_version.txt
index 2f585f1..a66a072 100644
--- a/build_tools/llvm_version.txt
+++ b/build_tools/llvm_version.txt
@@ -1 +1 @@
-2cf353b5e8560723409f3f9164bddec76f499963
+7189c4bb83d90c383741a9480fceca9c7bf74b27
diff --git a/docs/generated/chlo.md b/docs/generated/chlo.md
index 631f037..ade2255 100755
--- a/docs/generated/chlo.md
+++ b/docs/generated/chlo.md
@@ -1778,6 +1778,7 @@
<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>dimension</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute whose value is non-negative</td></tr>
+<tr><td><code>scan_dim_size</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute whose value is non-negative</td></tr>
<tr><td><code>is_reverse</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>is_associative</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>
diff --git a/stablehlo/dialect/ChloOps.cpp b/stablehlo/dialect/ChloOps.cpp
index fd2df07..637f424 100644
--- a/stablehlo/dialect/ChloOps.cpp
+++ b/stablehlo/dialect/ChloOps.cpp
@@ -805,6 +805,9 @@
// output dimension size will be dynamic as well.
int64_t dim = adaptor.getDimension();
int64_t dimSize = ShapedType::kDynamic;
+ if (adaptor.getScanDimSize().has_value()) {
+ dimSize = adaptor.getScanDimSize().value();
+ }
for (auto [i, type] : llvm::enumerate(adaptor.getInputs().getTypes())) {
auto inputType = dyn_cast<RankedTensorType>(type);
if (!inputType) {
@@ -814,9 +817,8 @@
return emitOptionalError(location, "scan dimension of operand ", i,
" is out of bounds");
}
- dimSize = inputType.getDimSize(dim);
if (dimSize == ShapedType::kDynamic) {
- break;
+ dimSize = inputType.getDimSize(dim);
}
}
@@ -837,34 +839,6 @@
return success();
}
-LogicalResult ScanOp::reifyReturnTypeShapes(
- OpBuilder& builder, ValueRange operands,
- SmallVectorImpl<Value>& reifiedReturnShapes) {
- ScanOp::Adaptor adaptor(operands, getOperation()->getAttrDictionary(),
- getOperation()->getPropertiesStorage());
- auto inputs = adaptor.getInputs();
- size_t k = adaptor.getInits().size();
- size_t numResults = getOperation()->getNumResults();
- size_t numOutputs = numResults - k;
-
- for (size_t i = 0; i < numOutputs; ++i) {
- size_t inputIdx = i;
- Value inputVal = (inputIdx < inputs.size()) ? inputs[inputIdx] : inputs[0];
- if (failed(hlo::deriveShapeFromOperand(&builder, getOperation(), inputVal,
- &reifiedReturnShapes))) {
- return failure();
- }
- }
-
- for (auto init : adaptor.getInits()) {
- if (failed(hlo::deriveShapeFromOperand(&builder, getOperation(), init,
- &reifiedReturnShapes))) {
- return failure();
- }
- }
- return success();
-}
-
LogicalResult ScanOp::verify() {
if (getInputs().empty() && getOutputs().empty()) {
return emitOpError() << "at least one of inputs or outputs must be present";
@@ -873,7 +847,7 @@
// Check that the scan dimension is in bounds for all operands. Also check
// that all operands have the same scan dimension size.
int64_t dim = getDimension();
- int64_t dimSize = ShapedType::kDynamic;
+ std::optional<uint64_t> dimSize = getScanDimSize();
for (auto [i, type] : llvm::enumerate(getInputs().getTypes())) {
auto inputType = dyn_cast<RankedTensorType>(type);
if (!inputType) {
@@ -886,13 +860,11 @@
if (inputType.isDynamicDim(dim)) {
continue;
}
- dimSize = inputType.getDimSize(dim);
- if (dimSize != ShapedType::kDynamic &&
- dimSize != inputType.getDimSize(dim)) {
- return emitOpError() << "scan dimension size of operand " << i
- << " does not match previous operands";
+ uint64_t currentDimSize = inputType.getDimSize(dim);
+ if (dimSize.has_value() && dimSize.value() != currentDimSize) {
+ return emitOpError() << "invalid scan dimension size of operand " << i;
}
- dimSize = inputType.getDimSize(dim);
+ dimSize = currentDimSize;
}
Block& bodyBlock = getBody().front();
@@ -936,20 +908,22 @@
parser.parseKeyword("inits") ||
parser.parseOperandList(inits, OpAsmParser::Delimiter::Paren) ||
parser.parseKeyword("dimension") || parser.parseEqual() ||
- parser.parseInteger(dimension) || parser.parseRegion(*body) ||
- parser.parseOptionalAttrDict(result.attributes) ||
- parser.parseColonType(funcType)) {
+ parser.parseInteger(dimension) ||
+ parser.parseOptionalAttrDictWithKeyword(result.attributes) ||
+ parser.parseRegion(*body) || parser.parseColonType(funcType)) {
return failure();
}
- size_t numInputs = inputs.size();
- size_t numCarries = inits.size();
- if (funcType.getInputs().size() != numInputs + numCarries) {
+ int32_t numInputs = inputs.size();
+ int32_t numCarries = inits.size();
+ int32_t numOutputs = funcType.getNumResults() - numCarries;
+ if (static_cast<int64_t>(funcType.getInputs().size()) !=
+ static_cast<int64_t>(numInputs + numCarries)) {
return parser.emitError(
parser.getNameLoc(),
"operand types must match the number of inputs and inits");
}
- if (funcType.getResults().size() < numCarries) {
+ if (numOutputs < 0) {
return parser.emitError(
parser.getNameLoc(),
"not enough result types to cover the required carries");
@@ -968,13 +942,10 @@
Builder& builder = parser.getBuilder();
result.addAttribute(ScanOp::getDimensionAttrName(result.name),
builder.getI64IntegerAttr(dimension));
- result.addAttribute(
- ScanOp::getOperandSegmentSizeAttr(),
- builder.getDenseI32ArrayAttr({(int32_t)numInputs, (int32_t)numCarries}));
- size_t numOutputs = funcType.getNumResults() - numCarries;
- result.addAttribute(
- ScanOp::getResultSegmentSizeAttr(),
- builder.getDenseI32ArrayAttr({(int32_t)numOutputs, (int32_t)numCarries}));
+ result.addAttribute(ScanOp::getOperandSegmentSizeAttr(),
+ builder.getDenseI32ArrayAttr({numInputs, numCarries}));
+ result.addAttribute(ScanOp::getResultSegmentSizeAttr(),
+ builder.getDenseI32ArrayAttr({numOutputs, numCarries}));
return success();
}
@@ -985,10 +956,12 @@
p << ") inits (";
p.printOperands(getInits());
p << ") dimension=" << getDimension() << " ";
+ p.printOptionalAttrDictWithKeyword(
+ getOperation()->getAttrs(),
+ /*elidedAttrs=*/{"dimension", "operandSegmentSizes",
+ "resultSegmentSizes"});
+ p << " ";
p.printRegion(getBody(), /*printEntryBlockArgs=*/true);
- p.printOptionalAttrDict(getOperation()->getAttrs(),
- /*elidedAttrs=*/{"dimension", "operandSegmentSizes",
- "resultSegmentSizes"});
p << " : ";
p.printFunctionalType(*this);
}
diff --git a/stablehlo/dialect/ChloOps.td b/stablehlo/dialect/ChloOps.td
index 7defd25..676e7a0 100644
--- a/stablehlo/dialect/ChloOps.td
+++ b/stablehlo/dialect/ChloOps.td
@@ -938,7 +938,7 @@
def CHLO_ScanOp : CHLO_Op<"scan", [
AttrSizedOperandSegments,
AttrSizedResultSegments,
- InferTensorTypeWithReify,
+ InferTensorType,
IsolatedFromAbove,
OpAsmOpInterface,
RecursiveMemoryEffects,
@@ -961,6 +961,7 @@
Variadic<HLO_AnyTensor>:$inputs,
Variadic<HLO_AnyTensor>:$inits,
ConfinedAttr<I64Attr, [IntNonNegative]>:$dimension,
+ OptionalAttr<ConfinedAttr<I64Attr, [IntNonNegative]>>:$scan_dim_size,
DefaultValuedOptionalAttr<BoolAttr, "false">:$is_reverse,
OptionalAttr<BoolAttr>:$is_associative
);
diff --git a/stablehlo/dialect/TypeInference.cpp b/stablehlo/dialect/TypeInference.cpp
index ca0ba28..8203ea3 100644
--- a/stablehlo/dialect/TypeInference.cpp
+++ b/stablehlo/dialect/TypeInference.cpp
@@ -2110,11 +2110,22 @@
inferredBounds[dim] = inferredDimAndBound.second;
}
}
+ Attribute witnessEncoding;
+ for (auto inputType : inputTypes) {
+ if (auto rankedType = dyn_cast<RankedTensorType>(inputType)) {
+ if (rankedType.getEncoding()) {
+ witnessEncoding = rankedType.getEncoding();
+ break;
+ }
+ }
+ }
+ if (!witnessEncoding) witnessEncoding = witnessType.getEncoding();
+
// concatenate_c5, concatenate_c6
inferredReturnTypes.push_back(RankedTensorType::get(
inferredSizes, witnessType.getElementType(),
boundsToEncoding(
- witnessType.getEncoding(),
+ witnessEncoding,
// Empty array as argument is an indicator to boundsToEncoding()
// that there are no bounds at all in inputs, thus sparsity
// attributes will be included in the return type
diff --git a/stablehlo/dialect/VhloAttrs.td b/stablehlo/dialect/VhloAttrs.td
index a8d3cc0..1d21552 100644
--- a/stablehlo/dialect/VhloAttrs.td
+++ b/stablehlo/dialect/VhloAttrs.td
@@ -127,16 +127,13 @@
}
// Corresponds to TensorConstant from the StableHLO spec.
-def VHLO_TensorDataV1 : AttrParameter<"::llvm::ArrayRef<char>", ""> {
- let allocator = "$_dst = $_allocator.copyInto($_self);";
-}
def VHLO_TensorAttrV1 : VHLO_AttrDef<"TensorV1", "0.9.0", "current"> {
let mnemonic = "tensor_v1";
- let parameters = (ins "::mlir::Type":$type, VHLO_TensorDataV1:$data);
+ let parameters = (ins "::mlir::Type":$type, "DenseElementsAttr":$data);
let genVerifyDecl = 1;
let extraClassDefinition = [{
LogicalResult TensorV1Attr::verify(
- llvm::function_ref<mlir::InFlightDiagnostic ()> errFn, mlir::Type type, ArrayRef<char>) {
+ llvm::function_ref<mlir::InFlightDiagnostic ()> errFn, mlir::Type type, DenseElementsAttr) {
if (!isFromVhlo(type)) errFn() << "expected VHLO type";
return success();
}
diff --git a/stablehlo/dialect/VhloBytecode.cpp b/stablehlo/dialect/VhloBytecode.cpp
index 01e18c7..73fa509 100644
--- a/stablehlo/dialect/VhloBytecode.cpp
+++ b/stablehlo/dialect/VhloBytecode.cpp
@@ -1049,13 +1049,13 @@
static void writeVhloTensorV1Attr(DialectBytecodeWriter& writer,
vhlo::TensorV1Attr attr) {
if (!isBooleanType(attr.getType())) {
- writer.writeOwnedBlob(attr.getData());
+ writer.writeOwnedBlob(attr.getData().getRawData());
return;
}
// Pack the data if i1
SmallVector<char> data;
- ArrayRef<char> rawData = attr.getData();
+ ArrayRef<char> rawData = attr.getData().getRawData();
auto numElements = cast<RankedTensorV1Type>(attr.getType()).getNumElements();
// If the attribute is a splat, we can just splat the value directly.
@@ -1085,7 +1085,18 @@
if (failed(reader.readType(type)) ||
failed(readVhloTensorV1Attr(reader, type, blob)))
return TensorV1Attr();
- return TensorV1Attr::get(getContext(), type, blob);
+ struct VhloToBuiltinPrintConverter : VhloTypeConverter {
+ VhloToBuiltinPrintConverter() : VhloTypeConverter() {
+ addVhloToBuiltinConversions();
+ }
+ Attribute convertEncoding(Attribute attr) const override { return attr; }
+ };
+
+ VhloToBuiltinPrintConverter conv;
+ auto builtinType = cast<ShapedType>(conv.convertType(type));
+ return TensorV1Attr::get(
+ getContext(), type,
+ DenseElementsAttr::getFromRawBuffer(builtinType, blob));
}
void VhloBytecodeInterface::write(TensorV1Attr attr,
diff --git a/stablehlo/dialect/VhloOps.cpp b/stablehlo/dialect/VhloOps.cpp
index 53ed964..886b47a 100644
--- a/stablehlo/dialect/VhloOps.cpp
+++ b/stablehlo/dialect/VhloOps.cpp
@@ -185,12 +185,7 @@
}
void TensorV1Attr::print(mlir::AsmPrinter& odsPrinter) const {
- odsPrinter << '<'
- << DenseTypedElementsAttr::getFromRawBuffer(
- llvm::cast<ShapedType>(
- convertTypeToBuiltinForPrint(getType())),
- getData())
- << '>';
+ odsPrinter << '<' << getData() << '>';
}
// Parse tensor elements using DenseTypedElementsAttr printing.
@@ -201,8 +196,7 @@
return TensorV1Attr();
}
return TensorV1Attr::get(parser.getContext(),
- convertTypeToVhloForParse(attr.getType()),
- attr.getRawData());
+ convertTypeToVhloForParse(attr.getType()), attr);
}
void printEscapedString(AsmPrinter& p, llvm::StringRef value) {
diff --git a/stablehlo/integrations/c/StablehloAttributes.cpp b/stablehlo/integrations/c/StablehloAttributes.cpp
index 3be3dd2..473e316 100644
--- a/stablehlo/integrations/c/StablehloAttributes.cpp
+++ b/stablehlo/integrations/c/StablehloAttributes.cpp
@@ -754,3 +754,136 @@
llvm::cast<mlir::stablehlo::ResultAccuracyAttr>(unwrap(attr)).getMode();
return wrap(modeAttr);
}
+
+//===----------------------------------------------------------------------===//
+// SubAxisInfo
+//===----------------------------------------------------------------------===//
+
+MlirAttribute stablehloSubAxisInfoAttrGet(MlirContext ctx, int64_t preSize,
+ int64_t size) {
+ return wrap(
+ mlir::stablehlo::SubAxisInfoAttr::get(unwrap(ctx), preSize, size));
+}
+
+bool stablehloAttributeIsASubAxisInfoAttr(MlirAttribute attr) {
+ return llvm::isa<mlir::stablehlo::SubAxisInfoAttr>(unwrap(attr));
+}
+
+int64_t stablehloSubAxisInfoAttrGetPreSize(MlirAttribute attr) {
+ return llvm::cast<mlir::stablehlo::SubAxisInfoAttr>(unwrap(attr))
+ .getPreSize();
+}
+
+int64_t stablehloSubAxisInfoAttrGetSize(MlirAttribute attr) {
+ return llvm::cast<mlir::stablehlo::SubAxisInfoAttr>(unwrap(attr)).getSize();
+}
+
+//===----------------------------------------------------------------------===//
+// AxisRef
+//===----------------------------------------------------------------------===//
+
+MlirAttribute stablehloAxisRefAttrGet(MlirContext ctx, MlirStringRef name,
+ MlirAttribute subAxisInfo) {
+ return wrap(mlir::stablehlo::AxisRefAttr::get(
+ unwrap(ctx), unwrap(name),
+ llvm::cast_or_null<mlir::stablehlo::SubAxisInfoAttr>(
+ unwrap(subAxisInfo))));
+}
+
+bool stablehloAttributeIsAnAxisRefAttr(MlirAttribute attr) {
+ return llvm::isa<mlir::stablehlo::AxisRefAttr>(unwrap(attr));
+}
+
+MlirStringRef stablehloAxisRefAttrGetName(MlirAttribute attr) {
+ return wrap(llvm::cast<mlir::stablehlo::AxisRefAttr>(unwrap(attr)).getName());
+}
+
+MlirAttribute stablehloAxisRefAttrGetSubAxisInfo(MlirAttribute attr) {
+ return wrap(
+ llvm::cast<mlir::stablehlo::AxisRefAttr>(unwrap(attr)).getSubAxisInfo());
+}
+
+//===----------------------------------------------------------------------===//
+// ReplicaGroupMeshAxes
+//===----------------------------------------------------------------------===//
+
+MlirAttribute stablehloReplicaGroupMeshAxesAttrGet(MlirContext ctx,
+ MlirAttribute mesh,
+ MlirAttribute axes) {
+ return wrap(mlir::stablehlo::ReplicaGroupMeshAxesAttr::get(
+ unwrap(ctx), unwrap(mesh), llvm::cast<mlir::ArrayAttr>(unwrap(axes))));
+}
+
+bool stablehloAttributeIsAReplicaGroupMeshAxesAttr(MlirAttribute attr) {
+ return llvm::isa<mlir::stablehlo::ReplicaGroupMeshAxesAttr>(unwrap(attr));
+}
+
+MlirAttribute stablehloReplicaGroupMeshAxesAttrGetMesh(MlirAttribute attr) {
+ return wrap(
+ llvm::cast<mlir::stablehlo::ReplicaGroupMeshAxesAttr>(unwrap(attr))
+ .getMesh());
+}
+
+MlirAttribute stablehloReplicaGroupMeshAxesAttrGetAxes(MlirAttribute attr) {
+ return wrap(
+ llvm::cast<mlir::stablehlo::ReplicaGroupMeshAxesAttr>(unwrap(attr))
+ .getAxes());
+}
+
+//===----------------------------------------------------------------------===//
+// MeshAxis
+//===----------------------------------------------------------------------===//
+
+MlirAttribute stablehloMeshAxisAttrGet(MlirContext ctx, MlirStringRef name,
+ int64_t size) {
+ return wrap(
+ mlir::stablehlo::MeshAxisAttr::get(unwrap(ctx), unwrap(name), size));
+}
+
+bool stablehloAttributeIsAMeshAxisAttr(MlirAttribute attr) {
+ return llvm::isa<mlir::stablehlo::MeshAxisAttr>(unwrap(attr));
+}
+
+MlirStringRef stablehloMeshAxisAttrGetName(MlirAttribute attr) {
+ return wrap(
+ llvm::cast<mlir::stablehlo::MeshAxisAttr>(unwrap(attr)).getName());
+}
+
+int64_t stablehloMeshAxisAttrGetSize(MlirAttribute attr) {
+ return llvm::cast<mlir::stablehlo::MeshAxisAttr>(unwrap(attr)).getSize();
+}
+
+//===----------------------------------------------------------------------===//
+// Mesh
+//===----------------------------------------------------------------------===//
+
+MlirAttribute stablehloMeshAttrGet(MlirContext ctx, MlirAttribute axes,
+ MlirAttribute deviceIds) {
+ auto axesAttr = llvm::cast<mlir::ArrayAttr>(unwrap(axes));
+ llvm::SmallVector<mlir::stablehlo::MeshAxisAttr> meshAxes;
+ for (auto attr : axesAttr) {
+ meshAxes.push_back(llvm::cast<mlir::stablehlo::MeshAxisAttr>(attr));
+ }
+ return wrap(mlir::stablehlo::MeshAttr::get(
+ unwrap(ctx), meshAxes,
+ llvm::cast_or_null<mlir::DenseIntElementsAttr>(unwrap(deviceIds))));
+}
+
+bool stablehloAttributeIsAMeshAttr(MlirAttribute attr) {
+ return llvm::isa<mlir::stablehlo::MeshAttr>(unwrap(attr));
+}
+
+MlirAttribute stablehloMeshAttrGetAxes(MlirAttribute attr) {
+ auto meshAttr = llvm::cast<mlir::stablehlo::MeshAttr>(unwrap(attr));
+ llvm::SmallVector<mlir::Attribute> attrs;
+ for (auto axis : meshAttr.getAxes()) {
+ attrs.push_back(axis);
+ }
+ return wrap(mlir::ArrayAttr::get(meshAttr.getContext(), attrs));
+}
+
+MlirAttribute stablehloMeshAttrGetDeviceIds(MlirAttribute attr) {
+ auto deviceIds =
+ llvm::cast<mlir::stablehlo::MeshAttr>(unwrap(attr)).getDeviceIds();
+ return wrap(deviceIds);
+}
diff --git a/stablehlo/integrations/c/StablehloAttributes.h b/stablehlo/integrations/c/StablehloAttributes.h
index 8ea0653..8ac2328 100644
--- a/stablehlo/integrations/c/StablehloAttributes.h
+++ b/stablehlo/integrations/c/StablehloAttributes.h
@@ -413,6 +413,73 @@
MLIR_CAPI_EXPORTED MlirAttribute
stablehloResultAccuracyAttrGetMode(MlirAttribute attr);
+//===----------------------------------------------------------------------===//
+// SubAxisInfo
+//===----------------------------------------------------------------------===//
+
+MLIR_CAPI_EXPORTED MlirAttribute stablehloSubAxisInfoAttrGet(MlirContext ctx,
+ int64_t preSize,
+ int64_t size);
+
+MLIR_CAPI_EXPORTED bool stablehloAttributeIsASubAxisInfoAttr(
+ MlirAttribute attr);
+MLIR_CAPI_EXPORTED int64_t
+stablehloSubAxisInfoAttrGetPreSize(MlirAttribute attr);
+MLIR_CAPI_EXPORTED int64_t stablehloSubAxisInfoAttrGetSize(MlirAttribute attr);
+
+//===----------------------------------------------------------------------===//
+// AxisRef
+//===----------------------------------------------------------------------===//
+
+MLIR_CAPI_EXPORTED MlirAttribute stablehloAxisRefAttrGet(
+ MlirContext ctx, MlirStringRef name, MlirAttribute subAxisInfo);
+
+MLIR_CAPI_EXPORTED bool stablehloAttributeIsAnAxisRefAttr(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirStringRef
+stablehloAxisRefAttrGetName(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirAttribute
+stablehloAxisRefAttrGetSubAxisInfo(MlirAttribute attr);
+
+//===----------------------------------------------------------------------===//
+// ReplicaGroupMeshAxes
+//===----------------------------------------------------------------------===//
+
+MLIR_CAPI_EXPORTED MlirAttribute stablehloReplicaGroupMeshAxesAttrGet(
+ MlirContext ctx, MlirAttribute mesh, MlirAttribute axes);
+
+MLIR_CAPI_EXPORTED bool stablehloAttributeIsAReplicaGroupMeshAxesAttr(
+ MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirAttribute
+stablehloReplicaGroupMeshAxesAttrGetMesh(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirAttribute
+stablehloReplicaGroupMeshAxesAttrGetAxes(MlirAttribute attr);
+
+//===----------------------------------------------------------------------===//
+// MeshAxis
+//===----------------------------------------------------------------------===//
+
+MLIR_CAPI_EXPORTED MlirAttribute stablehloMeshAxisAttrGet(MlirContext ctx,
+ MlirStringRef name,
+ int64_t size);
+
+MLIR_CAPI_EXPORTED bool stablehloAttributeIsAMeshAxisAttr(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirStringRef
+stablehloMeshAxisAttrGetName(MlirAttribute attr);
+MLIR_CAPI_EXPORTED int64_t stablehloMeshAxisAttrGetSize(MlirAttribute attr);
+
+//===----------------------------------------------------------------------===//
+// Mesh
+//===----------------------------------------------------------------------===//
+
+MLIR_CAPI_EXPORTED MlirAttribute stablehloMeshAttrGet(MlirContext ctx,
+ MlirAttribute axes,
+ MlirAttribute deviceIds);
+
+MLIR_CAPI_EXPORTED bool stablehloAttributeIsAMeshAttr(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirAttribute stablehloMeshAttrGetAxes(MlirAttribute attr);
+MLIR_CAPI_EXPORTED MlirAttribute
+stablehloMeshAttrGetDeviceIds(MlirAttribute attr);
+
#ifdef __cplusplus
}
#endif
diff --git a/stablehlo/integrations/python/StablehloModule.cpp b/stablehlo/integrations/python/StablehloModule.cpp
index ecd3975..7cd0fd1 100644
--- a/stablehlo/integrations/python/StablehloModule.cpp
+++ b/stablehlo/integrations/python/StablehloModule.cpp
@@ -663,6 +663,121 @@
});
mlir::python::nanobind_adaptors::mlir_attribute_subclass(
+ m, "SubAxisInfoAttr", stablehloAttributeIsASubAxisInfoAttr)
+ .def_classmethod(
+ "get",
+ [](nb::object cls, int64_t preSize, int64_t size, MlirContext ctx) {
+ return cls(stablehloSubAxisInfoAttrGet(ctx, preSize, size));
+ },
+ nb::arg("cls"), nb::arg("pre_size"), nb::arg("size"),
+ nb::arg("context").none() = nb::none(),
+ "Creates a SubAxisInfoAttr with the given pre_size and size.")
+ .def_property_readonly("pre_size",
+ [](MlirAttribute self) {
+ return stablehloSubAxisInfoAttrGetPreSize(self);
+ })
+ .def_property_readonly("size", [](MlirAttribute self) {
+ return stablehloSubAxisInfoAttrGetSize(self);
+ });
+
+ mlir::python::nanobind_adaptors::mlir_attribute_subclass(
+ m, "AxisRefAttr", stablehloAttributeIsAnAxisRefAttr)
+ .def_classmethod(
+ "get",
+ [](nb::object cls, const std::string& name,
+ std::optional<MlirAttribute> subAxisInfo, MlirContext ctx) {
+ return cls(stablehloAxisRefAttrGet(
+ ctx, mlirStringRefCreate(name.c_str(), name.size()),
+ subAxisInfo.has_value() ? *subAxisInfo
+ : MlirAttribute{nullptr}));
+ },
+ nb::arg("cls"), nb::arg("name"),
+ nb::arg("sub_axis_info").none() = nb::none(),
+ nb::arg("context").none() = nb::none(),
+ "Creates an AxisRefAttr with the given name and optional "
+ "sub_axis_info.")
+ .def_property_readonly(
+ "name",
+ [](MlirAttribute self) {
+ return toPyString(stablehloAxisRefAttrGetName(self));
+ })
+ .def_property_readonly("sub_axis_info", [](MlirAttribute self) {
+ MlirAttribute subAxisInfo = stablehloAxisRefAttrGetSubAxisInfo(self);
+ if (mlirAttributeIsNull(subAxisInfo)) {
+ return nb::cast(std::optional<MlirAttribute>(std::nullopt));
+ }
+ return nb::cast(std::optional<MlirAttribute>(subAxisInfo));
+ });
+
+ mlir::python::nanobind_adaptors::mlir_attribute_subclass(
+ m, "ReplicaGroupMeshAxesAttr",
+ stablehloAttributeIsAReplicaGroupMeshAxesAttr)
+ .def_classmethod(
+ "get",
+ [](nb::object cls, MlirAttribute mesh, MlirAttribute axes,
+ MlirContext ctx) {
+ return cls(stablehloReplicaGroupMeshAxesAttrGet(ctx, mesh, axes));
+ },
+ nb::arg("cls"), nb::arg("mesh"), nb::arg("axes"),
+ nb::arg("context").none() = nb::none(),
+ "Creates a ReplicaGroupMeshAxesAttr with the given mesh and "
+ "axes.")
+ .def_property_readonly(
+ "mesh",
+ [](MlirAttribute self) {
+ return stablehloReplicaGroupMeshAxesAttrGetMesh(self);
+ })
+ .def_property_readonly("axes", [](MlirAttribute self) {
+ return stablehloReplicaGroupMeshAxesAttrGetAxes(self);
+ });
+
+ mlir::python::nanobind_adaptors::mlir_attribute_subclass(
+ m, "MeshAxisAttr", stablehloAttributeIsAMeshAxisAttr)
+ .def_classmethod(
+ "get",
+ [](nb::object cls, const std::string& name, int64_t size,
+ MlirContext ctx) {
+ return cls(stablehloMeshAxisAttrGet(
+ ctx, mlirStringRefCreate(name.c_str(), name.size()), size));
+ },
+ nb::arg("cls"), nb::arg("name"), nb::arg("size"),
+ nb::arg("context").none() = nb::none(),
+ "Creates a MeshAxisAttr with the given name and size.")
+ .def_property_readonly(
+ "name",
+ [](MlirAttribute self) {
+ return toPyString(stablehloMeshAxisAttrGetName(self));
+ })
+ .def_property_readonly("size", [](MlirAttribute self) {
+ return stablehloMeshAxisAttrGetSize(self);
+ });
+
+ mlir::python::nanobind_adaptors::mlir_attribute_subclass(
+ m, "MeshAttr", stablehloAttributeIsAMeshAttr)
+ .def_classmethod(
+ "get",
+ [](nb::object cls, MlirAttribute axes,
+ std::optional<MlirAttribute> deviceIds, MlirContext ctx) {
+ return cls(stablehloMeshAttrGet(
+ ctx, axes,
+ deviceIds.has_value() ? *deviceIds : MlirAttribute{nullptr}));
+ },
+ nb::arg("cls"), nb::arg("axes"),
+ nb::arg("device_ids").none() = nb::none(),
+ nb::arg("context").none() = nb::none(),
+ "Creates a MeshAttr with the given axes and optional device_ids.")
+ .def_property_readonly(
+ "axes",
+ [](MlirAttribute self) { return stablehloMeshAttrGetAxes(self); })
+ .def_property_readonly("device_ids", [](MlirAttribute self) {
+ MlirAttribute deviceIds = stablehloMeshAttrGetDeviceIds(self);
+ if (mlirAttributeIsNull(deviceIds)) {
+ return nb::cast(std::optional<MlirAttribute>(std::nullopt));
+ }
+ return nb::cast(std::optional<MlirAttribute>(deviceIds));
+ });
+
+ mlir::python::nanobind_adaptors::mlir_attribute_subclass(
m, "ResultAccuracyModeAttr", stablehloAttributeIsAResultAccuracyModeAttr)
.def_classmethod(
"get",
diff --git a/stablehlo/integrations/python/tests/stablehlo.py b/stablehlo/integrations/python/tests/stablehlo.py
index ae0cdbd..a91902a 100644
--- a/stablehlo/integrations/python/tests/stablehlo.py
+++ b/stablehlo/integrations/python/tests/stablehlo.py
@@ -264,6 +264,41 @@
@run
+def test_replica_group_mesh_axes_attr():
+ sub_axis = stablehlo.SubAxisInfoAttr.get(pre_size=2, size=4)
+ assert sub_axis is not None
+ assert str(sub_axis) == "#stablehlo<sub_axis_info(2)4>"
+ assert sub_axis.pre_size == 2
+ assert sub_axis.size == 4
+
+ axis_ref = stablehlo.AxisRefAttr.get(name="x", sub_axis_info=sub_axis)
+ assert axis_ref is not None
+ assert (
+ str(axis_ref) == '#stablehlo.axis_ref<name = "x", sub_axis_info = (2)4>'
+ )
+ assert axis_ref.name == "x"
+ assert axis_ref.sub_axis_info == sub_axis
+
+ axis_ref_simple = stablehlo.AxisRefAttr.get(name="y")
+ assert axis_ref_simple is not None
+ assert str(axis_ref_simple) == '#stablehlo.axis_ref<name = "y">'
+ assert axis_ref_simple.name == "y"
+
+ mesh_name = ir.FlatSymbolRefAttr.get("mesh")
+ axes = ir.ArrayAttr.get([axis_ref, axis_ref_simple])
+ rg_attr = stablehlo.ReplicaGroupMeshAxesAttr.get(mesh=mesh_name, axes=axes)
+ assert rg_attr is not None
+ assert rg_attr.mesh == mesh_name
+ assert rg_attr.axes == axes
+ assert (
+ str(rg_attr)
+ == "#stablehlo.replica_group_mesh_axes<mesh = @mesh, axes ="
+ ' [#stablehlo.axis_ref<name = "x", sub_axis_info = (2)4>,'
+ ' #stablehlo.axis_ref<name = "y">]>'
+ )
+
+
+@run
def test_api_version():
api_version = stablehlo.get_api_version()
assert type(api_version) == int
@@ -271,7 +306,7 @@
def is_semver_format(version_str):
- return re.match("^\d+\.\d+\.\d+$", version_str)
+ return re.match(r"^\d+\.\d+\.\d+$", version_str)
@run
@@ -462,17 +497,18 @@
metadata_file
), f"Metadata file not found: {metadata_file}"
+
# Disabling to avoid jax dependency.
-#@run
-#def test_reference_api_with_probe():
+# @run
+# def test_reference_api_with_probe():
# from jax.interpreters import mlir
# import jax.numpy as jnp
# """Tests that probe files are created in the specified directory."""
# _run_probe_test("f32", np.asarray(2, np.float32))
# Disabling to avoid jax dependency.
-#@run
-#def test_reference_api_with_probe_bf16():
+# @run
+# def test_reference_api_with_probe_bf16():
# from jax.interpreters import mlir
# import jax.numpy as jnp
# """Tests that probe files are created for bf16 tensors."""
diff --git a/stablehlo/tests/infer_stablehlo.mlir b/stablehlo/tests/infer_stablehlo.mlir
index fea50ae..0730740 100644
--- a/stablehlo/tests/infer_stablehlo.mlir
+++ b/stablehlo/tests/infer_stablehlo.mlir
@@ -2076,3 +2076,15 @@
%1 = "hlo_test_infer.reify_return_type_shapes"(%0) : (tensor<?x?x?x?xf32>) -> tensor<4xindex>
func.return %1 : tensor<4xindex>
}
+
+// CHECK-LABEL: @concat_bounds_mixed_static_first
+func.func @concat_bounds_mixed_static_first(
+ %arg0: tensor<5x2xi32>,
+ %arg1: tensor<5x?xi32, #stablehlo.bounds<?, 4>>) -> tensor<?x?xindex> {
+ %result = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 1 : i64 } : (
+ tensor<5x2xi32>,
+ tensor<5x?xi32, #stablehlo.bounds<?, 4>>) -> tensor<?x?xi32>
+ // CHECK: types0 = tensor<5x?xi32, #stablehlo.bounds<?, 6>>
+ %1 = "hlo_test_infer.get_return_types"(%result) : (tensor<?x?xi32>) -> tensor<?x?xindex>
+ func.return %1 : tensor<?x?xindex>
+}
diff --git a/stablehlo/tests/ops_chlo.mlir b/stablehlo/tests/ops_chlo.mlir
index 0886643..75d6f0a 100644
--- a/stablehlo/tests/ops_chlo.mlir
+++ b/stablehlo/tests/ops_chlo.mlir
@@ -423,7 +423,49 @@
// CHECK-LABEL: func @scan
func.func @scan(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
// CHECK: chlo.scan
- %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 {
+ %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 attributes {} {
+ ^bb0(%input0: tensor<3xf32>, %carry0: tensor<3xf32>):
+ %2 = stablehlo.add %input0, %carry0 : tensor<3xf32>
+ stablehlo.return %2, %2 : tensor<3xf32>, tensor<3xf32>
+ } : (tensor<2x3xf32>, tensor<3xf32>) -> (tensor<2x3xf32>, tensor<3xf32>)
+ func.return %0 : tensor<2x3xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @scan_explicit_size
+func.func @scan_explicit_size(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
+ // CHECK: chlo.scan
+ // CHECK: scan_dim_size = 2 : i64
+ %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 attributes { scan_dim_size = 2 : i64 } {
+ ^bb0(%input0: tensor<3xf32>, %carry0: tensor<3xf32>):
+ %2 = stablehlo.add %input0, %carry0 : tensor<3xf32>
+ stablehlo.return %2, %2 : tensor<3xf32>, tensor<3xf32>
+ } : (tensor<2x3xf32>, tensor<3xf32>) -> (tensor<2x3xf32>, tensor<3xf32>)
+ func.return %0 : tensor<2x3xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @scan_is_reverse
+func.func @scan_is_reverse(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
+ // CHECK: chlo.scan
+ // CHECK: is_reverse = true
+ %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 attributes { is_reverse = true } {
+ ^bb0(%input0: tensor<3xf32>, %carry0: tensor<3xf32>):
+ %2 = stablehlo.add %input0, %carry0 : tensor<3xf32>
+ stablehlo.return %2, %2 : tensor<3xf32>, tensor<3xf32>
+ } : (tensor<2x3xf32>, tensor<3xf32>) -> (tensor<2x3xf32>, tensor<3xf32>)
+ func.return %0 : tensor<2x3xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @scan_is_associative
+func.func @scan_is_associative(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
+ // CHECK: chlo.scan
+ // CHECK: is_associative = true
+ %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 attributes { is_associative = true } {
^bb0(%input0: tensor<3xf32>, %carry0: tensor<3xf32>):
%2 = stablehlo.add %input0, %carry0 : tensor<3xf32>
stablehlo.return %2, %2 : tensor<3xf32>, tensor<3xf32>
@@ -447,6 +489,18 @@
// -----
+func.func @scan_explicit_size_mismatch(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
+ // expected-error @+1 {{'chlo.scan' op invalid scan dimension size of operand 0}}
+ %0, %1 = chlo.scan (%arg0) inits (%arg1) dimension = 0 attributes { scan_dim_size = 3 : i64 } {
+ ^bb0(%input0: tensor<3xf32>, %carry0: tensor<3xf32>):
+ %2 = stablehlo.add %input0, %carry0 : tensor<3xf32>
+ stablehlo.return %2, %2 : tensor<3xf32>, tensor<3xf32>
+ } : (tensor<2x3xf32>, tensor<3xf32>) -> (tensor<2x3xf32>, tensor<3xf32>)
+ func.return %0 : tensor<2x3xf32>
+}
+
+// -----
+
func.func @scan_size_mismatch(%arg0: tensor<2x3xf32>, %arg1: tensor<3xf32>) -> tensor<2x3xf32> {
// expected-error @+1 {{'chlo.scan' op expects 1 arguments in the body, but got 2}}
%0 = chlo.scan(%arg0) inits () dimension = 0 {
diff --git a/stablehlo/tests/ops_stablehlo_quantized.mlir b/stablehlo/tests/ops_stablehlo_quantized.mlir
index f15d073..1de2d2c 100644
--- a/stablehlo/tests/ops_stablehlo_quantized.mlir
+++ b/stablehlo/tests/ops_stablehlo_quantized.mlir
@@ -437,7 +437,7 @@
// -----
func.func @negative_clamp_quantization(%arg0: tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>) -> tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>> {
- // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}}, but got 'tensor<1x!quant.uniform<u8:f32:0, {1.000000e-01:-30}>>'}}
+ // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}}, but got 'tensor<1x!quant.uniform<ui8:f32:0, {1.000000e-01:-30}>>'}}
%0 = "stablehlo.clamp"(%arg0, %arg0, %arg0) : (tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>, tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>, tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>) -> tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>
func.return %0: tensor<1x!quant.uniform<ui8:f32:0, {0.1:-30}>>
}
@@ -888,7 +888,7 @@
// -----
func.func @illegal_storage_type_for_quantized_element_type(%arg0: tensor<4x!quant.uniform<si8:f32, 1.000000e+00>>) -> tensor<4xf32> {
- // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}}, but got 'tensor<4x!quant.uniform<i8:f32, 1.000000e+00>>}}
+ // expected-error-re@+1 {{operand #0 must be ranked tensor of {{.*}}, but got 'tensor<4x!quant.uniform<si8:f32, 1.000000e+00>>}}
%0 = "stablehlo.uniform_dequantize"(%arg0) : (tensor<4x!quant.uniform<si8:f32, 1.000000e+00>>) -> tensor<4xf32>
func.return %0 : tensor<4xf32>
}
diff --git a/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir b/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
index fba9aed..912b229 100644
--- a/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
+++ b/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
@@ -3,8 +3,8 @@
// -----
// CHECK-LABEL @compose_quantized_abs_op
-// CHECK: %[[abs0:.*]] = stablehlo.abs %arg0 : tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %[[abs0]] : tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK: %[[abs0:.*]] = stablehlo.abs %arg0 : tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %[[abs0]] : tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @compose_quantized_abs_op(%arg0: tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>> {
%0 = stablehlo.uniform_dequantize %arg0 : (tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>) -> tensor<16x16xf32>
%1 = stablehlo.abs %0 : tensor<16x16xf32>
@@ -15,8 +15,8 @@
// -----
// CHECK-LABEL @failed_to_match_uniform_quant_op_operand_not_defined_by_op
-// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %0 : tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %0 : tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @failed_to_match_uniform_quant_op_operand_not_defined_by_op(%arg0: tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>> {
%0 = stablehlo.uniform_quantize %arg0 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
func.return %0 : tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
@@ -30,8 +30,8 @@
// CHECK-NEXT: %2 = stablehlo.add %arg1, %arg2 : tensor<f32>
// CHECK-NEXT: stablehlo.return %2 : tensor<f32>
// CHECK-NEXT: }) : (tensor<4xf32>) -> tensor<4xf32>
-// CHECK-NEXT: %1 = stablehlo.uniform_quantize %0 : (tensor<4xf32>) -> tensor<4x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %1 : tensor<4x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %1 = stablehlo.uniform_quantize %0 : (tensor<4xf32>) -> tensor<4x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %1 : tensor<4x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @failed_to_match_op_with_region(%operand0 : tensor<4xf32>) -> (tensor<4x!quant.uniform<ui8:f32, 34.0:16>>) {
%0 = stablehlo.uniform_quantize %operand0 : (tensor<4xf32>) -> tensor<4x!quant.uniform<ui8:f32, 34.0:16>>
@@ -51,13 +51,13 @@
// -----
// CHECK-LABEL failed_to_match_varidic_op
-// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<8x2xf32>) -> tensor<8x2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<8x2x!quant.uniform<u8:f32, 3.400000e+01:16>>) -> tensor<8x2xf32>
-// CHECK-NEXT: %2 = stablehlo.uniform_quantize %arg1 : (tensor<2x2xf32>) -> tensor<2x2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: %3 = stablehlo.uniform_dequantize %2 : (tensor<2x2x!quant.uniform<u8:f32, 3.400000e+01:16>>) -> tensor<2x2xf32>
+// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<8x2xf32>) -> tensor<8x2x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<8x2x!quant.uniform<ui8:f32, 3.400000e+01:16>>) -> tensor<8x2xf32>
+// CHECK-NEXT: %2 = stablehlo.uniform_quantize %arg1 : (tensor<2x2xf32>) -> tensor<2x2x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %3 = stablehlo.uniform_dequantize %2 : (tensor<2x2x!quant.uniform<ui8:f32, 3.400000e+01:16>>) -> tensor<2x2xf32>
// CHECK-NEXT: %4:2 = "stablehlo.all_gather"(%1, %3) {{.*}} : (tensor<8x2xf32>, tensor<2x2xf32>) -> (tensor<8x8xf32>, tensor<2x4xf32>)
-// CHECK-NEXT: %5 = stablehlo.uniform_quantize %4#0 : (tensor<8x8xf32>) -> tensor<8x8x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %5, %4#1 : tensor<8x8x!quant.uniform<u8:f32, 3.400000e+01:16>>, tensor<2x4xf32>
+// CHECK-NEXT: %5 = stablehlo.uniform_quantize %4#0 : (tensor<8x8xf32>) -> tensor<8x8x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %5, %4#1 : tensor<8x8x!quant.uniform<ui8:f32, 3.400000e+01:16>>, tensor<2x4xf32>
func.func @failed_to_match_varidic_op(%arg0: tensor<8x2xf32>, %arg1: tensor<2x2xf32>) -> (tensor<8x8x!quant.uniform<ui8:f32, 34.0:16>>, tensor<2x4xf32>) {
%0 = stablehlo.uniform_quantize %arg0 : (tensor<8x2xf32>) -> tensor<8x2x!quant.uniform<ui8:f32, 34.0:16>>
%1 = stablehlo.uniform_dequantize %0 : (tensor<8x2x!quant.uniform<ui8:f32, 34.0:16>>) -> tensor<8x2xf32>
@@ -75,12 +75,12 @@
// -----
// CHECK-LABEL @failed_to_match_operand_of_compute_op_already_quantized
-// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x8x8x207xf32>) -> tensor<1x8x8x207x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<1x8x8x207x!quant.uniform<u8:f32, 3.400000e+01:16>>) -> tensor<1x8x8x207xf32>
+// CHECK: %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x8x8x207xf32>) -> tensor<1x8x8x207x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<1x8x8x207x!quant.uniform<ui8:f32, 3.400000e+01:16>>) -> tensor<1x8x8x207xf32>
// CHECK-NEXT: %2 = stablehlo.abs %arg1 : tensor<3x3x207x16x!quant.uniform<i8:f32, 5.000000e+00:20>>
// CHECK-NEXT: %3 = stablehlo.convolution(%1, %2) {{.*}} : (tensor<1x8x8x207xf32>, tensor<3x3x207x16x!quant.uniform<i8:f32, 5.000000e+00:20>>) -> tensor<1x8x8x16xf32>
-// CHECK-NEXT: %4 = stablehlo.uniform_quantize %3 : (tensor<1x8x8x16xf32>) -> tensor<1x8x8x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %4 : tensor<1x8x8x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %4 = stablehlo.uniform_quantize %3 : (tensor<1x8x8x16xf32>) -> tensor<1x8x8x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %4 : tensor<1x8x8x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @failed_to_match_operand_of_compute_op_already_quantized(%arg0: tensor<1x8x8x207xf32>, %arg1: tensor<3x3x207x16x!quant.uniform<i8:f32, 5.0:20>>) -> tensor<1x8x8x16x!quant.uniform<ui8:f32, 34.0:16>> {
%0 = stablehlo.uniform_quantize %arg0 : (tensor<1x8x8x207xf32>) -> tensor<1x8x8x207x!quant.uniform<ui8:f32, 34.0:16>>
%1 = stablehlo.uniform_dequantize %0 : (tensor<1x8x8x207x!quant.uniform<ui8:f32, 34.0:16>>) -> tensor<1x8x8x207xf32>
@@ -97,11 +97,11 @@
// -----
// CHECK-LABEL @failed_to_match_operand_not_defined_by_op
-// CHECK: %0 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>) -> tensor<16x16xf32>
+// CHECK: %0 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %1 = stablehlo.uniform_dequantize %0 : (tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>) -> tensor<16x16xf32>
// CHECK-NEXT: %2 = stablehlo.add %arg0, %1 : tensor<16x16xf32>
-// CHECK-NEXT: %3 = stablehlo.uniform_quantize %2 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %3 : tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %3 = stablehlo.uniform_quantize %2 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %3 : tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @failed_to_match_operand_not_defined_by_op(%arg0: tensor<16x16xf32>, %arg1: tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>> {
%1 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
%2 = stablehlo.uniform_dequantize %1 : (tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>) -> tensor<16x16xf32>
@@ -114,11 +114,11 @@
// CHECK-LABEL @failed_to_match_defining_op_is_not_a_uniform_dequantized_op
// CHECK: %0 = stablehlo.abs %arg0 : tensor<16x16xf32>
-// CHECK-NEXT: %1 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: %2 = stablehlo.uniform_dequantize %1 : (tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>) -> tensor<16x16xf32>
+// CHECK-NEXT: %1 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %2 = stablehlo.uniform_dequantize %1 : (tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>) -> tensor<16x16xf32>
// CHECK-NEXT: %3 = stablehlo.add %0, %2 : tensor<16x16xf32>
-// CHECK-NEXT: %4 = stablehlo.uniform_quantize %3 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
-// CHECK-NEXT: return %4 : tensor<16x16x!quant.uniform<u8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: %4 = stablehlo.uniform_quantize %3 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
+// CHECK-NEXT: return %4 : tensor<16x16x!quant.uniform<ui8:f32, 3.400000e+01:16>>
func.func @failed_to_match_defining_op_is_not_a_uniform_dequantized_op(%arg0: tensor<16x16xf32>, %arg1: tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>> {
%0 = stablehlo.abs %arg0 : tensor<16x16xf32>
%1 = stablehlo.uniform_quantize %arg1 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
diff --git a/stablehlo/transforms/StablehloLegalizeToVhlo.cpp b/stablehlo/transforms/StablehloLegalizeToVhlo.cpp
index 372535e..a00cac5 100644
--- a/stablehlo/transforms/StablehloLegalizeToVhlo.cpp
+++ b/stablehlo/transforms/StablehloLegalizeToVhlo.cpp
@@ -252,8 +252,8 @@
auto vhloType = typeConverter->convertType(attr.getType());
LLVM_DEBUG(llvm::dbgs() << "Converted " << vhloType << '\n');
if (!vhloType) return {};
- return vhlo::TensorV1Attr::get(attr.getContext(), vhloType,
- attr.getRawData());
+
+ return vhlo::TensorV1Attr::get(attr.getContext(), vhloType, attr);
}
if (auto attr = dyn_cast<DenseI64ArrayAttr>(stablehloAttr)) {
// Leverage the serialization for DenseElements.
diff --git a/stablehlo/transforms/VhloLegalizeToStablehlo.cpp b/stablehlo/transforms/VhloLegalizeToStablehlo.cpp
index 5f9d550..d3a0f9d 100644
--- a/stablehlo/transforms/VhloLegalizeToStablehlo.cpp
+++ b/stablehlo/transforms/VhloLegalizeToStablehlo.cpp
@@ -25,6 +25,7 @@
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/AttrTypeSubElements.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinOps.h"
@@ -173,8 +174,7 @@
auto builtinType =
cast<ShapedType>(typeConverter->convertType(attr.getType()));
if (!builtinType) return {};
- return DenseTypedElementsAttr::getFromRawBuffer(builtinType,
- attr.getData());
+ return attr.getData();
}
if (auto attr = dyn_cast<vhlo::TransposeV1Attr>(vhloAttr)) {
RETURN_CONVERTED_ENUM_ATTR(Transpose, V1);
@@ -619,7 +619,8 @@
typeConverter->convertType(tensorAttr.getType()));
if (!type) return specialFailure();
- auto elems = DenseElementsAttr::getFromRawBuffer(type, tensorAttr.getData());
+ auto elems = DenseElementsAttr::getFromRawBuffer(
+ type, tensorAttr.getData().getRawData());
stablehloAttrs.emplace_back(
vhloName, DenseArrayAttr::get(vhloAttr.getContext(),
@@ -806,7 +807,7 @@
bool isEmptyTensor(Attribute vhloAttr) {
auto attr = dyn_cast_or_null<vhlo::TensorV1Attr>(vhloAttr);
- return attr && attr.getData().empty();
+ return attr && attr.getData().getRawData().empty();
}
bool isEnum(Attribute vhloAttr, Attribute value) { return vhloAttr == value; }
@@ -1158,10 +1159,28 @@
}
void runOnOperation() override {
- // Upgraded VHLO should always be convertible to StableHLO.
- // Arbitrary VHLO might not be convertible if it uses deprecated features
- // which are no longer available in StableHLO.
- if (failed(applyPartialConversion(getOperation(), *target, patterns))) {
+ // Upgraded VHLO types should always be convertible to StableHLO.
+ mlir::AttrTypeReplacer replacer;
+ replacer.addReplacement([&](Type type) -> std::optional<Type> {
+ Type newType = converter.convertType(type);
+ if (newType && newType != type) return newType;
+ return std::nullopt;
+ });
+ // Only types are updated here, as attribute creation with StableHLO types
+ // would trigger VHLO assertion. This is used to avoid need for signature
+ // conversion which results in invalidating the blocks and triggering
+ // recomputing the order of operations, as well as inserting many
+ // unrealized conversion casts.
+ replacer.recursivelyReplaceElementsIn(getOperation(),
+ /*replaceAttrs=*/false,
+ /*replaceLocs=*/false,
+ /*replaceTypes=*/true);
+
+ ConversionConfig config;
+ config.allowPatternRollback = false;
+ config.foldingMode = mlir::DialectConversionFoldingMode::Never;
+ if (failed(applyPartialConversion(getOperation(), *target, patterns,
+ config))) {
return signalPassFailure();
}
diff --git a/stablehlo/transforms/VhloToVersion.cpp b/stablehlo/transforms/VhloToVersion.cpp
index c3c15d2..1d0fea2 100644
--- a/stablehlo/transforms/VhloToVersion.cpp
+++ b/stablehlo/transforms/VhloToVersion.cpp
@@ -365,12 +365,15 @@
auto shape = vhlo::RankedTensorV1Type::get(
builder.getContext(), {0},
vhlo::IntegerSI64V1Type::get(builder.getContext()), {});
- return vhlo::TensorV1Attr::get(builder.getContext(), shape, {});
+ auto denseElements = DenseIntElementsAttr::get(
+ RankedTensorType::get({0}, builder.getI64Type()),
+ llvm::ArrayRef<int64_t>{});
+ return vhlo::TensorV1Attr::get(builder.getContext(), shape, denseElements);
}
bool isEmptyTensor(Attribute attr) {
auto tensor = dyn_cast<TensorV1Attr>(attr);
- if (tensor) return tensor.getData().empty();
+ if (tensor) return tensor.getData().getRawData().empty();
return false;
}
@@ -401,7 +404,7 @@
RankedTensorV1Type::get(builder.getContext(), paddingShape,
IntegerSI64V1Type::get(builder.getContext()),
nullptr),
- denseElements.getRawData());
+ denseElements);
}
bool isDefaultResultAccuracy(Attribute attr) {
diff --git a/stablehlo/transforms/optimization/StablehloAggressiveFolder.cpp b/stablehlo/transforms/optimization/StablehloAggressiveFolder.cpp
index 2dac7e3..622f23c 100644
--- a/stablehlo/transforms/optimization/StablehloAggressiveFolder.cpp
+++ b/stablehlo/transforms/optimization/StablehloAggressiveFolder.cpp
@@ -209,11 +209,13 @@
Attribute res;
if (isa<IntegerType>(elemTy))
- res = constFoldBinaryOp<IntegerAttr, IntegerAttr::ValueType, void,
- IntResultType>(operands, resultType, folder);
+ res = constFoldBinaryOp<IntegerAttr, IntegerAttr, IntegerAttr::ValueType,
+ IntegerAttr::ValueType, void, IntResultType>(
+ operands, resultType, folder);
if (isa<FloatType>(elemTy))
- res = constFoldBinaryOp<FloatAttr, FloatAttr::ValueType, void,
- FloatResultType>(operands, resultType, folder);
+ res = constFoldBinaryOp<FloatAttr, FloatAttr, FloatAttr::ValueType,
+ FloatAttr::ValueType, void, FloatResultType>(
+ operands, resultType, folder);
if (res) return cast<TypedAttr>(res);
return nullptr;
diff --git a/third_party/llvm/build.patch b/third_party/llvm/build.patch
index a886a42..70c0216 100644
--- a/third_party/llvm/build.patch
+++ b/third_party/llvm/build.patch
@@ -1,5 +1,5 @@
diff --git a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel
-index a7e652c..5b8ac5e 100644
+index 3df64d8a483f..aa3643bd4a2c 100644
--- a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel
@@ -234,19 +234,19 @@ write_file(
@@ -34,7 +34,7 @@
)
BLAKE3_x86_64_ASM_SOURCE_PATTERNS = [
-@@ -308,7 +308,8 @@ cc_library(
+@@ -307,7 +307,8 @@ cc_library(
"@platforms//cpu:aarch64": [
"lib/Support/BLAKE3/blake3_neon.c",
],
@@ -44,7 +44,7 @@
"//conditions:default": [
],
}),
-@@ -337,8 +338,9 @@ cc_library(
+@@ -336,8 +337,9 @@ cc_library(
],
"//conditions:default": ["BLAKE3_USE_NEON=0"],
}) + select({
@@ -56,3 +56,17 @@
"//conditions:default": [
"BLAKE3_NO_AVX2",
"BLAKE3_NO_AVX512",
+@@ -4679,6 +4681,13 @@ cc_library(
+ srcs = glob(["lib/HTTP/*.cpp"]),
+ hdrs = glob(["include/llvm/HTTP/*.h"]),
+ includes = ["include"],
++ linkopts = select({
++ "@platforms//os:windows": [
++ "crypt32.lib",
++ "winhttp.lib",
++ ],
++ "//conditions:default": [],
++ }),
+ deps = [":Support"],
+ )
+
diff --git a/third_party/llvm/generated.patch b/third_party/llvm/generated.patch
index c6a2aaa..509398d 100644
--- a/third_party/llvm/generated.patch
+++ b/third_party/llvm/generated.patch
@@ -1,293 +1 @@
Auto generated patch. Do not edit or delete it, even if empty.
-diff -ruN --strip-trailing-cr a/llvm/lib/Transforms/Vectorize/VPlanTransforms.cpp b/llvm/lib/Transforms/Vectorize/VPlanTransforms.cpp
---- a/llvm/lib/Transforms/Vectorize/VPlanTransforms.cpp
-+++ b/llvm/lib/Transforms/Vectorize/VPlanTransforms.cpp
-@@ -5882,21 +5882,26 @@
- BackedgeVal,
- match_fn(m_VPInstruction<VPInstruction::ComputeReductionResult>())));
-
-- // If IV is an expression to sink, create a new select that tracks the
-- // underlying IV directly and update the backedge value.
-- if (IVOfExpressionToSink) {
-- auto *OrigFindLastSelectR = FindLastSelect->getDefiningRecipe();
-- VPBuilder LoopBuilder(OrigFindLastSelectR);
-- DebugLoc DL = OrigFindLastSelectR->getDebugLoc();
-- VPValue *SelectCond = OrigFindLastSelectR->getOperand(0);
-- if (OrigFindLastSelectR->getOperand(1) == PhiR) {
-- FindLastSelect = LoopBuilder.createSelect(SelectCond, PhiR,
-- IVOfExpressionToSink, DL);
-- } else {
-- assert(OrigFindLastSelectR->getOperand(2) == PhiR &&
-- "PhiR expected as operand 1 or 2");
-- FindLastSelect = LoopBuilder.createSelect(
-- SelectCond, IVOfExpressionToSink, PhiR, DL);
-+ VPValue *NewFindLastSelect = BackedgeVal;
-+ VPValue *SelectCond = Cond;
-+ if (!SentinelVal || IVOfExpressionToSink) {
-+ // When we need to create a new select, normalize the condition so that
-+ // PhiR is the last operand and include the header mask if needed.
-+ DebugLoc DL = FindLastSelect->getDefiningRecipe()->getDebugLoc();
-+ VPBuilder LoopBuilder(FindLastSelect->getDefiningRecipe());
-+ if (FindLastSelect->getDefiningRecipe()->getOperand(1) == PhiR)
-+ SelectCond = LoopBuilder.createNot(SelectCond);
-+
-+ // When tail folding, mask the condition with the header mask to prevent
-+ // propagating poison from inactive lanes in the last vector iteration.
-+ if (HeaderMask)
-+ SelectCond = LoopBuilder.createLogicalAnd(HeaderMask, SelectCond);
-+
-+ if (SelectCond != Cond || IVOfExpressionToSink) {
-+ NewFindLastSelect = LoopBuilder.createSelect(
-+ SelectCond,
-+ IVOfExpressionToSink ? IVOfExpressionToSink : FindLastExpression,
-+ PhiR, DL);
- }
- }
-
-@@ -5908,8 +5913,9 @@
- FastMathFlags());
- DebugLoc ExitDL = RdxResult->getDebugLoc();
- VPBuilder MiddleBuilder(RdxResult);
-- VPValue *ReducedIV = MiddleBuilder.createNaryOp(
-- VPInstruction::ComputeReductionResult, FindLastSelect, Flags, ExitDL);
-+ VPValue *ReducedIV =
-+ MiddleBuilder.createNaryOp(VPInstruction::ComputeReductionResult,
-+ NewFindLastSelect, Flags, ExitDL);
-
- // If IVOfExpressionToSink is an expression to sink, sink it now.
- VPValue *VectorRegionExitingVal = ReducedIV;
-@@ -5939,10 +5945,7 @@
- AnyOfPhi->insertAfter(PhiR);
-
- VPBuilder LoopBuilder(BackedgeVal->getDefiningRecipe());
-- VPValue *AnyOfCond = Cond;
-- if (FindLastSelect->getOperand(1) == PhiR)
-- AnyOfCond = LoopBuilder.createNot(Cond);
-- VPValue *OrVal = LoopBuilder.createOr(AnyOfPhi, AnyOfCond);
-+ VPValue *OrVal = LoopBuilder.createOr(AnyOfPhi, SelectCond);
- AnyOfPhi->setOperand(1, OrVal);
-
- VPIRFlags OrFlags(RecurKind::Or, /*IsOrdered=*/false,
-@@ -5963,7 +5966,7 @@
-
- auto *NewPhiR = new VPReductionPHIRecipe(
- cast<PHINode>(PhiR->getUnderlyingInstr()), RecurKind::FindIV, *StartVPV,
-- *FindLastSelect, RdxUnordered{1}, {},
-+ *NewFindLastSelect, RdxUnordered{1}, {},
- PhiR->hasUsesOutsideReductionChain());
- NewPhiR->insertBefore(PhiR);
- PhiR->replaceAllUsesWith(NewPhiR);
-diff -ruN --strip-trailing-cr a/llvm/test/Transforms/LoopVectorize/find-last-iv-sinkable-expr.ll b/llvm/test/Transforms/LoopVectorize/find-last-iv-sinkable-expr.ll
---- a/llvm/test/Transforms/LoopVectorize/find-last-iv-sinkable-expr.ll
-+++ b/llvm/test/Transforms/LoopVectorize/find-last-iv-sinkable-expr.ll
-@@ -73,7 +73,8 @@
- ; TF-NEXT: [[TMP1:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[INDEX]]
- ; TF-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP1]], <4 x i1> [[TMP0]], <4 x i64> poison)
- ; TF-NEXT: [[TMP2:%.*]] = icmp eq <4 x i64> [[WIDE_MASKED_LOAD]], splat (i64 42)
--; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP2]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP9:%.*]] = select <4 x i1> [[TMP0]], <4 x i1> [[TMP2]], <4 x i1> zeroinitializer
-+; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP9]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], splat (i64 4)
- ; TF-NEXT: [[TMP4:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-@@ -174,7 +175,8 @@
- ; TF-NEXT: [[TMP1:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[INDEX]]
- ; TF-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP1]], <4 x i1> [[TMP0]], <4 x i64> poison)
- ; TF-NEXT: [[TMP2:%.*]] = icmp eq <4 x i64> [[WIDE_MASKED_LOAD]], splat (i64 42)
--; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP2]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP9:%.*]] = select <4 x i1> [[TMP0]], <4 x i1> [[TMP2]], <4 x i1> zeroinitializer
-+; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP9]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], splat (i64 4)
- ; TF-NEXT: [[TMP4:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-@@ -392,7 +394,8 @@
- ; TF-NEXT: [[TMP1:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[INDEX]]
- ; TF-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP1]], <4 x i1> [[TMP0]], <4 x i64> poison)
- ; TF-NEXT: [[TMP2:%.*]] = icmp eq <4 x i64> [[WIDE_MASKED_LOAD]], splat (i64 42)
--; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP2]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP9:%.*]] = select <4 x i1> [[TMP0]], <4 x i1> [[TMP2]], <4 x i1> zeroinitializer
-+; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP9]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], splat (i64 4)
- ; TF-NEXT: [[TMP4:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-@@ -513,7 +516,7 @@
- ; TF: [[VECTOR_BODY]]:
- ; TF-NEXT: [[INDEX:%.*]] = phi i64 [ 0, %[[VECTOR_PH]] ], [ [[INDEX_NEXT:%.*]], %[[PRED_LOAD_CONTINUE8:.*]] ]
- ; TF-NEXT: [[VEC_IND:%.*]] = phi <4 x i64> [ <i64 10, i64 12, i64 14, i64 16>, %[[VECTOR_PH]] ], [ [[VEC_IND_NEXT:%.*]], %[[PRED_LOAD_CONTINUE8]] ]
--; TF-NEXT: [[VEC_PHI:%.*]] = phi <4 x i64> [ zeroinitializer, %[[VECTOR_PH]] ], [ [[TMP31:%.*]], %[[PRED_LOAD_CONTINUE8]] ]
-+; TF-NEXT: [[VEC_PHI:%.*]] = phi <4 x i64> [ zeroinitializer, %[[VECTOR_PH]] ], [ [[TMP36:%.*]], %[[PRED_LOAD_CONTINUE8]] ]
- ; TF-NEXT: [[TMP3:%.*]] = mul i64 [[INDEX]], 2
- ; TF-NEXT: [[OFFSET_IDX:%.*]] = add i64 10, [[TMP3]]
- ; TF-NEXT: [[BROADCAST_SPLATINSERT1:%.*]] = insertelement <4 x i64> poison, i64 [[INDEX]], i64 0
-@@ -562,13 +565,14 @@
- ; TF-NEXT: [[TMP28:%.*]] = phi <4 x i64> [ [[TMP22]], %[[PRED_LOAD_CONTINUE6]] ], [ [[TMP27]], %[[PRED_LOAD_IF7]] ]
- ; TF-NEXT: [[TMP29:%.*]] = icmp eq <4 x i64> [[TMP28]], splat (i64 42)
- ; TF-NEXT: [[TMP30:%.*]] = add <4 x i64> [[VEC_IND]], splat (i64 100)
--; TF-NEXT: [[TMP31]] = select <4 x i1> [[TMP29]], <4 x i64> [[TMP30]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP31:%.*]] = select <4 x i1> [[TMP29]], <4 x i64> [[TMP30]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP36]] = select <4 x i1> [[TMP4]], <4 x i64> [[TMP31]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add nuw nsw <4 x i64> [[VEC_IND]], splat (i64 8)
- ; TF-NEXT: [[TMP32:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
- ; TF-NEXT: br i1 [[TMP32]], label %[[MIDDLE_BLOCK:.*]], label %[[VECTOR_BODY]], !llvm.loop [[LOOP6:![0-9]+]]
- ; TF: [[MIDDLE_BLOCK]]:
--; TF-NEXT: [[TMP33:%.*]] = call i64 @llvm.vector.reduce.umax.v4i64(<4 x i64> [[TMP31]])
-+; TF-NEXT: [[TMP33:%.*]] = call i64 @llvm.vector.reduce.umax.v4i64(<4 x i64> [[TMP36]])
- ; TF-NEXT: [[TMP34:%.*]] = icmp ne i64 [[TMP33]], 0
- ; TF-NEXT: [[TMP35:%.*]] = select i1 [[TMP34]], i64 [[TMP33]], i64 -1
- ; TF-NEXT: br label %[[DONE:.*]]
-@@ -1213,7 +1217,8 @@
- ; TF-NEXT: [[TMP7:%.*]] = getelementptr i64, ptr [[TMP6]], i64 -3
- ; TF-NEXT: [[WIDE_MASKED_LOAD8:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP7]], <4 x i1> [[REVERSE]], <4 x i64> poison)
- ; TF-NEXT: [[REVERSE9:%.*]] = shufflevector <4 x i64> [[WIDE_MASKED_LOAD8]], <4 x i64> poison, <4 x i32> <i32 3, i32 2, i32 1, i32 0>
--; TF-NEXT: [[TMP8:%.*]] = icmp sle <4 x i64> [[REVERSE6]], [[REVERSE9]]
-+; TF-NEXT: [[TMP16:%.*]] = icmp sle <4 x i64> [[REVERSE6]], [[REVERSE9]]
-+; TF-NEXT: [[TMP8:%.*]] = select <4 x i1> [[TMP2]], <4 x i1> [[TMP16]], <4 x i1> zeroinitializer
- ; TF-NEXT: [[TMP9]] = select <4 x i1> [[TMP8]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[TMP10]] = or <4 x i1> [[VEC_PHI3]], [[TMP8]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
-@@ -1321,7 +1326,8 @@
- ; TF-NEXT: [[TMP1:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[INDEX]]
- ; TF-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP1]], <4 x i1> [[TMP0]], <4 x i64> poison)
- ; TF-NEXT: [[TMP2:%.*]] = icmp eq <4 x i64> [[WIDE_MASKED_LOAD]], splat (i64 42)
--; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP2]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP9:%.*]] = select <4 x i1> [[TMP0]], <4 x i1> [[TMP2]], <4 x i1> zeroinitializer
-+; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP9]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], splat (i64 4)
- ; TF-NEXT: [[TMP4:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-@@ -1424,7 +1430,8 @@
- ; TF-NEXT: [[TMP1:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[INDEX]]
- ; TF-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP1]], <4 x i1> [[TMP0]], <4 x i64> poison)
- ; TF-NEXT: [[TMP2:%.*]] = icmp eq <4 x i64> [[WIDE_MASKED_LOAD]], splat (i64 42)
--; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP2]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; TF-NEXT: [[TMP9:%.*]] = select <4 x i1> [[TMP0]], <4 x i1> [[TMP2]], <4 x i1> zeroinitializer
-+; TF-NEXT: [[TMP3]] = select <4 x i1> [[TMP9]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
- ; TF-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
- ; TF-NEXT: [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], splat (i64 4)
- ; TF-NEXT: [[TMP4:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-diff -ruN --strip-trailing-cr a/llvm/test/Transforms/LoopVectorize/iv-select-cmp-decreasing.ll b/llvm/test/Transforms/LoopVectorize/iv-select-cmp-decreasing.ll
---- a/llvm/test/Transforms/LoopVectorize/iv-select-cmp-decreasing.ll
-+++ b/llvm/test/Transforms/LoopVectorize/iv-select-cmp-decreasing.ll
-@@ -330,11 +330,15 @@
- ; IC4VF4-NEXT: [[TMP77:%.*]] = select <4 x i1> [[TMP101]], <4 x i16> <i16 7, i16 6, i16 5, i16 4>, <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: [[TMP70:%.*]] = select <4 x i1> [[TMP102]], <4 x i16> <i16 3, i16 2, i16 1, i16 0>, <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: [[TMP71:%.*]] = select <4 x i1> [[TMP103]], <4 x i16> <i16 -1, i16 -2, i16 -3, i16 -4>, <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP82:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP76]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP83:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP77]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP88:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP70]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP89:%.*]] = select <4 x i1> zeroinitializer, <4 x i16> [[TMP71]], <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: br label %[[MIDDLE_BLOCK:.*]]
- ; IC4VF4: [[MIDDLE_BLOCK]]:
--; IC4VF4-NEXT: [[RDX_MINMAX:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[TMP76]], <4 x i16> [[TMP77]])
--; IC4VF4-NEXT: [[RDX_MINMAX45:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX]], <4 x i16> [[TMP70]])
--; IC4VF4-NEXT: [[RDX_MINMAX46:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX45]], <4 x i16> [[TMP71]])
-+; IC4VF4-NEXT: [[RDX_MINMAX:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[TMP82]], <4 x i16> [[TMP83]])
-+; IC4VF4-NEXT: [[RDX_MINMAX31:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX]], <4 x i16> [[TMP88]])
-+; IC4VF4-NEXT: [[RDX_MINMAX46:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX31]], <4 x i16> [[TMP89]])
- ; IC4VF4-NEXT: [[TMP116:%.*]] = call i16 @llvm.vector.reduce.smin.v4i16(<4 x i16> [[RDX_MINMAX46]])
- ; IC4VF4-NEXT: [[RDX_SELECT_CMP:%.*]] = icmp ne i16 [[TMP116]], 32767
- ; IC4VF4-NEXT: [[RDX_SELECT:%.*]] = select i1 [[RDX_SELECT_CMP]], i16 [[TMP116]], i16 0
-@@ -592,11 +596,15 @@
- ; IC4VF4-NEXT: [[TMP77:%.*]] = select <4 x i1> [[TMP101]], <4 x i16> <i16 7, i16 6, i16 5, i16 4>, <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: [[TMP70:%.*]] = select <4 x i1> [[TMP102]], <4 x i16> <i16 3, i16 2, i16 1, i16 0>, <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: [[TMP71:%.*]] = select <4 x i1> [[TMP103]], <4 x i16> <i16 -1, i16 -2, i16 -3, i16 -4>, <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP82:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP76]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP83:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP77]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP88:%.*]] = select <4 x i1> splat (i1 true), <4 x i16> [[TMP70]], <4 x i16> splat (i16 32767)
-+; IC4VF4-NEXT: [[TMP89:%.*]] = select <4 x i1> zeroinitializer, <4 x i16> [[TMP71]], <4 x i16> splat (i16 32767)
- ; IC4VF4-NEXT: br label %[[MIDDLE_BLOCK:.*]]
- ; IC4VF4: [[MIDDLE_BLOCK]]:
--; IC4VF4-NEXT: [[RDX_MINMAX:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[TMP76]], <4 x i16> [[TMP77]])
--; IC4VF4-NEXT: [[RDX_MINMAX45:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX]], <4 x i16> [[TMP70]])
--; IC4VF4-NEXT: [[RDX_MINMAX46:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX45]], <4 x i16> [[TMP71]])
-+; IC4VF4-NEXT: [[RDX_MINMAX:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[TMP82]], <4 x i16> [[TMP83]])
-+; IC4VF4-NEXT: [[RDX_MINMAX31:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX]], <4 x i16> [[TMP88]])
-+; IC4VF4-NEXT: [[RDX_MINMAX46:%.*]] = call <4 x i16> @llvm.smin.v4i16(<4 x i16> [[RDX_MINMAX31]], <4 x i16> [[TMP89]])
- ; IC4VF4-NEXT: [[TMP116:%.*]] = call i16 @llvm.vector.reduce.smin.v4i16(<4 x i16> [[RDX_MINMAX46]])
- ; IC4VF4-NEXT: [[RDX_SELECT_CMP:%.*]] = icmp ne i16 [[TMP116]], 32767
- ; IC4VF4-NEXT: [[RDX_SELECT:%.*]] = select i1 [[RDX_SELECT_CMP]], i16 [[TMP116]], i16 0
-diff -ruN --strip-trailing-cr a/llvm/test/Transforms/LoopVectorize/iv-select-cmp-fold-tail.ll b/llvm/test/Transforms/LoopVectorize/iv-select-cmp-fold-tail.ll
---- a/llvm/test/Transforms/LoopVectorize/iv-select-cmp-fold-tail.ll
-+++ b/llvm/test/Transforms/LoopVectorize/iv-select-cmp-fold-tail.ll
-@@ -59,3 +59,80 @@
- exit:
- ret i32 %rdx.next
- }
-+
-+define i64 @select_decreasing_induction_icmp_non_const_start(ptr %a, ptr %b, i64 %rdx.start, i64 %n) {
-+; CHECK-LABEL: define i64 @select_decreasing_induction_icmp_non_const_start(
-+; CHECK-SAME: ptr [[A:%.*]], ptr [[B:%.*]], i64 [[RDX_START:%.*]], i64 [[N:%.*]]) {
-+; CHECK-NEXT: [[ENTRY:.*:]]
-+; CHECK-NEXT: [[TMP0:%.*]] = add i64 [[N]], 1
-+; CHECK-NEXT: [[UMIN:%.*]] = call i64 @llvm.umin.i64(i64 [[N]], i64 1)
-+; CHECK-NEXT: [[TMP1:%.*]] = sub i64 [[TMP0]], [[UMIN]]
-+; CHECK-NEXT: br label %[[VECTOR_PH:.*]]
-+; CHECK: [[VECTOR_PH]]:
-+; CHECK-NEXT: [[N_RND_UP:%.*]] = add i64 [[TMP1]], 3
-+; CHECK-NEXT: [[N_MOD_VF:%.*]] = urem i64 [[N_RND_UP]], 4
-+; CHECK-NEXT: [[N_VEC:%.*]] = sub i64 [[N_RND_UP]], [[N_MOD_VF]]
-+; CHECK-NEXT: [[TRIP_COUNT_MINUS_1:%.*]] = sub i64 [[TMP1]], 1
-+; CHECK-NEXT: [[BROADCAST_SPLATINSERT:%.*]] = insertelement <4 x i64> poison, i64 [[TRIP_COUNT_MINUS_1]], i64 0
-+; CHECK-NEXT: [[BROADCAST_SPLAT:%.*]] = shufflevector <4 x i64> [[BROADCAST_SPLATINSERT]], <4 x i64> poison, <4 x i32> zeroinitializer
-+; CHECK-NEXT: [[BROADCAST_SPLATINSERT1:%.*]] = insertelement <4 x i64> poison, i64 [[N]], i64 0
-+; CHECK-NEXT: [[BROADCAST_SPLAT2:%.*]] = shufflevector <4 x i64> [[BROADCAST_SPLATINSERT1]], <4 x i64> poison, <4 x i32> zeroinitializer
-+; CHECK-NEXT: [[INDUCTION:%.*]] = add nsw <4 x i64> [[BROADCAST_SPLAT2]], <i64 0, i64 -1, i64 -2, i64 -3>
-+; CHECK-NEXT: br label %[[VECTOR_BODY:.*]]
-+; CHECK: [[VECTOR_BODY]]:
-+; CHECK-NEXT: [[INDEX:%.*]] = phi i64 [ 0, %[[VECTOR_PH]] ], [ [[INDEX_NEXT:%.*]], %[[VECTOR_BODY]] ]
-+; CHECK-NEXT: [[VEC_IND:%.*]] = phi <4 x i64> [ [[INDUCTION]], %[[VECTOR_PH]] ], [ [[VEC_IND_NEXT:%.*]], %[[VECTOR_BODY]] ]
-+; CHECK-NEXT: [[VEC_PHI:%.*]] = phi <4 x i64> [ splat (i64 9223372036854775807), %[[VECTOR_PH]] ], [ [[TMP52:%.*]], %[[VECTOR_BODY]] ]
-+; CHECK-NEXT: [[VEC_PHI3:%.*]] = phi <4 x i1> [ zeroinitializer, %[[VECTOR_PH]] ], [ [[TMP50:%.*]], %[[VECTOR_BODY]] ]
-+; CHECK-NEXT: [[OFFSET_IDX:%.*]] = sub i64 [[N]], [[INDEX]]
-+; CHECK-NEXT: [[BROADCAST_SPLATINSERT4:%.*]] = insertelement <4 x i64> poison, i64 [[INDEX]], i64 0
-+; CHECK-NEXT: [[BROADCAST_SPLAT5:%.*]] = shufflevector <4 x i64> [[BROADCAST_SPLATINSERT4]], <4 x i64> poison, <4 x i32> zeroinitializer
-+; CHECK-NEXT: [[VEC_IV:%.*]] = add <4 x i64> [[BROADCAST_SPLAT5]], <i64 0, i64 1, i64 2, i64 3>
-+; CHECK-NEXT: [[TMP2:%.*]] = icmp ule <4 x i64> [[VEC_IV]], [[BROADCAST_SPLAT]]
-+; CHECK-NEXT: [[TMP38:%.*]] = add nsw i64 [[OFFSET_IDX]], -1
-+; CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds i64, ptr [[A]], i64 [[TMP38]]
-+; CHECK-NEXT: [[TMP5:%.*]] = getelementptr i64, ptr [[TMP39]], i64 -3
-+; CHECK-NEXT: [[REVERSE:%.*]] = shufflevector <4 x i1> [[TMP2]], <4 x i1> poison, <4 x i32> <i32 3, i32 2, i32 1, i32 0>
-+; CHECK-NEXT: [[WIDE_MASKED_LOAD:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP5]], <4 x i1> [[REVERSE]], <4 x i64> poison)
-+; CHECK-NEXT: [[TMP46:%.*]] = shufflevector <4 x i64> [[WIDE_MASKED_LOAD]], <4 x i64> poison, <4 x i32> <i32 3, i32 2, i32 1, i32 0>
-+; CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds i64, ptr [[B]], i64 [[TMP38]]
-+; CHECK-NEXT: [[TMP7:%.*]] = getelementptr i64, ptr [[TMP6]], i64 -3
-+; CHECK-NEXT: [[WIDE_MASKED_LOAD8:%.*]] = call <4 x i64> @llvm.masked.load.v4i64.p0(ptr align 8 [[TMP7]], <4 x i1> [[REVERSE]], <4 x i64> poison)
-+; CHECK-NEXT: [[TMP47:%.*]] = shufflevector <4 x i64> [[WIDE_MASKED_LOAD8]], <4 x i64> poison, <4 x i32> <i32 3, i32 2, i32 1, i32 0>
-+; CHECK-NEXT: [[TMP48:%.*]] = icmp sgt <4 x i64> [[TMP46]], [[TMP47]]
-+; CHECK-NEXT: [[TMP49:%.*]] = select <4 x i1> [[TMP2]], <4 x i1> [[TMP48]], <4 x i1> zeroinitializer
-+; CHECK-NEXT: [[TMP52]] = select <4 x i1> [[TMP49]], <4 x i64> [[VEC_IND]], <4 x i64> [[VEC_PHI]]
-+; CHECK-NEXT: [[TMP50]] = or <4 x i1> [[VEC_PHI3]], [[TMP49]]
-+; CHECK-NEXT: [[INDEX_NEXT]] = add i64 [[INDEX]], 4
-+; CHECK-NEXT: [[VEC_IND_NEXT]] = add nsw <4 x i64> [[VEC_IND]], splat (i64 -4)
-+; CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
-+; CHECK-NEXT: br i1 [[TMP51]], label %[[MIDDLE_BLOCK:.*]], label %[[VECTOR_BODY]], !llvm.loop [[LOOP3:![0-9]+]]
-+; CHECK: [[MIDDLE_BLOCK]]:
-+; CHECK-NEXT: [[TMP53:%.*]] = call i64 @llvm.vector.reduce.smin.v4i64(<4 x i64> [[TMP52]])
-+; CHECK-NEXT: [[TMP54:%.*]] = add nsw i64 [[TMP53]], -1
-+; CHECK-NEXT: [[TMP56:%.*]] = call i1 @llvm.vector.reduce.or.v4i1(<4 x i1> [[TMP50]])
-+; CHECK-NEXT: [[TMP57:%.*]] = freeze i1 [[TMP56]]
-+; CHECK-NEXT: [[RDX_SELECT:%.*]] = select i1 [[TMP57]], i64 [[TMP54]], i64 [[RDX_START]]
-+; CHECK-NEXT: br label %[[EXIT:.*]]
-+; CHECK: [[EXIT]]:
-+; CHECK-NEXT: ret i64 [[RDX_SELECT]]
-+;
-+entry:
-+ br label %loop
-+
-+loop:
-+ %iv = phi i64 [ %iv.next, %loop ], [ %n, %entry ]
-+ %rdx = phi i64 [ %cond, %loop ], [ %rdx.start, %entry ]
-+ %iv.next = add nsw i64 %iv, -1
-+ %gep.a.iv = getelementptr inbounds i64, ptr %a, i64 %iv.next
-+ %ld.a = load i64, ptr %gep.a.iv, align 8
-+ %gep.b.iv = getelementptr inbounds i64, ptr %b, i64 %iv.next
-+ %ld.b = load i64, ptr %gep.b.iv, align 8
-+ %cmp.a.b = icmp sgt i64 %ld.a, %ld.b
-+ %cond = select i1 %cmp.a.b, i64 %iv.next, i64 %rdx
-+ %exit.cond = icmp ugt i64 %iv, 1
-+ br i1 %exit.cond, label %loop, label %exit
-+
-+exit:
-+ ret i64 %cond
-+}
diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl
index f220caa..8ee1b2c 100644
--- a/third_party/llvm/workspace.bzl
+++ b/third_party/llvm/workspace.bzl
@@ -4,8 +4,8 @@
def repo(name):
"""Imports LLVM."""
- LLVM_COMMIT = "2cf353b5e8560723409f3f9164bddec76f499963"
- LLVM_SHA256 = "3a9b33b701b9ebd4110ef795d23be54fb5d762ca4e190480c4a06d71b7aa8064"
+ LLVM_COMMIT = "7189c4bb83d90c383741a9480fceca9c7bf74b27"
+ LLVM_SHA256 = "dd652f7b4e46001614ea3c1e10cbb117a73199f2d8f7e95c66f233bd0111fbd4"
tf_http_archive(
name = name,
