[Mlir-commits] [mlir] [MLIR][LLVM] Remove the type consistency pass (PR #93283)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri May 24 02:29:08 PDT 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-llvm
Author: Christian Ulmann (Dinistro)
<details>
<summary>Changes</summary>
This commit removes the LLVM dialect's type consistency pass. This pass was originally introduced to make type information on memory operations consistent. The main benefactor of this consistency where Mem2Reg and SROA, which were in the meantime improved to no longer require consitent type information.
Appart from providing a no longer required functionality, the pass had some fundamental flaws that lead to issues:
* It introduced trivial GEPs (only zero indices) that could fold be folded again.
* Aggressively splitting stores lead to substantial performance regressions in some cases. Subsequent memory coalescings were not able to recover this information, due to using non-trivial bit-fiddling.
---
Patch is 53.04 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/93283.diff
6 Files Affected:
- (modified) mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h (-1)
- (modified) mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td (-18)
- (removed) mlir/include/mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h (-73)
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt (-1)
- (removed) mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp (-575)
- (removed) mlir/test/Dialect/LLVMIR/type-consistency.mlir (-533)
``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
index 13e10b29c0743..a7bf8796c0277 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
@@ -13,7 +13,6 @@
#include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
#include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
-#include "mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h"
#include "mlir/Pass/Pass.h"
namespace mlir {
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
index 0242cfd9abb7d..11d1b94110714 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
@@ -43,24 +43,6 @@ def LLVMRequestCWrappers
let constructor = "::mlir::LLVM::createRequestCWrappersPass()";
}
-def LLVMTypeConsistency
- : Pass<"llvm-type-consistency", "::mlir::LLVM::LLVMFuncOp"> {
- let summary = "Rewrites to improve type consistency";
- let description = [{
- Set of rewrites to improve the coherency of types within an LLVM dialect
- program. This will adjust operations operating on pointers so they interpret
- their associated pointee type as consistently as possible.
- }];
- let constructor = "::mlir::LLVM::createTypeConsistencyPass()";
-
- let options = [
- Option<"maxVectorSplitSize", "max-vector-split-size", "unsigned",
- /*default=*/"512",
- "Maximum size in bits of a vector value in a load or store operation"
- " operating on multiple elements that should still be split">,
- ];
-}
-
def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> {
let summary = "Optimize NVVM IR";
let constructor = "::mlir::NVVM::createOptimizeForTargetPass()";
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h
deleted file mode 100644
index a4bb380b99b86..0000000000000
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h
+++ /dev/null
@@ -1,73 +0,0 @@
-//===- TypeConsistency.h - Rewrites to improve type consistency -----------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// Set of rewrites to improve the coherency of types within an LLVM dialect
-// program. This will adjust operations around a given pointer so they interpret
-// its pointee type as consistently as possible.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_DIALECT_LLVMIR_TRANSFORMS_TYPECONSISTENCY_H
-#define MLIR_DIALECT_LLVMIR_TRANSFORMS_TYPECONSISTENCY_H
-
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/IR/PatternMatch.h"
-#include "mlir/Pass/Pass.h"
-
-namespace mlir {
-namespace LLVM {
-
-#define GEN_PASS_DECL_LLVMTYPECONSISTENCY
-#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
-
-/// Creates a pass that adjusts operations operating on pointers so they
-/// interpret pointee types as consistently as possible.
-std::unique_ptr<Pass> createTypeConsistencyPass();
-
-/// Canonicalizes GEPs of which the base type and the pointer's type hint do not
-/// match. This is done by replacing the original GEP into a GEP with the type
-/// hint as a base type when an element of the hinted type aligns with the
-/// original GEP.
-class CanonicalizeAlignedGep : public OpRewritePattern<GEPOp> {
-public:
- using OpRewritePattern::OpRewritePattern;
-
- LogicalResult matchAndRewrite(GEPOp gep,
- PatternRewriter &rewriter) const override;
-};
-
-/// Splits stores which write into multiple adjacent elements of an aggregate
-/// through a pointer. Currently, integers and vector are split and stores
-/// are generated for every element being stored to in a type-consistent manner.
-/// This is done on a best-effort basis.
-class SplitStores : public OpRewritePattern<StoreOp> {
- unsigned maxVectorSplitSize;
-
-public:
- SplitStores(MLIRContext *context, unsigned maxVectorSplitSize)
- : OpRewritePattern(context), maxVectorSplitSize(maxVectorSplitSize) {}
-
- LogicalResult matchAndRewrite(StoreOp store,
- PatternRewriter &rewrite) const override;
-};
-
-/// Splits GEPs with more than two indices into multiple GEPs with exactly
-/// two indices. The created GEPs are then guaranteed to index into only
-/// one aggregate at a time.
-class SplitGEP : public OpRewritePattern<GEPOp> {
-public:
- using OpRewritePattern::OpRewritePattern;
-
- LogicalResult matchAndRewrite(GEPOp gepOp,
- PatternRewriter &rewriter) const override;
-};
-
-} // namespace LLVM
-} // namespace mlir
-
-#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_TYPECONSISTENCY_H
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
index c80494a440116..728885fcbeaf3 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
@@ -6,7 +6,6 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms
LegalizeForExport.cpp
OptimizeForNVVM.cpp
RequestCWrappers.cpp
- TypeConsistency.cpp
DEPENDS
MLIRLLVMPassIncGen
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
deleted file mode 100644
index 0a372ad0c52fc..0000000000000
--- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
+++ /dev/null
@@ -1,575 +0,0 @@
-//===- TypeConsistency.cpp - Rewrites to improve type consistency ---------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h"
-#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#include "llvm/ADT/TypeSwitch.h"
-
-namespace mlir {
-namespace LLVM {
-#define GEN_PASS_DEF_LLVMTYPECONSISTENCY
-#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
-} // namespace LLVM
-} // namespace mlir
-
-using namespace mlir;
-using namespace LLVM;
-
-//===----------------------------------------------------------------------===//
-// Utils
-//===----------------------------------------------------------------------===//
-
-/// Checks that a pointer value has a pointee type hint consistent with the
-/// expected type. Returns the type it actually hints to if it differs, or
-/// nullptr if the type is consistent or impossible to analyze.
-static Type isElementTypeInconsistent(Value addr, Type expectedType) {
- auto defOp = dyn_cast_or_null<GetResultPtrElementType>(addr.getDefiningOp());
- if (!defOp)
- return nullptr;
-
- Type elemType = defOp.getResultPtrElementType();
- if (!elemType)
- return nullptr;
-
- if (elemType == expectedType)
- return nullptr;
-
- return elemType;
-}
-
-//===----------------------------------------------------------------------===//
-// CanonicalizeAlignedGep
-//===----------------------------------------------------------------------===//
-
-/// Returns the amount of bytes the provided GEP elements will offset the
-/// pointer by. Returns nullopt if the offset could not be computed.
-static std::optional<uint64_t> gepToByteOffset(DataLayout &layout, GEPOp gep) {
-
- SmallVector<uint32_t> indices;
- // Ensures all indices are static and fetches them.
- for (auto index : gep.getIndices()) {
- IntegerAttr indexInt = llvm::dyn_cast_if_present<IntegerAttr>(index);
- if (!indexInt)
- return std::nullopt;
- int32_t gepIndex = indexInt.getInt();
- if (gepIndex < 0)
- return std::nullopt;
- indices.push_back(static_cast<uint32_t>(gepIndex));
- }
-
- uint64_t offset = indices[0] * layout.getTypeSize(gep.getElemType());
-
- Type currentType = gep.getElemType();
- for (uint32_t index : llvm::drop_begin(indices)) {
- bool shouldCancel =
- TypeSwitch<Type, bool>(currentType)
- .Case([&](LLVMArrayType arrayType) {
- if (arrayType.getNumElements() <= index)
- return true;
- offset += index * layout.getTypeSize(arrayType.getElementType());
- currentType = arrayType.getElementType();
- return false;
- })
- .Case([&](LLVMStructType structType) {
- ArrayRef<Type> body = structType.getBody();
- if (body.size() <= index)
- return true;
- for (uint32_t i = 0; i < index; i++) {
- if (!structType.isPacked())
- offset = llvm::alignTo(offset,
- layout.getTypeABIAlignment(body[i]));
- offset += layout.getTypeSize(body[i]);
- }
- currentType = body[index];
- return false;
- })
- .Default([](Type) { return true; });
-
- if (shouldCancel)
- return std::nullopt;
- }
-
- return offset;
-}
-
-/// Fills in `equivalentIndicesOut` with GEP indices that would be equivalent to
-/// offsetting a pointer by `offset` bytes, assuming the GEP has `base` as base
-/// type.
-static LogicalResult
-findIndicesForOffset(DataLayout &layout, Type base, uint64_t offset,
- SmallVectorImpl<GEPArg> &equivalentIndicesOut) {
-
- uint64_t baseSize = layout.getTypeSize(base);
- uint64_t rootIndex = offset / baseSize;
- if (rootIndex > std::numeric_limits<uint32_t>::max())
- return failure();
- equivalentIndicesOut.push_back(rootIndex);
-
- uint64_t distanceToStart = rootIndex * baseSize;
-
-#ifndef NDEBUG
- auto isWithinCurrentType = [&](Type currentType) {
- return offset < distanceToStart + layout.getTypeSize(currentType);
- };
-#endif
-
- Type currentType = base;
- while (distanceToStart < offset) {
- // While an index that does not perfectly align with offset has not been
- // reached...
-
- assert(isWithinCurrentType(currentType));
-
- bool shouldCancel =
- TypeSwitch<Type, bool>(currentType)
- .Case([&](LLVMArrayType arrayType) {
- // Find which element of the array contains the offset.
- uint64_t elemSize =
- layout.getTypeSize(arrayType.getElementType());
- uint64_t index = (offset - distanceToStart) / elemSize;
- equivalentIndicesOut.push_back(index);
- distanceToStart += index * elemSize;
-
- // Then, try to find where in the element the offset is. If the
- // offset is exactly the beginning of the element, the loop is
- // complete.
- currentType = arrayType.getElementType();
-
- // Only continue if the element in question can be indexed using
- // an i32.
- return index > std::numeric_limits<uint32_t>::max();
- })
- .Case([&](LLVMStructType structType) {
- ArrayRef<Type> body = structType.getBody();
- uint32_t index = 0;
-
- // Walk over the elements of the struct to find in which of them
- // the offset is.
- for (Type elem : body) {
- uint64_t elemSize = layout.getTypeSize(elem);
- if (!structType.isPacked()) {
- distanceToStart = llvm::alignTo(
- distanceToStart, layout.getTypeABIAlignment(elem));
- // If the offset is in padding, cancel the rewrite.
- if (offset < distanceToStart)
- return true;
- }
-
- if (offset < distanceToStart + elemSize) {
- // The offset is within this element, stop iterating the
- // struct and look within the current element.
- equivalentIndicesOut.push_back(index);
- currentType = elem;
- return false;
- }
-
- // The offset is not within this element, continue walking over
- // the struct.
- distanceToStart += elemSize;
- index++;
- }
-
- // The offset was supposed to be within this struct but is not.
- // This can happen if the offset points into final padding.
- // Anyway, nothing can be done.
- return true;
- })
- .Default([](Type) {
- // If the offset is within a type that cannot be split, no indices
- // will yield this offset. This can happen if the offset is not
- // perfectly aligned with a leaf type.
- // TODO: support vectors.
- return true;
- });
-
- if (shouldCancel)
- return failure();
- }
-
- return success();
-}
-
-/// Returns the consistent type for the GEP if the GEP is not type-consistent.
-/// Returns failure if the GEP is already consistent.
-static FailureOr<Type> getRequiredConsistentGEPType(GEPOp gep) {
- // GEP of typed pointers are not supported.
- if (!gep.getElemType())
- return failure();
-
- std::optional<Type> maybeBaseType = gep.getElemType();
- if (!maybeBaseType)
- return failure();
- Type baseType = *maybeBaseType;
-
- Type typeHint = isElementTypeInconsistent(gep.getBase(), baseType);
- if (!typeHint)
- return failure();
- return typeHint;
-}
-
-LogicalResult
-CanonicalizeAlignedGep::matchAndRewrite(GEPOp gep,
- PatternRewriter &rewriter) const {
- FailureOr<Type> typeHint = getRequiredConsistentGEPType(gep);
- if (failed(typeHint)) {
- // GEP is already canonical, nothing to do here.
- return failure();
- }
-
- DataLayout layout = DataLayout::closest(gep);
- std::optional<uint64_t> desiredOffset = gepToByteOffset(layout, gep);
- if (!desiredOffset)
- return failure();
-
- SmallVector<GEPArg> newIndices;
- if (failed(
- findIndicesForOffset(layout, *typeHint, *desiredOffset, newIndices)))
- return failure();
-
- rewriter.replaceOpWithNewOp<GEPOp>(
- gep, LLVM::LLVMPointerType::get(getContext()), *typeHint, gep.getBase(),
- newIndices, gep.getInbounds());
-
- return success();
-}
-
-namespace {
-/// Class abstracting over both array and struct types, turning each into ranges
-/// of their sub-types.
-class DestructurableTypeRange
- : public llvm::indexed_accessor_range<DestructurableTypeRange,
- DestructurableTypeInterface, Type,
- Type *, Type> {
-
- using Base = llvm::indexed_accessor_range<
- DestructurableTypeRange, DestructurableTypeInterface, Type, Type *, Type>;
-
-public:
- using Base::Base;
-
- /// Constructs a DestructurableTypeRange from either a LLVMStructType or
- /// LLVMArrayType.
- explicit DestructurableTypeRange(DestructurableTypeInterface base)
- : Base(base, 0, [&]() -> ptrdiff_t {
- return TypeSwitch<DestructurableTypeInterface, ptrdiff_t>(base)
- .Case([](LLVMStructType structType) {
- return structType.getBody().size();
- })
- .Case([](LLVMArrayType arrayType) {
- return arrayType.getNumElements();
- })
- .Default([](auto) -> ptrdiff_t {
- llvm_unreachable(
- "Only LLVMStructType or LLVMArrayType supported");
- });
- }()) {}
-
- /// Returns true if this is a range over a packed struct.
- bool isPacked() const {
- if (auto structType = dyn_cast<LLVMStructType>(getBase()))
- return structType.isPacked();
- return false;
- }
-
-private:
- static Type dereference(DestructurableTypeInterface base, ptrdiff_t index) {
- // i32 chosen because the implementations of ArrayType and StructType
- // specifically expect it to be 32 bit. They will fail otherwise.
- Type result = base.getTypeAtIndex(
- IntegerAttr::get(IntegerType::get(base.getContext(), 32), index));
- assert(result && "Should always succeed");
- return result;
- }
-
- friend Base;
-};
-} // namespace
-
-/// Returns the list of elements of `destructurableType` that are written to by
-/// a store operation writing `storeSize` bytes at `storeOffset`.
-/// `storeOffset` is required to cleanly point to an immediate element within
-/// the type. If the write operation were to write to any padding, write beyond
-/// the aggregate or partially write to a non-aggregate, failure is returned.
-static FailureOr<DestructurableTypeRange>
-getWrittenToFields(const DataLayout &dataLayout,
- DestructurableTypeInterface destructurableType,
- unsigned storeSize, unsigned storeOffset) {
- DestructurableTypeRange destructurableTypeRange(destructurableType);
-
- unsigned currentOffset = 0;
- for (; !destructurableTypeRange.empty();
- destructurableTypeRange = destructurableTypeRange.drop_front()) {
- Type type = destructurableTypeRange.front();
- if (!destructurableTypeRange.isPacked()) {
- unsigned alignment = dataLayout.getTypeABIAlignment(type);
- currentOffset = llvm::alignTo(currentOffset, alignment);
- }
-
- // currentOffset is guaranteed to be equal to offset since offset is either
- // 0 or stems from a type-consistent GEP indexing into just a single
- // aggregate.
- if (currentOffset == storeOffset)
- break;
-
- assert(currentOffset < storeOffset &&
- "storeOffset should cleanly point into an immediate field");
-
- currentOffset += dataLayout.getTypeSize(type);
- }
-
- size_t exclusiveEnd = 0;
- for (; exclusiveEnd < destructurableTypeRange.size() && storeSize > 0;
- exclusiveEnd++) {
- if (!destructurableTypeRange.isPacked()) {
- unsigned alignment =
- dataLayout.getTypeABIAlignment(destructurableTypeRange[exclusiveEnd]);
- // No padding allowed inbetween fields at this point in time.
- if (!llvm::isAligned(llvm::Align(alignment), currentOffset))
- return failure();
- }
-
- unsigned fieldSize =
- dataLayout.getTypeSize(destructurableTypeRange[exclusiveEnd]);
- if (fieldSize > storeSize) {
- // Partial writes into an aggregate are okay since subsequent pattern
- // applications can further split these up into writes into the
- // sub-elements.
- auto subAggregate = dyn_cast<DestructurableTypeInterface>(
- destructurableTypeRange[exclusiveEnd]);
- if (!subAggregate)
- return failure();
-
- // Avoid splitting redundantly by making sure the store into the
- // aggregate can actually be split.
- if (failed(getWrittenToFields(dataLayout, subAggregate, storeSize,
- /*storeOffset=*/0)))
- return failure();
-
- return destructurableTypeRange.take_front(exclusiveEnd + 1);
- }
- currentOffset += fieldSize;
- storeSize -= fieldSize;
- }
-
- // If the storeSize is not 0 at this point we are writing past the aggregate
- // as a whole. Abort.
- if (storeSize > 0)
- return failure();
- return destructurableTypeRange.take_front(exclusiveEnd);
-}
-
-/// Splits a store of the vector `value` into `address` at `storeOffset` into
-/// multiple stores of each element with the goal of each generated store
-/// becoming type-consistent through subsequent pattern applications.
-static void splitVectorStore(const DataLayout &dataLayout, Location loc,
- RewriterBase &rewriter, Value address,
- TypedValue<VectorType> value,
- unsigned storeOffset) {
- VectorType vectorType = value.getType();
- unsigned elementSize = dataLayout.getTypeSize(vectorType.getElementType());
-
- // Extract every element in the vector and store it in the given address.
- for (size_t index : llvm::seq<size_t>(0, vectorType.getNumElements())) {
- auto pos =
- r...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/93283
More information about the Mlir-commits
mailing list