[Mlir-commits] [mlir] Revert "[mlir][DLTI] Introduce an operation interface for scoped DLTI queries (#222776)" (PR #232698)
Andrzej Warzyński
llvmlistbot at llvm.org
Fri Oct 16 04:51:29 PDT 2026
https://github.com/banach-space created https://github.com/llvm/llvm-project/pull/232698
This reverts commit efd7e66eb884dba5341ab250e1bd245bab64e92c.
Broken bots:
* https://lab.llvm.org/buildbot/#/builders/240/builds/1016
* https://lab.llvm.org/buildbot/#/builders/241/builds/1040
>From 92c74fd6bb85bcc1d6915ff604904ab1829729be Mon Sep 17 00:00:00 2001
From: Andrzej Warzynski <andrzej.warzynski at arm.com>
Date: Fri, 16 Oct 2026 11:47:19 +0000
Subject: [PATCH] Revert "[mlir][DLTI] Introduce an operation interface for
scoped DLTI queries (#222776)"
This reverts commit efd7e66eb884dba5341ab250e1bd245bab64e92c.
Broken bots:
* https://lab.llvm.org/buildbot/#/builders/240/builds/1016
* https://lab.llvm.org/buildbot/#/builders/241/builds/1040
---
mlir/docs/DataLayout.md | 11 --
mlir/include/mlir/Dialect/DLTI/CMakeLists.txt | 5 -
mlir/include/mlir/Dialect/DLTI/DLTI.h | 16 +--
mlir/include/mlir/Dialect/DLTI/DLTIAttrs.td | 18 +--
.../mlir/Dialect/DLTI/DLTIOpInterfaces.td | 58 --------
.../DLTI/TransformOps/DLTITransformOps.td | 20 +--
mlir/include/mlir/Dialect/Func/IR/FuncOps.h | 1 -
mlir/include/mlir/Dialect/Func/IR/FuncOps.td | 12 +-
mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h | 1 -
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 15 +--
mlir/include/mlir/IR/BuiltinDialect.h | 2 +-
mlir/include/mlir/IR/BuiltinOps.td | 9 +-
.../mlir/Interfaces/DataLayoutInterfaces.td | 6 -
.../mlir/Target/LLVMIR/Transforms/Passes.td | 9 +-
mlir/lib/Dialect/DLTI/CMakeLists.txt | 1 -
mlir/lib/Dialect/DLTI/DLTI.cpp | 107 ++++-----------
mlir/lib/Dialect/Func/IR/CMakeLists.txt | 1 -
mlir/lib/Dialect/Func/IR/FuncOps.cpp | 53 --------
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 49 -------
mlir/lib/IR/BuiltinDialect.cpp | 31 -----
.../Transforms/TargetToTargetFeatures.cpp | 13 +-
.../dlti-query-op-interface-compat.mlir | 27 ----
...lti-query-op-interface-pre-dlti-gpu.mlirbc | Bin 199 -> 0 bytes
.../dlti-query-op-interface-pre-dlti.mlirbc | Bin 290 -> 0 bytes
.../dlti-query-op-interface-roundtrip.mlir | 12 --
mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir | 16 +--
.../ShardToMPI/convert-shard-to-mpi.mlir | 8 +-
.../ShardToMPI/convert-shardshape-to-mpi.mlir | 2 +-
.../Builtin/Bytecode/builtin_fixed_0.mlirbc | Bin 5508 -> 5470 bytes
mlir/test/Dialect/DLTI/invalid.mlir | 5 -
.../test/Dialect/DLTI/query-op-interface.mlir | 32 -----
mlir/test/Dialect/DLTI/query.mlir | 125 +++++-------------
mlir/test/Dialect/MPI/canonicalize.mlir | 2 +-
.../Shard/forward-sharding-propagation.mlir | 2 +-
...et-to-data-layout-and-target-features.mlir | 55 +++-----
.../lib/Dialect/DLTI/TestDataLayoutQuery.cpp | 70 +---------
mlir/test/lib/Dialect/Test/TestOps.td | 6 +-
mlir/unittests/Bytecode/BytecodeTest.cpp | 4 -
38 files changed, 139 insertions(+), 665 deletions(-)
delete mode 100644 mlir/include/mlir/Dialect/DLTI/DLTIOpInterfaces.td
delete mode 100644 mlir/test/Bytecode/dlti-query-op-interface-compat.mlir
delete mode 100644 mlir/test/Bytecode/dlti-query-op-interface-pre-dlti-gpu.mlirbc
delete mode 100644 mlir/test/Bytecode/dlti-query-op-interface-pre-dlti.mlirbc
delete mode 100644 mlir/test/Bytecode/dlti-query-op-interface-roundtrip.mlir
delete mode 100644 mlir/test/Dialect/DLTI/query-op-interface.mlir
diff --git a/mlir/docs/DataLayout.md b/mlir/docs/DataLayout.md
index 984784fb7b894c..a6ee6120e8b921 100644
--- a/mlir/docs/DataLayout.md
+++ b/mlir/docs/DataLayout.md
@@ -295,14 +295,3 @@ The [DLTI](../Dialects/DLTIDialect/) dialect provides the attributes implementin
attribute that can be used to attach the specification to a given operation. The
verifier of this attribute triggers those of the specification and checks the
compatibility of nested specifications.
-
-DLTI key-value queries use a separate interface from data layout scope
-calculation. `DLTIQueryOpInterface` does not require a single backing
-attribute: an operation can answer the key `"architecture"` using an inherent
-`architecture = "sm_90"` attribute, without storing a `dlti` attribute or a
-`#dlti.map`. Its default implementation answers keys from an operation's
-inherent `dlti` attribute.
-`func.func` and `gpu.module` implement this interface; `builtin.module`
-acquires it when the DLTI dialect is loaded. A query searches the target
-operation and then its ancestors until one answers the first key. Further keys
-are queried on the returned attribute through `DLTIQueryInterface`.
diff --git a/mlir/include/mlir/Dialect/DLTI/CMakeLists.txt b/mlir/include/mlir/Dialect/DLTI/CMakeLists.txt
index d42f4663b39902..a7695e67e2d8a5 100644
--- a/mlir/include/mlir/Dialect/DLTI/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/DLTI/CMakeLists.txt
@@ -7,8 +7,3 @@ set(LLVM_TARGET_DEFINITIONS DLTIAttrs.td)
mlir_tablegen(DLTIAttrs.h.inc -gen-attrdef-decls -attrdefs-dialect=dlti)
mlir_tablegen(DLTIAttrs.cpp.inc -gen-attrdef-defs -attrdefs-dialect=dlti)
add_mlir_dialect_tablegen_target(MLIRDLTIAttrsIncGen)
-
-set(LLVM_TARGET_DEFINITIONS DLTIOpInterfaces.td)
-mlir_tablegen(DLTIOpInterfaces.h.inc -gen-op-interface-decls)
-mlir_tablegen(DLTIOpInterfaces.cpp.inc -gen-op-interface-defs)
-add_mlir_generic_tablegen_target(MLIRDLTIOpInterfacesIncGen)
diff --git a/mlir/include/mlir/Dialect/DLTI/DLTI.h b/mlir/include/mlir/Dialect/DLTI/DLTI.h
index 41e42cd940ce0c..84cf6eabc5fa93 100644
--- a/mlir/include/mlir/Dialect/DLTI/DLTI.h
+++ b/mlir/include/mlir/Dialect/DLTI/DLTI.h
@@ -24,25 +24,19 @@ class DataLayoutEntryAttrStorage;
} // namespace mlir
namespace mlir {
namespace dlti {
-namespace detail {
-FailureOr<Attribute> queryDlti(Operation *op, DataLayoutEntryKey key);
-LogicalResult setDlti(Operation *op, DataLayoutEntryKey key, Attribute value);
-} // namespace detail
-
-/// Query the first key at `op` or the nearest ancestor that can answer it.
-/// Query later keys recursively on the returned attributes, without searching
-/// operation ancestors again.
+/// Perform a DLTI-query at `op`, recursively querying each key of `keys` on
+/// query interface-implementing attrs, starting from attr obtained from `op`.
FailureOr<Attribute> query(Operation *op, ArrayRef<DataLayoutEntryKey> keys,
bool emitError = false);
-/// As above, using each string in `keys` as a DLTI entry key.
+/// Perform a DLTI-query at `op` using each string in `keys` as a separate DLTI
+/// entry key, recursively querying on query interface-implementing attrs,
+/// starting from attr obtained from `op`.
FailureOr<Attribute> query(Operation *op, ArrayRef<StringRef> keys,
bool emitError = false);
} // namespace dlti
} // namespace mlir
-#include "mlir/Dialect/DLTI/DLTIOpInterfaces.h.inc"
-
#define GET_ATTRDEF_CLASSES
#include "mlir/Dialect/DLTI/DLTIAttrs.h.inc"
#include "mlir/Dialect/DLTI/DLTIDialect.h.inc"
diff --git a/mlir/include/mlir/Dialect/DLTI/DLTIAttrs.td b/mlir/include/mlir/Dialect/DLTI/DLTIAttrs.td
index 8e16363adf8c33..56bd704d306560 100644
--- a/mlir/include/mlir/Dialect/DLTI/DLTIAttrs.td
+++ b/mlir/include/mlir/Dialect/DLTI/DLTIAttrs.td
@@ -115,14 +115,14 @@ def DLTI_MapAttr : DLTIAttr<"Map", [DLTIQueryInterface]> {
A Data Layout and Target Information map is a list of entries effectively
encoding a dictionary, mapping DLTI-related keys to DLTI-related values.
- This attribute facilitates queries for arbitrary DLTI key-value
- associations. An operation can store it in its inherent `dlti` attribute
- and expose its entries through `DLTIQueryOpInterface`. Query results that
- implement `DLTIQueryInterface` can be queried recursively.
+ This attribute's main purpose is to facilate querying IR for arbitrary
+ key-value associations that encode DLTI. Facility functions exist to perform
+ recursive lookups on nested DLTI-map/query interface-implementing
+ attributes.
Consider the following flat encoding of a single-key dictionary
```
- #dlti.map<"CPU::cache::L1::size_in_bytes" = 65536 : i32>
+ #dlti.map<"CPU::cache::L1::size_in_bytes" = 65536 : i32>>
```
versus nested maps, which make it possible to obtain sub-dictionaries of
related information (with the following example making use of other
@@ -137,12 +137,12 @@ def DLTI_MapAttr : DLTIAttr<"Map", [DLTIQueryInterface]> {
With the flat encoding, the implied structure of the key is ignored, that is
the only successful query (as expressed in the Transform Dialect) is:
`transform.dlti.query ["CPU::cache::L1::size_in_bytes"] at %op`,
- where `%op` handles an operation with the flat `#dlti.map` stored as its
- inherent `dlti` attribute.
+ where `%op` is a handle to an operation which associates the flat-encoding
+ `#dlti.map` attribute.
For querying nested dictionaries, the relevant keys need to be separately
- provided. If `%op` handles an operation with the nested
- `#dlti.target_system_spec` above as its inherent `dlti` attribute, then
+ provided. That is, if `%op` is an handle to an op which has the nesting
+ `#dlti.target_system_spec`-attribute from above attached, then
`transform.dlti.query ["CPU","cache","L1","size_in_bytes"] at %op` gives
back the first leaf value contained. To access the other leaf, we need to do
`transform.dlti.query ["CPU","cache","L1d","size_in_bytes"] at %op`.
diff --git a/mlir/include/mlir/Dialect/DLTI/DLTIOpInterfaces.td b/mlir/include/mlir/Dialect/DLTI/DLTIOpInterfaces.td
deleted file mode 100644
index 074aff85c80b5a..00000000000000
--- a/mlir/include/mlir/Dialect/DLTI/DLTIOpInterfaces.td
+++ /dev/null
@@ -1,58 +0,0 @@
-//===- DLTIOpInterfaces.td - DLTI operation interfaces -----*- tablegen -*-===//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_DIALECT_DLTI_DLTIOPINTERFACES
-#define MLIR_DIALECT_DLTI_DLTIOPINTERFACES
-
-include "mlir/Interfaces/DataLayoutInterfaces.td"
-
-def DLTIQueryOpInterface : OpInterface<"DLTIQueryOpInterface"> {
- let cppNamespace = "::mlir";
-
- let description = [{
- Operation interface for scoped DLTI key-value associations.
-
- Queries and updates operate on one key at a time. The interface does not
- prescribe how an operation stores values. For example, an operation with
- an inherent `architecture = "sm_90"` attribute could return that string
- attribute for the key `"architecture"`, without having a `dlti` attribute
- or a `#dlti.map`.
-
- The default implementations use the operation's inherent `dlti` attribute.
- An update preserves other entries in a `#dlti.map`; setting a null value
- removes an entry. Updating a non-map `dlti` attribute or using a null or
- empty string key fails.
- }];
-
- let methods = [
- InterfaceMethod<
- /*description=*/"Returns the DLTI value associated with the key.",
- /*retTy=*/"::mlir::FailureOr<::mlir::Attribute>",
- /*methodName=*/"queryDlti",
- /*args=*/(ins "::mlir::DataLayoutEntryKey":$key),
- /*methodBody=*/"",
- /*defaultImplementation=*/[{
- return ::mlir::dlti::detail::queryDlti($_op.getOperation(), key);
- }]
- >,
- InterfaceMethod<
- /*description=*/"Creates, replaces, or removes a DLTI key-value "
- "association.",
- /*retTy=*/"::mlir::LogicalResult",
- /*methodName=*/"setDlti",
- /*args=*/(ins "::mlir::DataLayoutEntryKey":$key,
- "::mlir::Attribute":$value),
- /*methodBody=*/"",
- /*defaultImplementation=*/[{
- return ::mlir::dlti::detail::setDlti($_op.getOperation(), key, value);
- }]
- >
- ];
-}
-
-#endif // MLIR_DIALECT_DLTI_DLTIOPINTERFACES
diff --git a/mlir/include/mlir/Dialect/DLTI/TransformOps/DLTITransformOps.td b/mlir/include/mlir/Dialect/DLTI/TransformOps/DLTITransformOps.td
index 748074a8e2c067..f25bb383912d45 100644
--- a/mlir/include/mlir/Dialect/DLTI/TransformOps/DLTITransformOps.td
+++ b/mlir/include/mlir/Dialect/DLTI/TransformOps/DLTITransformOps.td
@@ -19,25 +19,25 @@ def QueryOp : Op<Transform_Dialect, "dlti.query", [
TransformOpInterface, TransformEachOpTrait,
DeclareOpInterfaceMethods<MemoryEffectsOpInterface>
]> {
- let summary = "Return attribute (as param) associated to key via DLTI";
+ let summary = "Return attribute (as param) associated to key via DTLI";
let description = [{
This op queries data layout and target information associated to payload
IR by way of the DLTI dialect.
- A lookup of the first key is performed at `target` or its ancestors by way
- of `DLTIQueryOpInterface`. Ancestors that do not implement the interface or
- cannot answer that key are skipped. Each key should be either a (quoted)
- string or a type. If more than one key is provided, the lookup continues
- recursively on the returned attributes by way of `DLTIQueryInterface`.
- Operation ancestor lookup is not restarted for these later keys. For
- example if the payload IR is
+ A lookup is performed for the given `keys` at `target` op - or its closest
+ interface-implementing ancestor - by way of the `DLTIQueryInterface`, which
+ returns an attribute for a key. Each key should be either a (quoted) string
+ or a type. If more than one key is provided, the lookup continues
+ recursively, now on the returned attributes, with the condition that these
+ implement the above interface. For example if the payload IR is
```
- module attributes {dlti = #dlti.map<"A" = #dlti.map<"B" = 42 : i32>>} {
+ module attributes {#dlti.map = #dlti.map<#dlti.dl_entry<"A",
+ #dlti.map<#dlti.dl_entry<"B", 42: int>>>} {
func.func private @f()
}
```
- and we have that `%func` is a Transform handle to op `@f`, then
+ and we have that `%func` is a Tranform handle to op `@f`, then
`transform.dlti.query ["A", "B"] at %func` returns 42 as a param and
`transform.dlti.query ["A"] at %func` returns the `#dlti.map` attribute
containing just the key "B" and its value. Using `["B"]` or `["A","C"]` as
diff --git a/mlir/include/mlir/Dialect/Func/IR/FuncOps.h b/mlir/include/mlir/Dialect/Func/IR/FuncOps.h
index e18a7dc216574a..e6efeb5ee53f95 100644
--- a/mlir/include/mlir/Dialect/Func/IR/FuncOps.h
+++ b/mlir/include/mlir/Dialect/Func/IR/FuncOps.h
@@ -10,7 +10,6 @@
#define MLIR_DIALECT_FUNC_IR_OPS_H
#include "mlir/Bytecode/BytecodeOpInterface.h"
-#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinTypes.h"
diff --git a/mlir/include/mlir/Dialect/Func/IR/FuncOps.td b/mlir/include/mlir/Dialect/Func/IR/FuncOps.td
index ba5531c46728e4..a99147b380eb33 100644
--- a/mlir/include/mlir/Dialect/Func/IR/FuncOps.td
+++ b/mlir/include/mlir/Dialect/Func/IR/FuncOps.td
@@ -10,7 +10,6 @@
#define MLIR_DIALECT_FUNC_IR_FUNCOPS_TD
include "mlir/IR/EnumAttr.td"
-include "mlir/Dialect/DLTI/DLTIOpInterfaces.td"
include "mlir/IR/OpAsmInterface.td"
include "mlir/IR/SymbolInterfaces.td"
include "mlir/Interfaces/CallInterfaces.td"
@@ -257,7 +256,7 @@ def ConstantOp : Func_Op<"constant",
def FuncOp : Func_Op<"func", [
AffineScope, AutomaticAllocationScope,
SymbolName, SymbolVisibility, FunctionOpInterface, IsolatedFromAbove,
- OpAsmOpInterface, DLTIQueryOpInterface
+ OpAsmOpInterface
]> {
let summary = "An operation with a name containing a single `SSACFG` region";
let description = [{
@@ -272,10 +271,7 @@ def FuncOp : Func_Op<"func", [
the region.
Only dialect attribute names may be specified in the attribute dictionaries
- for function arguments, results, or the function itself. The function may
- also have an inherent `dlti` attribute implementing `DLTIQueryInterface`.
- Its values are available to scoped DLTI queries on the function and nested
- operations, unless a closer scope answers the same key.
+ for function arguments, results, or the function itself.
Example:
@@ -306,9 +302,7 @@ def FuncOp : Func_Op<"func", [
OptionalAttr<StrAttr>:$sym_visibility,
OptionalAttr<DictArrayAttr>:$arg_attrs,
OptionalAttr<DictArrayAttr>:$res_attrs,
- UnitAttr:$no_inline,
- OptionalAttr<DLTIQueryAttr>:$dlti);
- let useCustomPropertiesEncoding = 1;
+ UnitAttr:$no_inline);
let regions = (region AnyRegion:$body);
let builders = [OpBuilder<(ins
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
index 34ff868db7d446..0d6ea9024c8422 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
@@ -15,7 +15,6 @@
#define MLIR_DIALECT_GPU_IR_GPUDIALECT_H
#include "mlir/Bytecode/BytecodeOpInterface.h"
-#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/DLTI/Traits.h"
#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
#include "mlir/Dialect/GPU/IR/GPUDialectDecl.h"
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index d51bac0317d105..ab5c9e0f476608 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -14,7 +14,6 @@
#define GPU_OPS
include "mlir/Dialect/DLTI/DLTIBase.td"
-include "mlir/Dialect/DLTI/DLTIOpInterfaces.td"
include "mlir/Dialect/GPU/IR/GPUBase.td"
include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
include "mlir/Dialect/GPU/IR/CompilationAttrs.td"
@@ -1554,8 +1553,7 @@ def GPU_InitializeNamedBarrierOp
}
def GPU_GPUModuleOp : GPU_Op<"module", [
- IsolatedFromAbove, DataLayoutOpInterface, DLTIQueryOpInterface,
- HasDefaultDLTIDataLayout,
+ IsolatedFromAbove, DataLayoutOpInterface, HasDefaultDLTIDataLayout,
NoRegionArguments, SymbolTable, SymbolName, SymbolVisibility,
Symbol] # GraphRegionNoTerminator.traits> {
let summary = "A top level compilation unit containing code to be run on a GPU.";
@@ -1582,11 +1580,6 @@ def GPU_GPUModuleOp : GPU_Op<"module", [
attribute will be used during the `gpu-module-to-binary` pass to specify the
`OffloadingTranslationAttr` used when creating the `gpu.binary` operation.
- A GPU module may also have an inherent `dlti` attribute implementing
- `DLTIQueryInterface`. It provides values for scoped DLTI queries on the
- module and its nested operations. In custom assembly, write it as
- `dlti = #dlti.map<...>` before the optional attribute dictionary.
-
```
gpu.module @symbol_name {
gpu.func {}
@@ -1614,14 +1607,12 @@ def GPU_GPUModuleOp : GPU_Op<"module", [
SymbolNameAttr:$sym_name,
OptionalAttr<StrAttr>:$sym_visibility,
OptionalAttr<GPUNonEmptyTargetArrayAttr>:$targets,
- OptionalAttr<OffloadingTranslationAttr>:$offloadingHandler,
- OptionalAttr<DLTIQueryAttr>:$dlti);
- let useCustomPropertiesEncoding = 1;
+ OptionalAttr<OffloadingTranslationAttr>:$offloadingHandler);
let regions = (region SizedRegion<1>:$bodyRegion);
let assemblyFormat = [{
($sym_visibility^)? $sym_name
(`<` $offloadingHandler^ `>`)?
- ($targets^)? (`dlti` `=` $dlti^)? attr-dict-with-keyword $bodyRegion
+ ($targets^)? attr-dict-with-keyword $bodyRegion
}];
// We need to ensure the block inside the region is properly terminated;
diff --git a/mlir/include/mlir/IR/BuiltinDialect.h b/mlir/include/mlir/IR/BuiltinDialect.h
index b7c7a271b07fcf..d06b63da6b892d 100644
--- a/mlir/include/mlir/IR/BuiltinDialect.h
+++ b/mlir/include/mlir/IR/BuiltinDialect.h
@@ -26,7 +26,7 @@ struct BuiltinDialectVersion : public mlir::DialectVersion {
int64_t getVersion() const { return version; }
- static BuiltinDialectVersion getCurrentVersion() { return {1}; }
+ static BuiltinDialectVersion getCurrentVersion() { return {0}; }
bool operator<(const BuiltinDialectVersion &other) const {
return version < other.version;
diff --git a/mlir/include/mlir/IR/BuiltinOps.td b/mlir/include/mlir/IR/BuiltinOps.td
index 36d381dbb8d393..fb39d50906493c 100644
--- a/mlir/include/mlir/IR/BuiltinOps.td
+++ b/mlir/include/mlir/IR/BuiltinOps.td
@@ -52,16 +52,10 @@ def ModuleOp : Builtin_Op<"module", [
func.func @foo()
}
```
-
- A module may have an inherent `dlti` attribute implementing
- `DLTIQueryInterface`, for example `dlti = #dlti.map<"chip" = "sm_90">`.
- When the DLTI dialect is loaded, the module implements
- `DLTIQueryOpInterface` and provides values for scoped DLTI queries.
}];
let arguments = (ins OptionalAttr<SymbolNameAttr>:$sym_name,
- OptionalAttr<StrAttr>:$sym_visibility,
- OptionalAttr<DLTIQueryAttr>:$dlti);
+ OptionalAttr<StrAttr>:$sym_visibility);
let regions = (region SizedRegion<1>:$bodyRegion);
let assemblyFormat = "($sym_name^)? prop-dict attr-dict-with-keyword $bodyRegion";
@@ -96,7 +90,6 @@ def ModuleOp : Builtin_Op<"module", [
}
}];
let hasVerifier = 1;
- let useCustomPropertiesEncoding = 1;
// We need to ensure the block inside the region is properly terminated;
// the auto-generated builders do not guarantee that.
diff --git a/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td b/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
index a4520dc3df0576..9d35f78b218f8f 100644
--- a/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
+++ b/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
@@ -43,12 +43,6 @@ def DLTIQueryInterface : AttrInterface<"DLTIQueryInterface"> {
];
}
-def ImplementsDLTIQueryInterface : AttrConstraint<
- CPred<"::llvm::isa<::mlir::DLTIQueryInterface>($_self)">,
- "implementing the `DLTIQueryInterface` attr interface">;
-
-def DLTIQueryAttr : ConfinedAttr<AnyAttr, [ImplementsDLTIQueryInterface]>;
-
def DataLayoutEntryInterface : AttrInterface<"DataLayoutEntryInterface"> {
let cppNamespace = "::mlir";
diff --git a/mlir/include/mlir/Target/LLVMIR/Transforms/Passes.td b/mlir/include/mlir/Target/LLVMIR/Transforms/Passes.td
index 7002333b99976c..c858e9f4c34cb0 100644
--- a/mlir/include/mlir/Target/LLVMIR/Transforms/Passes.td
+++ b/mlir/include/mlir/Target/LLVMIR/Transforms/Passes.td
@@ -28,13 +28,12 @@ def LLVMTargetToDataLayout : Pass<"llvm-target-to-data-layout"> {
}
def LLVMTargetToTargetFeatures : Pass<"llvm-target-to-target-features"> {
- let summary = "Derive and store the target's DLTI features";
- let dependentDialects = ["mlir::DLTIDialect"];
+ let summary = "Update attached #llvm.target's features per the described target";
let description = [{
Obtain the TargetMachine specified by the attached #llvm.target's attributes
- and obtain from it the full list of features of the selected target. Stores
- the resulting features under the `features` key of the operation's DLTI
- map.
+ and obtain from it the full list of features of the selected target. Updates
+ the attached #llvm.target so that its features reflect the full list of
+ features.
}];
let options = [
Option<"initializeLLVMTargets", "initialize-llvm-targets", "bool",
diff --git a/mlir/lib/Dialect/DLTI/CMakeLists.txt b/mlir/lib/Dialect/DLTI/CMakeLists.txt
index 5c366122f25a1b..7ae80f9e913bb3 100644
--- a/mlir/lib/Dialect/DLTI/CMakeLists.txt
+++ b/mlir/lib/Dialect/DLTI/CMakeLists.txt
@@ -6,7 +6,6 @@ add_mlir_dialect_library(MLIRDLTIDialect
DEPENDS
MLIRDLTIIncGen
MLIRDLTIAttrsIncGen
- MLIRDLTIOpInterfacesIncGen
LINK_LIBS PUBLIC
MLIRIR
diff --git a/mlir/lib/Dialect/DLTI/DLTI.cpp b/mlir/lib/Dialect/DLTI/DLTI.cpp
index 021255099c284f..da572f123fdc15 100644
--- a/mlir/lib/Dialect/DLTI/DLTI.cpp
+++ b/mlir/lib/Dialect/DLTI/DLTI.cpp
@@ -16,13 +16,10 @@
#include "mlir/IR/DialectImplementation.h"
#include "llvm/ADT/MapVector.h"
-#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/TypeSwitch.h"
using namespace mlir;
-#include "mlir/Dialect/DLTI/DLTIOpInterfaces.cpp.inc"
-
#include "mlir/Dialect/DLTI/DLTIDialect.cpp.inc"
#define GET_ATTRDEF_CLASSES
@@ -519,61 +516,21 @@ void TargetSystemSpecAttr::print(AsmPrinter &printer) const {
// DLTIDialect
//===----------------------------------------------------------------------===//
-static constexpr llvm::StringLiteral kDltiAttrName = "dlti";
-
-FailureOr<Attribute> mlir::dlti::detail::queryDlti(Operation *op,
- DataLayoutEntryKey key) {
- auto queryable = op->getInherentAttrOfType<DLTIQueryInterface>(kDltiAttrName);
- if (!queryable)
- return failure();
- return queryable.query(key);
-}
-
-LogicalResult mlir::dlti::detail::setDlti(Operation *op, DataLayoutEntryKey key,
- Attribute value) {
- if (key.isNull())
- return failure();
- if (auto stringKey = dyn_cast<StringAttr>(key);
- stringKey && stringKey.getValue().empty())
- return failure();
-
- std::optional<Attribute> inherent = op->getInherentAttr(kDltiAttrName);
- if (!inherent)
- return failure();
- auto current = dyn_cast_or_null<DLTIQueryInterface>(*inherent);
- if (current && !isa<MapAttr>(current))
- return failure();
-
- SmallVector<DataLayoutEntryInterface> entries;
- if (auto map = dyn_cast_or_null<MapAttr>(current))
- llvm::append_range(entries, map.getEntries());
-
- auto makeEntry = [&]() -> DataLayoutEntryInterface {
- if (auto type = dyn_cast<Type>(key))
- return DataLayoutEntryAttr::get(type, value);
- return DataLayoutEntryAttr::get(cast<StringAttr>(key), value);
- };
- auto entry = llvm::find_if(entries, [&](DataLayoutEntryInterface entry) {
- return entry.getKey() == key;
- });
- if (entry != entries.end()) {
- if (value)
- *entry = makeEntry();
- else
- entries.erase(entry);
- } else if (value) {
- entries.push_back(makeEntry());
- } else {
- return success();
- }
+/// Retrieve the first `DLTIQueryInterface`-implementing attribute that is
+/// attached to `op` or such an attr on as close as possible an ancestor. The
+/// op the attribute is attached to is returned as well.
+static std::pair<DLTIQueryInterface, Operation *>
+getClosestQueryable(Operation *op) {
+ DLTIQueryInterface queryable = {};
+
+ // Search op and its ancestors for the first attached DLTIQueryInterface attr.
+ do {
+ for (NamedAttribute attr : op->getAttrs())
+ if ((queryable = dyn_cast<DLTIQueryInterface>(attr.getValue())))
+ break;
+ } while (!queryable && (op = op->getParentOp()));
- StringAttr attrName = StringAttr::get(op->getContext(), kDltiAttrName);
- if (entries.empty()) {
- op->setInherentAttr(attrName, {});
- return success();
- }
- op->setInherentAttr(attrName, MapAttr::get(op->getContext(), entries));
- return success();
+ return std::pair(queryable, op);
}
FailureOr<Attribute>
@@ -589,37 +546,26 @@ dlti::query(Operation *op, ArrayRef<DataLayoutEntryKey> keys, bool emitError) {
return failure();
}
- Attribute currentAttr;
- Operation *queryOp = op;
- for (; queryOp; queryOp = queryOp->getParentOp()) {
- auto queryable = dyn_cast<DLTIQueryOpInterface>(queryOp);
- if (!queryable)
- continue;
- FailureOr<Attribute> result = queryable.queryDlti(keys.front());
- if (succeeded(result)) {
- currentAttr = *result;
- break;
- }
- }
+ auto [queryable, queryOp] = getClosestQueryable(op);
+ Operation *reportOp = (queryOp ? queryOp : op);
- if (!currentAttr) {
+ if (!queryable) {
if (emitError) {
auto diag = op->emitError() << "target op of failed DLTI query";
- diag.attachNote(op->getLoc())
- << "no DLTI-queryable operation on the target or any of its "
- "ancestors could answer key "
- << keyToStr(keys.front());
+ diag.attachNote(reportOp->getLoc())
+ << "no DLTI-queryable attrs on target op or any of its ancestors";
}
return failure();
}
- for (auto &&[idx, key] : llvm::enumerate(keys.drop_front())) {
+ Attribute currentAttr = queryable;
+ for (auto &&[idx, key] : llvm::enumerate(keys)) {
if (auto map = dyn_cast<DLTIQueryInterface>(currentAttr)) {
auto maybeAttr = map.query(key);
if (failed(maybeAttr)) {
if (emitError) {
auto diag = op->emitError() << "target op of failed DLTI query";
- diag.attachNote(queryOp->getLoc())
+ diag.attachNote(reportOp->getLoc())
<< "key " << keyToStr(key)
<< " has no DLTI-mapping per attr: " << map;
}
@@ -630,12 +576,12 @@ dlti::query(Operation *op, ArrayRef<DataLayoutEntryKey> keys, bool emitError) {
if (emitError) {
std::string commaSeparatedKeys;
llvm::interleave(
- keys.take_front(idx + 1), // All prior keys.
+ keys.take_front(idx), // All prior keys.
[&](auto key) { commaSeparatedKeys += keyToStr(key); },
[&]() { commaSeparatedKeys += ","; });
auto diag = op->emitError() << "target op of failed DLTI query";
- diag.attachNote(queryOp->getLoc())
+ diag.attachNote(reportOp->getLoc())
<< "got non-DLTI-queryable attribute upon looking up keys ["
<< commaSeparatedKeys << "] at op";
}
@@ -661,10 +607,6 @@ FailureOr<Attribute> dlti::query(Operation *op, ArrayRef<StringRef> keys,
}
namespace {
-struct ModuleDLTIQueryOpInterface
- : DLTIQueryOpInterface::ExternalModel<ModuleDLTIQueryOpInterface,
- ModuleOp> {};
-
class TargetDataLayoutInterface : public DataLayoutDialectInterface {
public:
using DataLayoutDialectInterface::DataLayoutDialectInterface;
@@ -703,7 +645,6 @@ void DLTIDialect::initialize() {
#include "mlir/Dialect/DLTI/DLTIAttrs.cpp.inc"
>();
addInterfaces<TargetDataLayoutInterface>();
- ModuleOp::attachInterface<ModuleDLTIQueryOpInterface>(*getContext());
}
LogicalResult DLTIDialect::verifyOperationAttribute(Operation *op,
diff --git a/mlir/lib/Dialect/Func/IR/CMakeLists.txt b/mlir/lib/Dialect/Func/IR/CMakeLists.txt
index 987c5954c0a488..329301c6fbafd2 100644
--- a/mlir/lib/Dialect/Func/IR/CMakeLists.txt
+++ b/mlir/lib/Dialect/Func/IR/CMakeLists.txt
@@ -10,7 +10,6 @@ add_mlir_dialect_library(MLIRFuncDialect
LINK_LIBS PUBLIC
MLIRCallInterfaces
MLIRControlFlowInterfaces
- MLIRDLTIDialect
MLIRFunctionInterfaces
MLIRInferTypeOpInterface
MLIRIR
diff --git a/mlir/lib/Dialect/Func/IR/FuncOps.cpp b/mlir/lib/Dialect/Func/IR/FuncOps.cpp
index 79816f84dd8825..6c41eeb36e9cb9 100644
--- a/mlir/lib/Dialect/Func/IR/FuncOps.cpp
+++ b/mlir/lib/Dialect/Func/IR/FuncOps.cpp
@@ -8,7 +8,6 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Bytecode/Encoding.h"
#include "mlir/Conversion/ConvertToEmitC/ToEmitCInterface.h"
#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
#include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
@@ -31,34 +30,11 @@
using namespace mlir;
using namespace mlir::func;
-namespace {
-struct FuncDialectBytecodeInterface : BytecodeDialectInterface {
- using BytecodeDialectInterface::BytecodeDialectInterface;
-
- void writeVersion(DialectBytecodeWriter &writer) const override {
- writer.writeVarInt(1);
- }
-
- std::unique_ptr<DialectVersion>
- readVersion(DialectBytecodeReader &reader) const override {
- uint64_t version;
- if (failed(reader.readVarInt(version)))
- return nullptr;
- if (version != 1) {
- reader.emitError("unsupported func dialect bytecode version");
- return nullptr;
- }
- return std::make_unique<DialectVersion>();
- }
-};
-} // namespace
-
//===----------------------------------------------------------------------===//
// FuncDialect
//===----------------------------------------------------------------------===//
void FuncDialect::initialize() {
- addInterfaces<FuncDialectBytecodeInterface>();
addOperations<
#define GET_OP_LIST
#include "mlir/Dialect/Func/IR/FuncOps.cpp.inc"
@@ -163,35 +139,6 @@ bool ConstantOp::isBuildableWith(Attribute value, Type type) {
// FuncOp
//===----------------------------------------------------------------------===//
-LogicalResult FuncOp::readProperties(DialectBytecodeReader &reader,
- OperationState &state) {
- auto &properties = state.getOrAddProperties<Properties>();
- // Preserve the pre-DLTI field order for unversioned bytecode.
- if (failed(reader.readOptionalAttribute(properties.arg_attrs)) ||
- failed(reader.readAttribute(properties.function_type)) ||
- failed(reader.readOptionalAttribute(properties.no_inline)) ||
- failed(reader.readOptionalAttribute(properties.res_attrs)) ||
- failed(reader.readAttribute(properties.sym_name)) ||
- failed(reader.readOptionalAttribute(properties.sym_visibility)))
- return failure();
-
- if (failed(reader.getDialectVersion<FuncDialect>()))
- return success();
- return reader.readOptionalAttribute(properties.dlti);
-}
-
-void FuncOp::writeProperties(DialectBytecodeWriter &writer) {
- Properties &properties = getProperties();
- writer.writeOptionalAttribute(properties.arg_attrs);
- writer.writeAttribute(properties.function_type);
- writer.writeOptionalAttribute(properties.no_inline);
- writer.writeOptionalAttribute(properties.res_attrs);
- writer.writeAttribute(properties.sym_name);
- writer.writeOptionalAttribute(properties.sym_visibility);
- if (writer.getBytecodeVersion() >= bytecode::kDialectVersioning)
- writer.writeOptionalAttribute(properties.dlti);
-}
-
FuncOp FuncOp::create(Location location, StringRef name, FunctionType type,
ArrayRef<NamedAttribute> attrs) {
OpBuilder builder(location->getContext());
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index ec31892edf1919..7b47c20515b131 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -12,7 +12,6 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Bytecode/Encoding.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Bufferization/IR/BufferDeallocationOpInterface.h"
#include "mlir/Dialect/Math/IR/Math.h"
@@ -50,28 +49,6 @@ using namespace mlir::gpu;
#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
-namespace {
-struct GPUDialectBytecodeInterface : BytecodeDialectInterface {
- using BytecodeDialectInterface::BytecodeDialectInterface;
-
- void writeVersion(DialectBytecodeWriter &writer) const override {
- writer.writeVarInt(1);
- }
-
- std::unique_ptr<DialectVersion>
- readVersion(DialectBytecodeReader &reader) const override {
- uint64_t version;
- if (failed(reader.readVarInt(version)))
- return nullptr;
- if (version != 1) {
- reader.emitError("unsupported gpu dialect bytecode version");
- return nullptr;
- }
- return std::make_unique<DialectVersion>();
- }
-};
-} // namespace
-
//===----------------------------------------------------------------------===//
// GPU Device Mapping Attributes
//===----------------------------------------------------------------------===//
@@ -306,7 +283,6 @@ struct GPUInlinerInterface : public DialectInlinerInterface {
} // namespace
void GPUDialect::initialize() {
- addInterfaces<GPUDialectBytecodeInterface>();
addTypes<AsyncTokenType>();
addTypes<MMAMatrixType>();
addTypes<NamedBarrierType>();
@@ -2007,31 +1983,6 @@ LogicalResult gpu::ReturnOp::verify() {
// GPUModuleOp
//===----------------------------------------------------------------------===//
-LogicalResult GPUModuleOp::readProperties(DialectBytecodeReader &reader,
- OperationState &state) {
- auto &properties = state.getOrAddProperties<Properties>();
- // Preserve the pre-DLTI field order for unversioned bytecode.
- if (failed(reader.readOptionalAttribute(properties.offloadingHandler)) ||
- failed(reader.readAttribute(properties.sym_name)) ||
- failed(reader.readOptionalAttribute(properties.sym_visibility)) ||
- failed(reader.readOptionalAttribute(properties.targets)))
- return failure();
-
- if (failed(reader.getDialectVersion<GPUDialect>()))
- return success();
- return reader.readOptionalAttribute(properties.dlti);
-}
-
-void GPUModuleOp::writeProperties(DialectBytecodeWriter &writer) {
- Properties &properties = getProperties();
- writer.writeOptionalAttribute(properties.offloadingHandler);
- writer.writeAttribute(properties.sym_name);
- writer.writeOptionalAttribute(properties.sym_visibility);
- writer.writeOptionalAttribute(properties.targets);
- if (writer.getBytecodeVersion() >= bytecode::kDialectVersioning)
- writer.writeOptionalAttribute(properties.dlti);
-}
-
void GPUModuleOp::build(OpBuilder &builder, OperationState &result,
StringRef name, ArrayAttr targets,
Attribute offloadingHandler) {
diff --git a/mlir/lib/IR/BuiltinDialect.cpp b/mlir/lib/IR/BuiltinDialect.cpp
index 080b79c36a2e41..c692a4dda518cd 100644
--- a/mlir/lib/IR/BuiltinDialect.cpp
+++ b/mlir/lib/IR/BuiltinDialect.cpp
@@ -103,7 +103,6 @@ struct BuiltinOpAsmDialectInterface : public OpAsmDialectInterface {
} // namespace
void BuiltinDialect::initialize() {
- setVersion(BuiltinDialectVersion::getCurrentVersion());
registerTypes();
registerAttributes();
registerLocationAttributes();
@@ -136,36 +135,6 @@ ModuleOp ModuleOp::create(Location loc, std::optional<StringRef> name) {
return ModuleOp::create(builder, loc, name);
}
-LogicalResult ModuleOp::readProperties(DialectBytecodeReader &reader,
- OperationState &state) {
- auto &properties = state.getOrAddProperties<Properties>();
- if (failed(reader.readOptionalAttribute(properties.sym_name)) ||
- failed(reader.readOptionalAttribute(properties.sym_visibility)))
- return failure();
-
- FailureOr<const DialectVersion *> maybeVersion =
- reader.getDialectVersion<BuiltinDialect>();
- if (failed(maybeVersion) ||
- static_cast<const BuiltinDialectVersion *>(*maybeVersion)->getVersion() <
- 1)
- return success();
- return reader.readOptionalAttribute(properties.dlti);
-}
-
-void ModuleOp::writeProperties(DialectBytecodeWriter &writer) {
- Properties &properties = getProperties();
- writer.writeOptionalAttribute(properties.sym_name);
- writer.writeOptionalAttribute(properties.sym_visibility);
-
- FailureOr<const DialectVersion *> maybeVersion =
- writer.getDialectVersion<BuiltinDialect>();
- if (succeeded(maybeVersion) &&
- static_cast<const BuiltinDialectVersion *>(*maybeVersion)->getVersion() <
- 1)
- return;
- writer.writeOptionalAttribute(properties.dlti);
-}
-
DataLayoutSpecInterface ModuleOp::getDataLayoutSpec() {
// Take the first and only (if present) attribute that implements the
// interface. This needs a linear search, but is called only once per data
diff --git a/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp b/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
index 0d4e994478be61..0cc5f5aabb6308 100644
--- a/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
+++ b/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
@@ -70,12 +70,11 @@ struct TargetToTargetFeaturesPass
auto fullTargetFeaturesAttr =
LLVM::TargetFeaturesAttr::get(&getContext(), plussedFeaturesRefs);
- auto queryOp = dyn_cast<DLTIQueryOpInterface>(op);
- if (!queryOp ||
- failed(queryOp.setDlti(StringAttr::get(&getContext(), "features"),
- fullTargetFeaturesAttr))) {
- op->emitError() << "failed to update DLTI target features";
- return signalPassFailure();
- }
+ auto updatedTargetAttr =
+ LLVM::TargetAttr::get(&getContext(), targetAttr.getTriple(),
+ targetAttr.getChip(), fullTargetFeaturesAttr);
+
+ op->setDiscardableAttr(LLVM::LLVMDialect::getTargetAttrName(),
+ updatedTargetAttr);
}
};
diff --git a/mlir/test/Bytecode/dlti-query-op-interface-compat.mlir b/mlir/test/Bytecode/dlti-query-op-interface-compat.mlir
deleted file mode 100644
index 51b7791bb2651c..00000000000000
--- a/mlir/test/Bytecode/dlti-query-op-interface-compat.mlir
+++ /dev/null
@@ -1,27 +0,0 @@
-// Verify that bytecode written before func.func and gpu.module gained a dlti
-// property can still be read. The fixture was produced from this input by
-// mlir-opt built at fc696f791e06 (the PR's parent) using -emit-bytecode.
-// RUN: mlir-opt %S/dlti-query-op-interface-pre-dlti.mlirbc | FileCheck %s --check-prefix=OLD
-// RUN: mlir-opt %S/dlti-query-op-interface-pre-dlti-gpu.mlirbc | FileCheck %s --check-prefix=OLD-GPU
-// RUN: mlir-opt -emit-bytecode %s | mlir-opt | FileCheck %s --check-prefix=NEW
-
-// The GPU-only fixture uses the same parent revision and this input:
-// module attributes {gpu.container_module} {
-// gpu.module @legacy_gpu {}
-// }
-
-// OLD: module attributes {gpu.container_module}
-// OLD: func.func @legacy()
-// OLD: gpu.module @legacy_gpu
-// OLD-GPU: module attributes {gpu.container_module}
-// OLD-GPU: gpu.module @legacy_gpu
-// NEW: module attributes {gpu.container_module}
-// NEW: func.func @legacy()
-// NEW: gpu.module @legacy_gpu
-module attributes {gpu.container_module} {
- func.func @legacy() {
- return
- }
- gpu.module @legacy_gpu {
- }
-}
diff --git a/mlir/test/Bytecode/dlti-query-op-interface-pre-dlti-gpu.mlirbc b/mlir/test/Bytecode/dlti-query-op-interface-pre-dlti-gpu.mlirbc
deleted file mode 100644
index 315113175c617079b01b4a9bd5dd75d1e73b6059..0000000000000000000000000000000000000000
GIT binary patch
literal 0
HcmV?d00001
literal 199
zcmXX=OA5j;6wLVgj3tyJMY`|+E~Hh&E4UC8v=@+CLk*;9Dec0ec`C191Ys8Qd2g8W
z=~G1K)3Mm?^KHJZ%^<K~FoP6%q%jbyM9D-YlNC!MmK??@JW(lF1{vcT5 at fI>FkK>&
zfh1t+6~sGA*M~Tq-$!Ey<3!z!qG=za)nYR=-R9xltm}34EM2WLA7`zzue`CQ{~hO5
a>xR-e-QWIiR at Y_qy8Vh~iJyrO3itx<b19|(
diff --git a/mlir/test/Bytecode/dlti-query-op-interface-pre-dlti.mlirbc b/mlir/test/Bytecode/dlti-query-op-interface-pre-dlti.mlirbc
deleted file mode 100644
index 3559fd2e57ff8f7a9749211cca905a63eff3bc48..0000000000000000000000000000000000000000
GIT binary patch
literal 0
HcmV?d00001
literal 290
zcmXYr(MrTH5JfYa#u!2gB}g7b7ZIPN?IP|E2qMBRtos8?O|vD?rZpxH{c3-~6n(h|
z=FZ$Z``ynm-|ya!PtVn++H^J`SW_U30T!50NfI6u2E`?(k|&iiS*usVaq>wn6kOvR
zNa778g at SvkV9g0Y_D0}?VW8O-1H%V_X_jLk1YOaCBs8NK;35qG-L9k(kE0Nb7oXAg
z!8&w~PNQy$Xc*hr8{|!h-l3tlz7FO))Gu=|dTdSIrv36O8b6ogR9Y8|KR4QxdK{)E
mRLjK at RO!)$#yaCq|B=1v8a<;woR+T3JDNVvfiJ?K0QdzN*gHf3
diff --git a/mlir/test/Bytecode/dlti-query-op-interface-roundtrip.mlir b/mlir/test/Bytecode/dlti-query-op-interface-roundtrip.mlir
deleted file mode 100644
index 454e1c034e30cc..00000000000000
--- a/mlir/test/Bytecode/dlti-query-op-interface-roundtrip.mlir
+++ /dev/null
@@ -1,12 +0,0 @@
-// RUN: mlir-opt -emit-bytecode %s | mlir-opt | FileCheck %s
-
-// CHECK: module attributes {dlti = #dlti.map<"module" = 1 : i32>, gpu.container_module}
-module attributes {dlti = #dlti.map<"module" = 1 : i32>, gpu.container_module} {
- // CHECK: func.func @new() attributes {dlti = #dlti.map<"func" = 2 : i32>}
- func.func @new() attributes {dlti = #dlti.map<"func" = 2 : i32>} {
- return
- }
- // CHECK: gpu.module @new_gpu dlti = #dlti.map<"gpu" = 3 : i32>
- gpu.module @new_gpu dlti = #dlti.map<"gpu" = 3 : i32> {
- }
-}
diff --git a/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir b/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
index 5beb280ef760b9..6445f5afd80002 100644
--- a/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
+++ b/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
@@ -3,7 +3,7 @@
// RUN: mlir-opt -split-input-file -convert-to-llvm="dynamic=true" %s | FileCheck %s
// COM: Test MPICH ABI
-// CHECK-LABEL: module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">} {
+// CHECK-LABEL: module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">} {
// CHECK-DAG: llvm.func @MPI_Finalize() -> i32
// CHECK-DAG: llvm.func @MPI_Reduce_scatter_block(!llvm.ptr, !llvm.ptr, i32, i32, i32, i32) -> i32
// CHECK-DAG: llvm.func @MPI_Allreduce(!llvm.ptr, !llvm.ptr, i32, i32, i32, i32) -> i32
@@ -14,7 +14,7 @@
// CHECK-DAG: llvm.func @MPI_Send(!llvm.ptr, i32, i32, i32, i32, i32) -> i32
// CHECK-DAG: llvm.func @MPI_Comm_rank(i32, !llvm.ptr) -> i32
// CHECK-DAG: llvm.func @MPI_Init(!llvm.ptr, !llvm.ptr) -> i32
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">} {
+module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">} {
// CHECK-LABEL: llvm.func @test_init_finalize_mpich
func.func @test_init_finalize_mpich() {
@@ -155,7 +155,7 @@ module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">} {
// -----
// COM: Test OpenMPI ABI
-// CHECK-LABEL: module attributes {dlti = #dlti.map<"MPI:Implementation" = "OpenMPI">} {
+// CHECK-LABEL: module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "OpenMPI">} {
// CHECK-DAG: llvm.func @MPI_Finalize() -> i32
// CHECK-DAG: llvm.func @MPI_Comm_split(!llvm.ptr, i32, i32, !llvm.ptr) -> i32
// CHECK-DAG: llvm.func @MPI_Reduce_scatter_block(!llvm.ptr, !llvm.ptr, i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> i32
@@ -169,7 +169,7 @@ module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">} {
// CHECK-DAG: llvm.func @MPI_Comm_rank(!llvm.ptr, !llvm.ptr) -> i32
// CHECK-DAG: llvm.mlir.global external @ompi_mpi_comm_world() {addr_space = 0 : i32} : !llvm.struct<"ompi_communicator_t", opaque>
// CHECK-DAG: llvm.func @MPI_Init(!llvm.ptr, !llvm.ptr) -> i32
-module attributes { dlti = #dlti.map<"MPI:Implementation" = "OpenMPI"> } {
+module attributes { dlti.map = #dlti.map<"MPI:Implementation" = "OpenMPI"> } {
// CHECK-LABEL: llvm.func @test_init_finalize_openmpi
func.func @test_init_finalize_openmpi() {
@@ -315,7 +315,7 @@ module attributes { dlti = #dlti.map<"MPI:Implementation" = "OpenMPI"> } {
// -----
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH", "MPI:comm_world_size" = 4, "MPI:comm_world_rank" = 1> } {
+module attributes {mpi.dlti = #dlti.map<"MPI:Implementation" = "MPICH", "MPI:comm_world_size" = 4, "MPI:comm_world_rank" = 1> } {
// CHECK-LABEL: llvm.func @test_fold
func.func @test_fold(%arg0: memref<100xf32>) {
// CHECK: [[v0:%.*]] = llvm.mlir.constant(1140850688 : i64) : i64
@@ -332,7 +332,7 @@ module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH", "MPI:comm_wo
// COM: Test that an index type that already matches the MPI element count width
// COM: is used as is, both for the offset and for the extents.
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">,
+module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">,
dlti.dl_spec = #dlti.dl_spec<index = 32 : i32>} {
// CHECK-LABEL: llvm.func @test_send_index32
func.func @test_send_index32(%arg0: memref<100xf32>, %rank: i32) {
@@ -357,7 +357,7 @@ module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">,
// COM: Test that an index type narrower than the MPI element count is zero
// COM: extended.
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">,
+module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">,
dlti.dl_spec = #dlti.dl_spec<index = 16 : i32>} {
// CHECK-LABEL: llvm.func @test_send_index16
func.func @test_send_index16(%arg0: memref<100xf32>, %rank: i32) {
@@ -380,7 +380,7 @@ module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">,
// COM: Test that a rank-zero memref, whose descriptor carries no extents, uses
// COM: the element count of one directly.
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "MPICH">} {
+module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">} {
// CHECK-LABEL: llvm.func @test_send_rank_zero
func.func @test_send_rank_zero(%arg0: memref<f32>, %rank: i32) {
// CHECK: [[v0:%.*]] = llvm.insertvalue {{.*}}[2] : !llvm.struct<(ptr, ptr, i64)>
diff --git a/mlir/test/Conversion/ShardToMPI/convert-shard-to-mpi.mlir b/mlir/test/Conversion/ShardToMPI/convert-shard-to-mpi.mlir
index 02612061676d5e..062f05b5c5e131 100644
--- a/mlir/test/Conversion/ShardToMPI/convert-shard-to-mpi.mlir
+++ b/mlir/test/Conversion/ShardToMPI/convert-shard-to-mpi.mlir
@@ -69,7 +69,7 @@ func.func @neighbors_dim2(%arg0 : tensor<120x120x120xi8>) -> (index, index) {
// -----
// CHECK: shard.grid @grid0
-module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
+module attributes { mpi.dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
shard.grid @grid0(shape = 3x4x5)
func.func @process_multi_index() -> (index, index, index) {
// CHECK-DAG: %[[c4:.*]] = arith.constant 4 : index
@@ -107,7 +107,7 @@ module {
}
// -----
-module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 7> } {
+module attributes { mpi.dlti = #dlti.map<"MPI:comm_world_rank" = 7> } {
shard.grid @grid0(shape = 3x4x5)
// CHECK-LABEL: func.func @allreduce_tensor(
func.func @allreduce_tensor(
@@ -291,7 +291,7 @@ func.func @update_halo_1d_first(
}
// -----
-module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 1> } {
+module attributes { mpi.dlti = #dlti.map<"MPI:comm_world_rank" = 1> } {
shard.grid @grid0(shape = 4)
// CHECK-LABEL: func @update_halo_1d_with_zero
func.func @update_halo_1d_with_zero (
@@ -316,7 +316,7 @@ module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 1> } {
}
// -----
-module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
+module attributes { mpi.dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
shard.grid @grid0(shape = 3x4x5)
// CHECK-LABEL: func @update_halo_3d
func.func @update_halo_3d(
diff --git a/mlir/test/Conversion/ShardToMPI/convert-shardshape-to-mpi.mlir b/mlir/test/Conversion/ShardToMPI/convert-shardshape-to-mpi.mlir
index 6337a1f0a178b3..9729d2bfb384e2 100644
--- a/mlir/test/Conversion/ShardToMPI/convert-shardshape-to-mpi.mlir
+++ b/mlir/test/Conversion/ShardToMPI/convert-shardshape-to-mpi.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-opt %s --convert-shard-to-mpi -canonicalize | FileCheck %s
-module attributes { dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
+module attributes { mpi.dlti = #dlti.map<"MPI:comm_world_rank" = 24> } {
// CHECK: shard.grid @grid0
shard.grid @grid0(shape = 3x4x5)
diff --git a/mlir/test/Dialect/Builtin/Bytecode/builtin_fixed_0.mlirbc b/mlir/test/Dialect/Builtin/Bytecode/builtin_fixed_0.mlirbc
index ee79226875ffae00eeceaf8a7bc78b06e147e216..4f661b5a1d233336116b303958bc724529f88b5a 100644
GIT binary patch
delta 144
zcmZqCzNf|S>+?Q{H!ZUwHHCqZpOtZ=kR;b;1Fj=nMUSmn85vo(GqPT1Wc|#@ZpXyP
z9tXr_K>QDg=P@&~>#;Dh?*n2!R-il(hp{oTe*@w^c1HGdK>QDg)i at a0(}1`Qh?fEJ
ZGawe at 1eycH$AI`55Q}j!vior{0s!B*8p{9x
delta 184
zcmcbo)uJur>+?Q{H!ZUwHHCpul$DvCnRz3>Bp2&x5ZJ82b%d+lipiRtk&%&oI}l$7
z;?InXoOVn=A`V29frx(~VjeS)&|?7-`#=OAE0C}Q5n*gV;v0zQ0~*A64n+I|5o#Pj
jA`L{efrw=w;u(k#;{+0AAmSLOAV?6z6ypLCen0{MjB_5M
diff --git a/mlir/test/Dialect/DLTI/invalid.mlir b/mlir/test/Dialect/DLTI/invalid.mlir
index c0d93ce686354a..2436e4f7484f5a 100644
--- a/mlir/test/Dialect/DLTI/invalid.mlir
+++ b/mlir/test/Dialect/DLTI/invalid.mlir
@@ -10,11 +10,6 @@
// -----
-// expected-error at below {{attribute 'dlti' failed to satisfy constraint}}
-module attributes {dlti = 42 : i32} {}
-
-// -----
-
// expected-error at below {{'dlti.dl_spec' is expected to be a #dlti.dl_spec attribute}}
"test.unknown_op"() { dlti.dl_spec = 42 } : () -> ()
diff --git a/mlir/test/Dialect/DLTI/query-op-interface.mlir b/mlir/test/Dialect/DLTI/query-op-interface.mlir
deleted file mode 100644
index 74df84ba8a5458..00000000000000
--- a/mlir/test/Dialect/DLTI/query-op-interface.mlir
+++ /dev/null
@@ -1,32 +0,0 @@
-// RUN: mlir-opt %s -test-dlti-query-op-interface --mlir-print-op-generic | FileCheck %s
-// RUN: mlir-opt %s -emit-bytecode | mlir-opt --mlir-print-op-generic | FileCheck --check-prefix=BYTECODE %s
-
-// BYTECODE: "builtin.module"() <{dlti = #dlti.map<"module" = 1 : i32>}> ({
-module attributes {dlti = #dlti.map<"module" = 1 : i32>} {
- // CHECK: "test.op_with_data_layout"() <{dlti = #dlti.map<"replaced" = 2 : i32, i32 = 32 : i32, "inserted" = 1 : i32>}> ({
- // CHECK-NOT: test.dlti_action
- "test.op_with_data_layout"() ({
- }) {dlti = #dlti.map<"replaced" = 0 : i32,
- "removed" = 0 : i32,
- i32 = 0 : i32>,
- test.dlti_action = "update"} : () -> ()
-
- // CHECK: "test.op_with_data_layout"() ({
- // CHECK-NOT: dlti =
- // CHECK-NOT: test.dlti_action
- "test.op_with_data_layout"() ({
- }) {dlti = #dlti.map<"only" = 0 : i32>,
- test.dlti_action = "clear"} : () -> ()
-
- // CHECK: "test.op_with_data_layout"() <{dlti = #dlti.dl_spec<"immutable" = 0 : i32>}> ({
- // CHECK: }) {test.set_dlti_failed}
- "test.op_with_data_layout"() ({
- }) {dlti = #dlti.dl_spec<"immutable" = 0 : i32>,
- test.dlti_action = "fail"} : () -> ()
-
- // CHECK: "test.op_with_data_layout"() <{dlti = #dlti.map<"kept" = 0 : i32>}> ({
- // CHECK-NOT: test.dlti_action
- "test.op_with_data_layout"() ({
- }) {dlti = #dlti.map<"kept" = 0 : i32>,
- test.dlti_action = "invalid"} : () -> ()
-}
diff --git a/mlir/test/Dialect/DLTI/query.mlir b/mlir/test/Dialect/DLTI/query.mlir
index 08aa97e7cefaed..3825cee6f16164 100644
--- a/mlir/test/Dialect/DLTI/query.mlir
+++ b/mlir/test/Dialect/DLTI/query.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt -transform-interpreter -canonicalize -split-input-file -verify-diagnostics %s | FileCheck %s
// expected-remark @below {{attr associated to "test.id" = 42 : i32}}
-module attributes { dlti = #dlti.map<"test.id" = 42 : i32> } {
+module attributes { test.dlti = #dlti.map<"test.id" = 42 : i32> } {
func.func private @f()
}
@@ -18,7 +18,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-remark @below {{i32 present in set : unit}}
-module attributes { dlti = #dlti.map<i32 = unit> } {
+module attributes { test.dlti = #dlti.map<i32 = unit> } {
func.func private @f()
}
@@ -34,65 +34,8 @@ module attributes {transform.with_named_sequence} {
// -----
-module {
- // expected-remark @below {{associated attr 7 : i32}}
- gpu.module @gpu dlti = #dlti.map<"test.id" = 7 : i32> {
- }
-}
-
-module attributes {transform.with_named_sequence} {
- transform.named_sequence @__transform_main(%arg: !transform.any_op) {
- %gpu_module = transform.structured.match ops{["gpu.module"]} in %arg : (!transform.any_op) -> !transform.any_op
- %param = transform.dlti.query ["test.id"] at %gpu_module : (!transform.any_op) -> !transform.any_param
- transform.debug.emit_param_as_remark %param, "associated attr" at %gpu_module : !transform.any_param, !transform.any_op
- transform.yield
- }
-}
-
-// -----
-
-// A nearer interface-bearing operation that cannot answer the first key does
-// not hide an association on an ancestor.
-module attributes { dlti = #dlti.map<"test.id" = 42 : i32> } {
- // expected-remark @below {{associated attr 42 : i32}}
- func.func private @f() attributes {
- dlti = #dlti.map<"other" = 24 : i32>
- }
-}
-
-module attributes {transform.with_named_sequence} {
- transform.named_sequence @__transform_main(%arg: !transform.any_op) {
- %func = transform.structured.match ops{["func.func"]} in %arg : (!transform.any_op) -> !transform.any_op
- %param = transform.dlti.query ["test.id"] at %func : (!transform.any_op) -> !transform.any_param
- transform.debug.emit_param_as_remark %param, "associated attr" at %func : !transform.any_param, !transform.any_op
- transform.yield
- }
-}
-
-// -----
-
-// Once the first key succeeds, later keys query only the returned attribute.
-module attributes { dlti = #dlti.map<"root" = #dlti.map<"leaf" = 42 : i32>> } {
- // expected-error @below {{target op of failed DLTI query}}
- // expected-note @below {{key "leaf" has no DLTI-mapping per attr: #dlti.map<"other" = 24 : i32>}}
- func.func private @f() attributes {
- dlti = #dlti.map<"root" = #dlti.map<"other" = 24 : i32>>
- }
-}
-
-module attributes {transform.with_named_sequence} {
- transform.named_sequence @__transform_main(%arg: !transform.any_op) {
- %func = transform.structured.match ops{["func.func"]} in %arg : (!transform.any_op) -> !transform.any_op
- // expected-error @below {{'transform.dlti.query' op failed to apply}}
- %param = transform.dlti.query ["root", "leaf"] at %func : (!transform.any_op) -> !transform.any_param
- transform.yield
- }
-}
-
-// -----
-
// expected-remark @below {{attr associated to i32's "width_in_bits" = 32 : i32}}
-module attributes { dlti = #dlti.map<i32 = #dlti.map<"width_in_bits" = 32 : i32>> } {
+module attributes { test.dlti = #dlti.map<i32 = #dlti.map<"width_in_bits" = 32 : i32>> } {
func.func private @f()
}
@@ -110,7 +53,7 @@ module attributes {transform.with_named_sequence} {
// expected-remark @below {{width in bits of i32 = 32 : i64}}
// expected-remark @below {{width in bits of f64 = 64 : i64}}
-module attributes { dlti = #dlti.map<"width_in_bits" = #dlti.map<i32 = 32, f64 = 64>> } {
+module attributes { test.dlti = #dlti.map<"width_in_bits" = #dlti.map<i32 = 32, f64 = 64>> } {
func.func private @f()
}
@@ -129,7 +72,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-remark @below {{attr associated to "test.id" = 42 : i32}}
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
func.func private @f()
}
@@ -145,9 +88,9 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
// expected-remark @below {{attr associated to "test.id" = 24 : i32}}
- func.func private @f() attributes { dlti = #dlti.dl_spec<"test.id" = 24 : i32>}
+ func.func private @f() attributes { test.dlti = #dlti.dl_spec<"test.id" = 24 : i32>}
}
module attributes {transform.with_named_sequence} {
@@ -162,8 +105,8 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-remark @below {{attr associated to "test.id" = 42 : i32}}
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
- func.func private @f() attributes { dlti = #dlti.dl_spec<"test.id" = 24 : i32> }
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+ func.func private @f() attributes { test.dlti = #dlti.dl_spec<"test.id" = 24 : i32> }
}
module attributes {transform.with_named_sequence} {
@@ -178,7 +121,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
func.func @matmul_tensors(
%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>, %arg2: tensor<?x?xf32>)
-> tensor<?x?xf32> {
@@ -201,10 +144,10 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
func.func @matmul_tensors(
%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>, %arg2: tensor<?x?xf32>)
- -> tensor<?x?xf32> attributes {dlti = #dlti.dl_spec<"test.id" = 24 : i32> } {
+ -> tensor<?x?xf32> attributes {test.dlti = #dlti.dl_spec<"test.id" = 24 : i32> } {
// expected-remark @below {{associated attr 24 : i32}}
%0 = linalg.matmul ins(%arg0, %arg1: tensor<?x?xf32>, tensor<?x?xf32>)
outs(%arg2: tensor<?x?xf32>)
@@ -225,7 +168,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-remark @below {{associated attr 42 : i32}}
-module attributes { dlti =
+module attributes { test.dlti =
#dlti.target_system_spec<"CPU" =
#dlti.target_device_spec<"test.id" = 42 : i32>> } {
func.func private @f()
@@ -243,7 +186,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
+module attributes { test.dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
"GPU" = #dlti.target_device_spec<"test.id" = 43 : i32>> } {
// expected-remark @below {{associated attr 43 : i32}}
func.func private @f()
@@ -260,10 +203,10 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
+module attributes { test.dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
"GPU" = #dlti.target_device_spec<"test.id" = 43 : i32>> } {
// expected-remark @below {{associated attr 24 : i32}}
- func.func private @f() attributes { dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 24 : i32>> }
+ func.func private @f() attributes { test.dlti = #dlti.target_system_spec<"CPU" = #dlti.target_device_spec<"test.id" = 24 : i32>> }
}
module attributes {transform.with_named_sequence} {
@@ -277,7 +220,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<
+module attributes { test.dlti = #dlti.target_system_spec<
"CPU" = #dlti.target_device_spec<
"cache::L1::size_in_bytes" = 65536 : i32,
"cache::L1d::size_in_bytes" = 32768 : i32>> } {
@@ -301,7 +244,7 @@ module attributes {transform.with_named_sequence} {
#l1_size = #dlti.map<"size_in_bytes" = 65536 : i32>
#l1d_size = #dlti.map<"size_in_bytes" = 32768 : i32>
-module attributes { dlti =
+module attributes { test.dlti =
#dlti.target_system_spec<"CPU" =
#dlti.target_device_spec<"cache" =
#dlti.map<"L1" = #l1_size,
@@ -324,7 +267,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<
+module attributes { test.dlti = #dlti.target_system_spec<
"CPU" = #dlti.target_device_spec<"inner_most_tile_size" = 42 : i32>> } {
// CHECK-LABEL: func @matmul_tensors
func.func @matmul_tensors(
@@ -355,11 +298,11 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<
+// expected-note @below {{key "NPU" has no DLTI-mapping per attr: #dlti.target_system_spec}}
+module attributes { test.dlti = #dlti.target_system_spec<
"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
"GPU" = #dlti.target_device_spec<"test.id" = 43 : i32>> } {
// expected-error @below {{target op of failed DLTI query}}
- // expected-note @below {{no DLTI-queryable operation on the target or any of its ancestors could answer key "NPU"}}
func.func private @f()
}
@@ -375,7 +318,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-note @below {{key "unspecified" has no DLTI-mapping per attr: #dlti.target_device_spec}}
-module attributes { dlti = #dlti.target_system_spec<
+module attributes { test.dlti = #dlti.target_system_spec<
"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
"GPU" = #dlti.target_device_spec<"test.id" = 43 : i32>> } {
// expected-error @below {{target op of failed DLTI query}}
@@ -393,11 +336,11 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.target_system_spec<
+// expected-note @below {{key "test.id" has no DLTI-mapping per attr: #dlti.target_system_spec}}
+module attributes { test.dlti = #dlti.target_system_spec<
"CPU" = #dlti.target_device_spec<"test.id" = 42 : i32>,
"GPU" = #dlti.target_device_spec<"test.id" = 43 : i32>> } {
// expected-error @below {{target op of failed DLTI query}}
- // expected-note @below {{no DLTI-queryable operation on the target or any of its ancestors could answer key "test.id"}}
func.func private @f()
}
@@ -412,9 +355,9 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
+// expected-note @below {{key "CPU" has no DLTI-mapping per attr: #dlti.dl_spec}}
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32> } {
// expected-error @below {{target op of failed DLTI query}}
- // expected-note @below {{no DLTI-queryable operation on the target or any of its ancestors could answer key "CPU"}}
func.func private @f()
}
@@ -430,7 +373,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-note @below {{got non-DLTI-queryable attribute upon looking up keys ["CPU"]}}
-module attributes { dlti = #dlti.dl_spec<"CPU" = 42 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<"CPU" = 42 : i32> } {
// expected-error @below {{target op of failed DLTI query}}
func.func private @f()
}
@@ -447,7 +390,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-note @below {{got non-DLTI-queryable attribute upon looking up keys [i32]}}
-module attributes { dlti = #dlti.dl_spec<i32 = 32 : i32> } {
+module attributes { test.dlti = #dlti.dl_spec<i32 = 32 : i32> } {
// expected-error @below {{target op of failed DLTI query}}
func.func private @f()
}
@@ -465,7 +408,7 @@ module attributes {transform.with_named_sequence} {
module {
// expected-error @below {{target op of failed DLTI query}}
- // expected-note @below {{no DLTI-queryable operation on the target or any of its ancestors could answer key "CPU"}}
+ // expected-note @below {{no DLTI-queryable attrs on target op or any of its ancestors}}
func.func private @f()
}
@@ -481,7 +424,7 @@ module attributes {transform.with_named_sequence} {
// -----
// expected-note @below {{key i64 has no DLTI-mapping per attr: #dlti.map<i32 = 32 : i64>}}
-module attributes { dlti = #dlti.map<"width_in_bits" = #dlti.map<i32 = 32>>} {
+module attributes { test.dlti = #dlti.map<"width_in_bits" = #dlti.map<i32 = 32>>} {
// expected-error @below {{target op of failed DLTI query}}
func.func private @f()
}
@@ -497,7 +440,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<"test.id" = 42 : i32>} {
+module attributes { test.dlti = #dlti.dl_spec<"test.id" = 42 : i32>} {
func.func private @f()
}
@@ -512,7 +455,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.map<#dlti.dl_entry<"test.id", 42 : i32>>} {
+module attributes { test.dlti = #dlti.map<#dlti.dl_entry<"test.id", 42 : i32>>} {
// expected-error @below {{target op of failed DLTI query}}
// expected-note @below {{no keys provided to attempt query with}}
func.func private @f()
@@ -529,7 +472,7 @@ module attributes {transform.with_named_sequence} {
// -----
-module attributes { dlti = #dlti.dl_spec<#dlti.dl_entry<"test.id", 42 : i32>>} {
+module attributes { test.dlti = #dlti.dl_spec<#dlti.dl_entry<"test.id", 42 : i32>>} {
func.func private @f()
}
@@ -540,4 +483,4 @@ module attributes {transform.with_named_sequence} {
%param = transform.dlti.query ["test.id"] at %funcs : (!transform.any_op) -> !transform.param<i64>
transform.yield
}
-}
+}
\ No newline at end of file
diff --git a/mlir/test/Dialect/MPI/canonicalize.mlir b/mlir/test/Dialect/MPI/canonicalize.mlir
index 76e9398a2c6083..3523d46e212196 100644
--- a/mlir/test/Dialect/MPI/canonicalize.mlir
+++ b/mlir/test/Dialect/MPI/canonicalize.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-opt %s -canonicalize | mlir-opt | FileCheck %s
-module attributes {dlti = #dlti.map<"MPI:comm_world_size" = 12, "MPI:comm_world_rank" = 5> } {
+module attributes {mpi.dlti = #dlti.map<"MPI:comm_world_size" = 12, "MPI:comm_world_rank" = 5> } {
// CHECK-LABEL: func.func @mpi_test
func.func @mpi_test(%ref : memref<100xf32>) -> (i32, i32) {
%comm = mpi.comm_world : !mpi.comm
diff --git a/mlir/test/Dialect/Shard/forward-sharding-propagation.mlir b/mlir/test/Dialect/Shard/forward-sharding-propagation.mlir
index b6f6732fd7d9d8..a5faafb438da0e 100644
--- a/mlir/test/Dialect/Shard/forward-sharding-propagation.mlir
+++ b/mlir/test/Dialect/Shard/forward-sharding-propagation.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt --pass-pipeline="builtin.module(func.func(sharding-propagation{traversal=forward}))" %s | FileCheck %s
#map = affine_map<(d0, d1) -> (d0, d1)>
-module attributes {dlti = #dlti.map<"MPI:Implementation" = "mpich", "MPI:comm_world_rank" = 0 : i32>} {
+module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "mpich", "MPI:comm_world_rank" = 0 : i32>} {
shard.grid "private" @grid(shape = 1)
func.func @test_forward() -> (tensor<6x6xi32>, tensor<6x6xi32>, tensor<i32>) attributes {llvm.emit_c_interface} {
%c1_i32 = arith.constant 1 : i32
diff --git a/mlir/test/Target/LLVMIR/target-to-data-layout-and-target-features.mlir b/mlir/test/Target/LLVMIR/target-to-data-layout-and-target-features.mlir
index 5b042a80269c36..b6b29768605416 100644
--- a/mlir/test/Target/LLVMIR/target-to-data-layout-and-target-features.mlir
+++ b/mlir/test/Target/LLVMIR/target-to-data-layout-and-target-features.mlir
@@ -13,15 +13,14 @@
// TARGET-FEATURES: module attributes
// TARGET-FEATURES-NOT: dlti.dl_spec
-// TARGET-FEATURES-SAME: dlti = #dlti.map<"features" = #llvm.target_features<[
+// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
+// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
+// TARGET-FEATURES-SAME: chip = ""
+// TARGET-FEATURES-SAME: features = <[
// TARGET-FEATURES-SAME: +64bit
// TARGET-FEATURES-NOT: +avx
// TARGET-FEATURES-SAME: +sse
// TARGET-FEATURES-NOT: +mmx
-// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
-// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
-// TARGET-FEATURES-SAME: chip = ""
-// TARGET-FEATURES-NOT: features =
module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
chip = ""> } {
@@ -39,15 +38,14 @@ module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
// TARGET-FEATURES: module attributes
// TARGET-FEATURES-NOT: dlti.dl_spec
-// TARGET-FEATURES-SAME: dlti = #dlti.map<"features" = #llvm.target_features<[
+// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
+// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
+// TARGET-FEATURES-SAME: chip = ""
+// TARGET-FEATURES-SAME: features = <[
// TARGET-FEATURES-SAME: +64bit
// TARGET-FEATURES-NOT: +avx
// TARGET-FEATURES-SAME: +mmx
// TARGET-FEATURES-SAME: +sse
-// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
-// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
-// TARGET-FEATURES-SAME: chip = ""
-// TARGET-FEATURES-SAME: features = <["+mmx", "+sse"]>
module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
chip = "",
@@ -66,17 +64,16 @@ module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
// TARGET-FEATURES: module attributes
// TARGET-FEATURES-NOT: dlti.dl_spec
-// TARGET-FEATURES-SAME: dlti = #dlti.map<"features" = #llvm.target_features<[
+// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
+// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
+// TARGET-FEATURES-SAME: chip = "skylake"
+// TARGET-FEATURES-SAME: features = <[
// TARGET-FEATURES-SAME: +64bit
// TARGET-FEATURES-SAME: +avx
// TARGET-FEATURES-SAME: +avx2
// TARGET-FEATURES-NOT: +avx512f
// TARGET-FEATURES-SAME: +mmx
// TARGET-FEATURES-SAME: +sse
-// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
-// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
-// TARGET-FEATURES-SAME: chip = "skylake"
-// TARGET-FEATURES-NOT: features =
module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
chip = "skylake"> } {
@@ -94,16 +91,15 @@ module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
// TARGET-FEATURES: module attributes
// TARGET-FEATURES-NOT: dlti.dl_spec
-// TARGET-FEATURES-SAME: dlti = #dlti.map<"features" = #llvm.target_features<[
+// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
+// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
+// TARGET-FEATURES-SAME: chip = "skylake"
+// TARGET-FEATURES-SAME: features = <[
// TARGET-FEATURES-SAME: +64bit
// TARGET-FEATURES-NOT: +avx
// TARGET-FEATURES-NOT: +avx2
// TARGET-FEATURES-SAME: +mmx
// TARGET-FEATURES-NOT: +sse
-// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
-// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
-// TARGET-FEATURES-SAME: chip = "skylake"
-// TARGET-FEATURES-SAME: features = <["-sse", "-avx"]>
module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
chip = "skylake",
@@ -122,26 +118,19 @@ module attributes { llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
// DATA-LAYOUT-SAME: features = <["-mmx", "+avx512f"]>
// TARGET-FEATURES: module attributes
-// TARGET-FEATURES-SAME: dlti = #dlti.map<
-// TARGET-FEATURES-SAME: "MPI:comm_world_size" = 4 : i64,
-// TARGET-FEATURES-SAME: "features" = #llvm.target_features<[
+// TARGET-FEATURES-SAME: #dlti.dl_spec<index = 32 : i64>
+// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
+// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
+// TARGET-FEATURES-SAME: chip = "skylake"
+// TARGET-FEATURES-SAME: features = <[
// TARGET-FEATURES-SAME: +64bit
// TARGET-FEATURES-SAME: +avx
// TARGET-FEATURES-SAME: +avx2
// TARGET-FEATURES-SAME: +avx512f
// TARGET-FEATURES-NOT: +mmx
// TARGET-FEATURES-SAME: +sse
-// TARGET-FEATURES-SAME: "triple" = "preserved"
-// TARGET-FEATURES-SAME: #dlti.dl_spec<index = 32 : i64>
-// TARGET-FEATURES-SAME: llvm.target = #llvm.target<
-// TARGET-FEATURES-SAME: triple = "x86_64-unknown-linux"
-// TARGET-FEATURES-SAME: chip = "skylake"
-// TARGET-FEATURES-SAME: features = <["-mmx", "+avx512f"]>
-module attributes { dlti = #dlti.map<"MPI:comm_world_size" = 4,
- "features" = #llvm.target_features<["+old"]>,
- "triple" = "preserved">,
- dlti.dl_spec = #dlti.dl_spec<index = 32>,
+module attributes { dlti.dl_spec = #dlti.dl_spec<index = 32>,
llvm.target = #llvm.target<triple = "x86_64-unknown-linux",
chip = "skylake",
features = <["-mmx", "+avx512f"]>> } {
diff --git a/mlir/test/lib/Dialect/DLTI/TestDataLayoutQuery.cpp b/mlir/test/lib/Dialect/DLTI/TestDataLayoutQuery.cpp
index 9fb89d3b4528b0..bf53a0875571f6 100644
--- a/mlir/test/lib/Dialect/DLTI/TestDataLayoutQuery.cpp
+++ b/mlir/test/lib/Dialect/DLTI/TestDataLayoutQuery.cpp
@@ -107,82 +107,14 @@ struct TestDataLayoutQuery
: legalIntWidths)
});
- });
- }
-};
-
-struct TestDLTIQueryOpInterface
- : public PassWrapper<TestDLTIQueryOpInterface, OperationPass<>> {
- MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestDLTIQueryOpInterface)
-
- StringRef getArgument() const final { return "test-dlti-query-op-interface"; }
- StringRef getDescription() const final {
- return "Test the DLTI query operation interface";
- }
-
- void runOnOperation() override {
- WalkResult result = getOperation()->walk([&](Operation *op) {
- auto action =
- op->getDiscardableAttrOfType<StringAttr>("test.dlti_action");
- if (!action)
- return WalkResult::advance();
-
- auto queryOp = dyn_cast<DLTIQueryOpInterface>(op);
- if (!queryOp) {
- op->emitError("expected DLTIQueryOpInterface");
- return WalkResult::interrupt();
- }
- Builder builder(op->getContext());
- auto key = [&](StringRef value) { return builder.getStringAttr(value); };
- if (action.getValue() == "update") {
- if (failed(queryOp.setDlti(key("inserted"),
- builder.getI32IntegerAttr(1))) ||
- failed(queryOp.setDlti(key("replaced"),
- builder.getI32IntegerAttr(2))) ||
- failed(queryOp.setDlti(builder.getI32Type(),
- builder.getI32IntegerAttr(32))) ||
- failed(queryOp.setDlti(key("removed"), Attribute())) ||
- failed(queryOp.setDlti(key("absent"), Attribute()))) {
- op->emitError("failed to update DLTI map");
- return WalkResult::interrupt();
- }
- } else if (action.getValue() == "clear") {
- if (failed(queryOp.setDlti(key("only"), Attribute())) ||
- failed(queryOp.setDlti(key("absent"), Attribute()))) {
- op->emitError("failed to clear DLTI map");
- return WalkResult::interrupt();
- }
- } else if (action.getValue() == "fail") {
- if (succeeded(
- queryOp.setDlti(key("new"), builder.getI32IntegerAttr(1)))) {
- op->emitError("unexpectedly updated immutable DLTI representation");
- return WalkResult::interrupt();
- }
- op->setDiscardableAttr("test.set_dlti_failed", builder.getUnitAttr());
- } else if (action.getValue() == "invalid") {
- Attribute value = builder.getI32IntegerAttr(1);
- if (succeeded(queryOp.setDlti(key(""), value)) ||
- succeeded(queryOp.setDlti(DataLayoutEntryKey(), value)) ||
- succeeded(queryOp.setDlti(key(""), Attribute()))) {
- op->emitError("unexpectedly accepted an invalid DLTI key");
- return WalkResult::interrupt();
- }
- }
- op->removeDiscardableAttr("test.dlti_action");
- return WalkResult::advance();
});
- if (result.wasInterrupted())
- signalPassFailure();
}
};
} // namespace
namespace mlir {
namespace test {
-void registerTestDataLayoutQuery() {
- PassRegistration<TestDataLayoutQuery>();
- PassRegistration<TestDLTIQueryOpInterface>();
-}
+void registerTestDataLayoutQuery() { PassRegistration<TestDataLayoutQuery>(); }
} // namespace test
} // namespace mlir
diff --git a/mlir/test/lib/Dialect/Test/TestOps.td b/mlir/test/lib/Dialect/Test/TestOps.td
index f95d4419a83ad1..a6b05b55804b7e 100644
--- a/mlir/test/lib/Dialect/Test/TestOps.td
+++ b/mlir/test/lib/Dialect/Test/TestOps.td
@@ -24,7 +24,6 @@ include "mlir/IR/SymbolInterfaces.td"
include "mlir/Interfaces/CallInterfaces.td"
include "mlir/Interfaces/ControlFlowInterfaces.td"
include "mlir/Interfaces/DataLayoutInterfaces.td"
-include "mlir/Dialect/DLTI/DLTIOpInterfaces.td"
include "mlir/Interfaces/DestinationStyleOpInterface.td"
include "mlir/Interfaces/InferIntRangeInterface.td"
include "mlir/Interfaces/InferTypeOpInterface.td"
@@ -3168,13 +3167,10 @@ def MakeTupleOp: TEST_Op<"make_tuple"> {
//===----------------------------------------------------------------------===//
def OpWithDataLayoutOp : TEST_Op<"op_with_data_layout",
- [IsolatedFromAbove, HasDefaultDLTIDataLayout,
- DataLayoutOpInterface,
- DLTIQueryOpInterface]> {
+ [IsolatedFromAbove, HasDefaultDLTIDataLayout, DataLayoutOpInterface]> {
let summary =
"An op that uses DataLayout implementation from the Target dialect";
let regions = (region VariadicRegion<AnyRegion>:$regions);
- let arguments = (ins OptionalAttr<DLTIQueryAttr>:$dlti);
}
def DataLayoutQueryOp : TEST_Op<"data_layout_query"> {
diff --git a/mlir/unittests/Bytecode/BytecodeTest.cpp b/mlir/unittests/Bytecode/BytecodeTest.cpp
index b6eb84fb4e60b3..7e11aa14fda441 100644
--- a/mlir/unittests/Bytecode/BytecodeTest.cpp
+++ b/mlir/unittests/Bytecode/BytecodeTest.cpp
@@ -129,10 +129,6 @@ TEST(Bytecode, AlignmentFailure) {
OwningOpRef<Operation *> module =
parseSourceString<Operation *>(irWithResources, parseConfig);
ASSERT_TRUE(module);
- // Ensure that the resource section needs an explicit alignment marker. The
- // test below exercises the reader's handling of that marker.
- module.get()->setDiscardableAttr("bytecode.padding",
- builder.getStringAttr("padding"));
// Write the module to bytecode.
std::string serializedBytecode;
More information about the Mlir-commits
mailing list