Integrate LLVM at llvm/llvm-project@20f79f8caa3a
Updates LLVM usage to match
[20f79f8caa3a](https://github.com/llvm/llvm-project/commit/20f79f8caa3a)
PiperOrigin-RevId: 411170625
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt
index 690f318..6c965b3 100644
--- a/SUBMODULE_VERSIONS.txt
+++ b/SUBMODULE_VERSIONS.txt
@@ -4,7 +4,7 @@
aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest
88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing
acd6f6f014c25e46363e718381e0b35205df2d83 third_party/libyaml
-d61840c168a34339aa8e602adf5aba98924a6f61 third_party/llvm-project
+20f79f8caa3a333a34021f0028e828f97d79c2a1 third_party/llvm-project
d0ee2ca2145e07e237dacb7ec7ded84f4a7842fc third_party/mlir-hlo
3f701faace7addc75d16dea8a6cd769fa5b3f260 third_party/musl
4c7697dbe973ed01ae6fbec37d186ebd05982e1f third_party/pybind11
diff --git a/integrations/tensorflow/iree_tf_compiler/MHLO/FlattenTuplesInCFG.cpp b/integrations/tensorflow/iree_tf_compiler/MHLO/FlattenTuplesInCFG.cpp
index 840f198..0fcf381 100644
--- a/integrations/tensorflow/iree_tf_compiler/MHLO/FlattenTuplesInCFG.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/MHLO/FlattenTuplesInCFG.cpp
@@ -53,11 +53,11 @@
for (const auto &oldAttr : oldOp->getAttrs()) {
// Don't copy segment attributes as these correspond to the number operands,
// which may be different.
- if (oldAttr.first == "operand_segment_sizes" ||
- oldAttr.first == "result_segment_sizes")
+ if (oldAttr.getName() == "operand_segment_sizes" ||
+ oldAttr.getName() == "result_segment_sizes")
continue;
- newOp->setAttr(oldAttr.first, oldAttr.second);
+ newOp->setAttr(oldAttr.getName(), oldAttr.getValue());
}
}
@@ -237,8 +237,8 @@
BlockAndValueMapping mapping;
for (auto attr : oldFunction->getAttrs()) {
- if (attr.first != oldFunction.getTypeAttrName()) {
- newFunction->setAttr(attr.first, attr.second);
+ if (attr.getName() != oldFunction.getTypeAttrName()) {
+ newFunction->setAttr(attr.getName(), attr.getValue());
}
}
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/StripMetadata.cpp b/integrations/tensorflow/iree_tf_compiler/TF/StripMetadata.cpp
index b44db44..f1fc99c 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/StripMetadata.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/TF/StripMetadata.cpp
@@ -17,11 +17,11 @@
namespace TF {
static bool isTFAttr(NamedAttribute &namedAttr) {
- auto name = namedAttr.first.strref();
+ auto name = namedAttr.getName().strref();
if (name.startswith("tf.") || name.startswith("tf_")) {
return true;
}
- StringRef attrNamespace = namedAttr.second.getDialect().getNamespace();
+ StringRef attrNamespace = namedAttr.getValue().getDialect().getNamespace();
return attrNamespace == mlir::TF::TensorFlowDialect::getDialectNamespace() ||
attrNamespace == mlir::tf_executor::TensorFlowExecutorDialect::
getDialectNamespace() ||
@@ -48,7 +48,7 @@
moduleOp->getAttrs(),
[](NamedAttribute namedAttr) { return isTFAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- moduleOp->removeAttr(namedAttr.first);
+ moduleOp->removeAttr(namedAttr.getName());
}
}
};
@@ -70,7 +70,7 @@
funcOp->getAttrs(),
[](NamedAttribute namedAttr) { return isTFAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp->removeAttr(namedAttr.first);
+ funcOp->removeAttr(namedAttr.getName());
}
for (int i = 0; i < funcOp.getNumArguments(); ++i) {
@@ -78,7 +78,7 @@
funcOp.getArgAttrs(i),
[](NamedAttribute namedAttr) { return isTFAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp.removeArgAttr(i, namedAttr.first);
+ funcOp.removeArgAttr(i, namedAttr.getName());
}
}
@@ -87,7 +87,7 @@
funcOp.getResultAttrs(i),
[](NamedAttribute namedAttr) { return isTFAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp.removeResultAttr(i, namedAttr.first);
+ funcOp.removeResultAttr(i, namedAttr.getName());
}
}
}
diff --git a/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp b/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp
index 0ba99e6..af1dbe6 100644
--- a/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/TFL/StripMetadata.cpp
@@ -16,12 +16,12 @@
static bool isTFLAttr(NamedAttribute &namedAttr) {
// NOTE: tflite mixes tf and tfl, for some reason.
- auto name = namedAttr.first.strref();
+ auto name = namedAttr.getName().strref();
if (name.startswith("tf.") || name.startswith("tf_") ||
name.startswith("tfl.") || name.startswith("tfl_")) {
return true;
}
- StringRef attrNamespace = namedAttr.second.getDialect().getNamespace();
+ StringRef attrNamespace = namedAttr.getValue().getDialect().getNamespace();
return attrNamespace == "tf" || attrNamespace == "tfl";
}
@@ -35,7 +35,7 @@
moduleOp->getAttrs(),
[](NamedAttribute namedAttr) { return isTFLAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- moduleOp->removeAttr(namedAttr.first);
+ moduleOp->removeAttr(namedAttr.getName());
}
}
};
@@ -50,7 +50,7 @@
funcOp->getAttrs(),
[](NamedAttribute namedAttr) { return isTFLAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp->removeAttr(namedAttr.first);
+ funcOp->removeAttr(namedAttr.getName());
}
for (int i = 0; i < funcOp.getNumArguments(); ++i) {
@@ -58,7 +58,7 @@
funcOp.getArgAttrs(i),
[](NamedAttribute namedAttr) { return isTFLAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp.removeArgAttr(i, namedAttr.first);
+ funcOp.removeArgAttr(i, namedAttr.getName());
}
}
@@ -67,7 +67,7 @@
funcOp.getResultAttrs(i),
[](NamedAttribute namedAttr) { return isTFLAttr(namedAttr); }));
for (auto namedAttr : stripAttrs) {
- funcOp.removeResultAttr(i, namedAttr.first);
+ funcOp.removeResultAttr(i, namedAttr.getName());
}
}
}
diff --git a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp
index 03ff4d7..4e102d2 100644
--- a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp
+++ b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp
@@ -169,8 +169,8 @@
SmallVector<NamedAttribute, 4> attrs;
auto abiAttr = entryFuncOp->getAttr("iree.abi");
if (abiAttr) {
- attrs.push_back(std::make_pair(
- Identifier::get("iree.abi", entryFuncOp.getContext()), abiAttr));
+ attrs.emplace_back(StringAttr::get(entryFuncOp.getContext(), "iree.abi"),
+ abiAttr);
}
if (!attrs.empty()) {
auto reflectionAttr = DictionaryAttr::get(&getContext(), attrs);
diff --git a/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
index 2025fbd..c96e44d 100644
--- a/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
+++ b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
@@ -97,7 +97,7 @@
ConversionPatternRewriter &rewriter,
SmallVectorImpl<NamedAttribute> &newAttrs) {
for (auto attr : attrs) {
- if (auto fpAttr = attr.second.dyn_cast<DenseFPElementsAttr>()) {
+ if (auto fpAttr = attr.getValue().dyn_cast<DenseFPElementsAttr>()) {
std::vector<llvm::APFloat> args;
if (!fpAttr.getType().getElementType().isF32()) continue;
for (llvm::APFloat f : fpAttr.getValues<APFloat>()) {
@@ -107,16 +107,15 @@
}
auto tensorType = RankedTensorType::get(fpAttr.getType().getShape(),
rewriter.getF16Type());
- newAttrs.push_back(std::make_pair(
- attr.first, DenseElementsAttr::get(tensorType, args)));
- } else if (auto typeAttr = attr.second.dyn_cast<TypeAttr>()) {
+ newAttrs.emplace_back(attr.getName(),
+ DenseElementsAttr::get(tensorType, args));
+ } else if (auto typeAttr = attr.getValue().dyn_cast<TypeAttr>()) {
if (isIllegalType(typeAttr.getValue())) {
if (auto tensorType =
typeAttr.getValue().dyn_cast<RankedTensorType>()) {
Type newType = RankedTensorType::get(tensorType.getShape(),
rewriter.getF16Type());
- newAttrs.push_back(
- std::make_pair(attr.first, TypeAttr::get(newType)));
+ newAttrs.emplace_back(attr.getName(), TypeAttr::get(newType));
}
}
} else {
diff --git a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index d89fd17..62eeb4e 100644
--- a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -434,8 +434,8 @@
// MLIR implicitly.
SmallVector<NamedAttribute, 4> funcAttrs;
for (auto attr : stdFuncOp->getAttrs()) {
- if (attr.first == SymbolTable::getSymbolAttrName() ||
- attr.first == mlir::function_like_impl::getTypeAttrName()) {
+ if (attr.getName() == SymbolTable::getSymbolAttrName() ||
+ attr.getName() == mlir::function_like_impl::getTypeAttrName()) {
continue;
}
funcAttrs.push_back(attr);
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
index 116680f..e7c8e95 100644
--- a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
@@ -192,8 +192,8 @@
// Construct newFunc with all attributes except return type & symbol name.
SmallVector<NamedAttribute, 4> funcAttrs;
for (auto attr : funcOp->getAttrs()) {
- if (attr.first == SymbolTable::getSymbolAttrName() ||
- attr.first == mlir::function_like_impl::getTypeAttrName()) {
+ if (attr.getName() == SymbolTable::getSymbolAttrName() ||
+ attr.getName() == mlir::function_like_impl::getTypeAttrName()) {
continue;
}
funcAttrs.push_back(attr);
diff --git a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
index dff4182..339e4d5 100644
--- a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
@@ -308,6 +308,12 @@
// No tiled loops means we cannot tile (and distribute) at all. Use just one
// single thread to run everything.
auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize;
+ if (auto linalgOp = dyn_cast<linalg::GenericOp>(op)) {
+ auto untiledResultShape = getUntiledResultShape(linalgOp, 0);
+ if (untiledResultShape.empty()) {
+ pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistribute;
+ }
+ }
std::array<int64_t, 3> workgroupSize = {1, 1, 1};
return setOpConfigAndEntryPointFnTranslation(funcOp, op, {}, {}, pipeline,
workgroupSize);
@@ -388,6 +394,7 @@
};
// Whether we can try to use the vectorization pipeline.
+ auto untiledResultShape = getUntiledResultShape(linalgOp, 0);
bool vectorizable =
!linalgOp.hasIndexSemantics() &&
// Skip vectorization for non-minor identity inputs as it generates
@@ -400,7 +407,8 @@
// TODO: Lowering of integers other than i32 may require emulation.
// This is currently not supported for vector operation.
llvm::all_of(linalgOp->getOperands(), has32BitElementType) &&
- llvm::none_of(getUntiledResultShape(linalgOp, 0), ShapedType::isDynamic);
+ !untiledResultShape.empty() &&
+ llvm::none_of(untiledResultShape, ShapedType::isDynamic);
LLVM_DEBUG({
llvm::dbgs() << "Linalg op " << linalgOp << "\n partitioned loops: [";
diff --git a/iree/compiler/Codegen/Transforms/Transforms.cpp b/iree/compiler/Codegen/Transforms/Transforms.cpp
index b4b6056..7d65ba4 100644
--- a/iree/compiler/Codegen/Transforms/Transforms.cpp
+++ b/iree/compiler/Codegen/Transforms/Transforms.cpp
@@ -48,12 +48,12 @@
entryPointOp.workgroup_local_memoryAttr(), 1);
// Copy over all attributes
for (auto attr : entryPointOp->getAttrs()) {
- if (attr.first != entryPointOp.sym_nameAttrName() &&
- attr.first != entryPointOp.ordinalAttrName() &&
- attr.first != entryPointOp.interfaceAttrName() &&
- attr.first != entryPointOp.workgroup_sizeAttrName() &&
- attr.first != entryPointOp.workgroup_local_memoryAttrName()) {
- clonedOp->setAttr(attr.first, attr.second);
+ if (attr.getName() != entryPointOp.sym_nameAttrName() &&
+ attr.getName() != entryPointOp.ordinalAttrName() &&
+ attr.getName() != entryPointOp.interfaceAttrName() &&
+ attr.getName() != entryPointOp.workgroup_sizeAttrName() &&
+ attr.getName() != entryPointOp.workgroup_local_memoryAttrName()) {
+ clonedOp->setAttr(attr.getName(), attr.getValue());
}
}
Region *region = clonedOp.getBody();
diff --git a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
index 8989fbc..359919c 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
@@ -132,8 +132,8 @@
if (!compare_ranges(
lhs.getAttrs(), rhs.getAttrs(),
[&](const NamedAttribute &lhs, const NamedAttribute &rhs) {
- if (lhs.first == "function_ref" ||
- lhs.first == SymbolTable::getSymbolAttrName()) {
+ if (lhs.getName() == "function_ref" ||
+ lhs.getName() == SymbolTable::getSymbolAttrName()) {
return true;
}
return lhs == rhs;
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
index 55f6cdf..6c0e24e 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
@@ -165,7 +165,7 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getIdentifier("executable_targets"),
+ configItems.emplace_back(b.getStringAttr("executable_targets"),
getExecutableTargets(context));
auto configAttr = b.getDictionaryAttr(configItems);
@@ -323,8 +323,7 @@
SmallVector<NamedAttribute> configItems;
// Add some configurations to the `hal.executable.target` attribute.
auto addConfig = [&](StringRef name, Attribute value) {
- configItems.emplace_back(
- std::make_pair(Identifier::get(name, context), value));
+ configItems.emplace_back(StringAttr::get(context, name), value);
};
// Set target arch
addConfig("target_arch", StringAttr::get(context, clTargetChip));
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
index 84999f9..6827555 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
@@ -550,8 +550,7 @@
// Add some configurations to the `hal.executable.target` attribute.
SmallVector<NamedAttribute> config;
auto addConfig = [&](StringRef name, Attribute value) {
- config.emplace_back(
- std::make_pair(Identifier::get(name, context), value));
+ config.emplace_back(StringAttr::get(context, name), value);
};
// Set target triple.
diff --git a/iree/compiler/Dialect/VM/Analysis/ValueLiveness.cpp b/iree/compiler/Dialect/VM/Analysis/ValueLiveness.cpp
index 9ffe386..06d60ad 100644
--- a/iree/compiler/Dialect/VM/Analysis/ValueLiveness.cpp
+++ b/iree/compiler/Dialect/VM/Analysis/ValueLiveness.cpp
@@ -120,7 +120,7 @@
// Markup all ops with their attributes.
for (auto &opAttrs : livenessAttrs) {
for (auto nameAttr : opAttrs.getSecond().getAttrs()) {
- opAttrs.getFirst()->setAttr(nameAttr.first, nameAttr.second);
+ opAttrs.getFirst()->setAttr(nameAttr.getName(), nameAttr.getValue());
}
}
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.cpp b/iree/compiler/Dialect/VM/IR/VMOps.cpp
index 35213d0..e4008de 100644
--- a/iree/compiler/Dialect/VM/IR/VMOps.cpp
+++ b/iree/compiler/Dialect/VM/IR/VMOps.cpp
@@ -102,8 +102,8 @@
SmallVector<NamedAttribute> attrs(existingAttr.begin(), existingAttr.end());
bool didFind = false;
for (size_t i = 0; i < attrs.size(); ++i) {
- if (attrs[i].first == name) {
- attrs[i].second = value;
+ if (attrs[i].getName() == name) {
+ attrs[i].setValue(value);
didFind = true;
break;
}
diff --git a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp
index 199070d..a4ee999 100644
--- a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp
+++ b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp
@@ -500,8 +500,8 @@
funcOp->getAttrOfType<DictionaryAttr>("iree.reflection")) {
SmallVector<iree_vm_ReflectionAttrDef_ref_t, 4> reflectionAttrRefs;
for (auto reflectionAttr : reflectionAttrs) {
- auto key = reflectionAttr.first.strref();
- auto value = reflectionAttr.second.dyn_cast<StringAttr>();
+ auto key = reflectionAttr.getName().strref();
+ auto value = reflectionAttr.getValue().dyn_cast<StringAttr>();
if (!value || key.empty()) continue;
// NOTE: if we actually want to keep these we should dedupe them (as the
// keys and likely several of the values are shared across all functions).
diff --git a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
index 5cf8e42..5d99e15 100644
--- a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
+++ b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
@@ -99,12 +99,12 @@
llvm::isa<IREE::Util::GlobalOp>(oldOp)) {
for (auto attr : oldOp->getAttrs()) {
auto newAttr =
- convertAttribute(oldOp->getLoc(), attr.second, typeConverter);
+ convertAttribute(oldOp->getLoc(), attr.getValue(), typeConverter);
if (!newAttr) {
return oldOp->emitOpError()
- << "failed to convert attribute " << attr.first;
+ << "failed to convert attribute " << attr.getName();
}
- state.addAttribute(attr.first, newAttr);
+ state.addAttribute(attr.getName(), newAttr);
}
} else {
state.attributes = llvm::to_vector<4>(oldOp->getAttrs());
diff --git a/third_party/llvm-project b/third_party/llvm-project
index d61840c..20f79f8 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit d61840c168a34339aa8e602adf5aba98924a6f61
+Subproject commit 20f79f8caa3a333a34021f0028e828f97d79c2a1
