Refactor IREECodegenAttrs to use typed array parameters (#15032)
This reworks the attributes in IREECodegenAttrs so that typed [Optional]ArrayRefParameters are used instead of untyped ArrayAttrs.
Doing this removes the manual type-checking and manual conversions to/from SmallVector<T>.
This did require a small restructuring of the lowering_config attribute, so instead of having multiple nested arrays IntegerAttrs for tiling levels, it now has one array of TilingLevelAttrs. The TilingLevelAttr then contains the tile sizes and tile interchange values for a particular level.
The TilingLevelAttr has a custom assembly format that keeps its syntax as close as possible to the previous syntax, to avoid as much test churn as possible. We plan to follow up this patch with another change that extends the TilingLevelAttr to support scalable sizes (using the canonical MLIR syntax), which will need a custom parser anyway, so we don't see this as a dealbreaker.
Example configs:
```
// No tile interchange (no change in syntax)
#iree_codegen.lowering_config<tile_sizes = [[2, 64, 64, 64],
[1, 1, 1, 4],
[0, 0, 0, 0]]>
// With tile interchange
// Before:
#iree_codegen.lowering_config<
tile_sizes = [[32, 64, 0], [8, 32, 0], [0, 0, 16]],
tile_interchange = [[1, 0, 2], [], []]>
// After:
#iree_codegen.lowering_config<tile_sizes = [
{sizes=[32, 64, 0], interchange=[1, 0, 2]}, [8, 32, 0], [0, 0, 16]]>
```
---
These changes makes the `lowering_config` easier to extend for a follow on patch that starts adding scalable tile sizes.diff --git a/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir b/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir
index 4f0093a..d1facc1 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_to_workgroups.mlir
@@ -1698,7 +1698,7 @@
// -----
-#config = #iree_codegen.lowering_config<tile_sizes = [[32, 64, 0], [8, 32, 0], [0, 0, 16]], tile_interchange = [[1, 0, 2], [], []]>
+#config = #iree_codegen.lowering_config<tile_sizes = [{sizes=[32, 64, 0], interchange=[1, 0, 2]}, [8, 32, 0], [0, 0, 16]]>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.cpp b/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.cpp
index 07be9b3..bc43581 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.cpp
@@ -23,48 +23,14 @@
namespace mlir {
namespace iree_compiler {
-//===----------------------------------------------------------------------===//
-// Utility function for common code patterns.
-//===----------------------------------------------------------------------===//
-
-static bool checkIntegerArrayAttr(ArrayAttr arrayAttr) {
- return !llvm::any_of(
- arrayAttr, [](Attribute attr) { return !llvm::isa<IntegerAttr>(attr); });
-}
-
-/// Returns an `ArrayAttr` where each element is an `IntegerAttr` of `IndexType`
-/// whose values is obtained from `values`.
-static ArrayAttr getIndexIntegerArrayAttr(MLIRContext *context,
- ArrayRef<int64_t> values) {
- auto attrs =
- llvm::map_to_vector(values, [&context](int64_t value) -> Attribute {
- return IntegerAttr::get(IndexType::get(context), APInt(64, value));
- });
- return ArrayAttr::get(context, attrs);
-}
-
/// Returns an `ArrayAttr` where each element is an `IntegerAttr` of 64-bit
/// integer type whose values is obtained from `values`.
-static ArrayAttr getI64IntegerArrayAttr(MLIRContext *context,
- ArrayRef<int64_t> values) {
- auto attrs =
- llvm::map_to_vector(values, [&context](int64_t value) -> Attribute {
- return IntegerAttr::get(IntegerType::get(context, 64),
- APInt(64, value));
- });
- return ArrayAttr::get(context, attrs);
-}
-
-/// Assumes that `arrayAttr` is a list of `IntegerAttr`s and returns the values
-/// in these attributes as a vector.
-static SmallVector<int64_t> getIntegerVals(ArrayAttr arrayAttr) {
- if (!arrayAttr)
- return {};
- SmallVector<int64_t> values(arrayAttr.size());
- for (auto [index, attr] : llvm::enumerate(arrayAttr)) {
- values[index] = llvm::cast<IntegerAttr>(attr).getInt();
- }
- return values;
+static ArrayAttr getIndexArrayAttr(MLIRContext *context,
+ ArrayRef<int64_t> values) {
+ return ArrayAttr::get(
+ context, llvm::map_to_vector(values, [&](int64_t value) -> Attribute {
+ return IntegerAttr::get(IndexType::get(context), APInt(64, value));
+ }));
}
namespace IREE {
@@ -76,23 +42,17 @@
LogicalResult
ExportConfigAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- ArrayAttr workgroupSize) {
- if (!workgroupSize) {
- return success();
- }
+ ArrayRef<int64_t> workgroupSize) {
if (workgroupSize.size() > 3) {
return emitError() << "expected workgroup size to have atmost 3 entries";
}
- if (!llvm::all_of(workgroupSize, [](Attribute attr) {
- auto intAttr = llvm::dyn_cast<IntegerAttr>(attr);
- return intAttr && intAttr.getType().isIndex();
- })) {
- return emitError()
- << "expected workgroup size to contain values of index type";
- }
return success();
}
+ArrayAttr ExportConfigAttr::getWorkgroupSizeIndexArray() {
+ return getIndexArrayAttr(getContext(), getWorkgroupSize());
+}
+
//===----------------------------------------------------------------------===//
// iree_codegen.translation_info
//===----------------------------------------------------------------------===//
@@ -127,6 +87,71 @@
}
//===----------------------------------------------------------------------===//
+// iree_codegen.lowering_config_level
+//===----------------------------------------------------------------------===//
+
+void LoweringConfigTilingLevelAttr::print(mlir::AsmPrinter &printer) const {
+ auto tileInterchange = getInterchange();
+ if (tileInterchange.empty()) {
+ printer << '[';
+ printer.printStrippedAttrOrType(getSizes());
+ printer << ']';
+ } else {
+ printer << "{sizes = [";
+ printer.printStrippedAttrOrType(getSizes());
+ printer << "], interchange = [";
+ printer.printStrippedAttrOrType(tileInterchange);
+ printer << "]}";
+ }
+}
+
+Attribute LoweringConfigTilingLevelAttr::parse(mlir::AsmParser &parser,
+ mlir::Type) {
+ auto loc = parser.getCurrentLocation();
+ auto parseListOfI64 = [&](bool prefixChecked =
+ false) -> FailureOr<SmallVector<int64_t>> {
+ if (!prefixChecked && parser.parseLSquare())
+ return failure();
+ if (parser.parseOptionalRSquare().succeeded()) {
+ // Empty list.
+ return SmallVector<int64_t>();
+ }
+ auto list = FieldParser<SmallVector<int64_t>>::parse(parser);
+ if (failed(list)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "failed to parse list of i64s");
+ return failure();
+ }
+ if (parser.parseRSquare())
+ return failure();
+ return list;
+ };
+ if (parser.parseOptionalLSquare().succeeded()) {
+ // Case 1: Simple list of tile sizes, e.g.:
+ // [0, 32, 16]
+ auto tileSizes = parseListOfI64(/*prefixChecked=*/true);
+ if (failed(tileSizes))
+ return {};
+ return parser.getChecked<LoweringConfigTilingLevelAttr>(
+ loc, parser.getContext(), *tileSizes, ArrayRef<int64_t>{});
+ }
+ // Case 2: sizes and interchange, e.g.:
+ // {sizes = [0, 32, 16], interchange = [0, 1, 2]}
+ if (parser.parseLBrace() || parser.parseKeyword("sizes") ||
+ parser.parseEqual())
+ return {};
+ auto tileSizes = parseListOfI64();
+ if (failed(tileSizes) || parser.parseComma() ||
+ parser.parseKeyword("interchange") || parser.parseEqual())
+ return {};
+ auto tileInterchange = parseListOfI64();
+ if (failed(tileInterchange) || parser.parseRBrace())
+ return {};
+ return parser.getChecked<LoweringConfigTilingLevelAttr>(
+ loc, parser.getContext(), *tileSizes, *tileInterchange);
+}
+
+//===----------------------------------------------------------------------===//
// iree_codegen.lowering_config
//===----------------------------------------------------------------------===//
@@ -134,81 +159,48 @@
TileSizesListTypeRef tileSizes,
TileSizesListTypeRef tileInterchange,
ArrayRef<int64_t> nativeVectorSize) {
- auto attrList = [&](TileSizesListTypeRef lst) {
- return llvm::map_to_vector(lst, [&](ArrayRef<int64_t> sizes) -> Attribute {
- return getI64IntegerArrayAttr(context, sizes);
- });
- };
- ArrayAttr tileSizesAttr = ArrayAttr::get(context, attrList(tileSizes));
- ArrayAttr tileInterchangeAttr =
- ArrayAttr::get(context, attrList(tileInterchange));
- ArrayAttr nativeVectorSizeAttr =
- getI64IntegerArrayAttr(context, nativeVectorSize);
- return get(context, tileSizesAttr, tileInterchangeAttr, nativeVectorSizeAttr);
+ SmallVector<LoweringConfigTilingLevelAttr> tilinglevels;
+ for (auto [level, sizes] : llvm::enumerate(tileSizes)) {
+ ArrayRef<int64_t> interchange = level < tileInterchange.size()
+ ? tileInterchange[level]
+ : ArrayRef<int64_t>{};
+ tilinglevels.push_back(
+ LoweringConfigTilingLevelAttr::get(context, sizes, interchange));
+ }
+ return get(context,
+ LoweringConfigTilingLevelsAttr::get(context, tilinglevels),
+ nativeVectorSize);
}
TileSizesListType LoweringConfigAttr::getTileSizeVals() {
- auto tileSizesAttr = getTileSizes();
- if (!tileSizesAttr)
- return {};
TileSizesListType tileSizes;
- for (auto attr : tileSizesAttr) {
- auto vals = getIntegerVals(llvm::cast<ArrayAttr>(attr));
- tileSizes.emplace_back(std::move(vals));
- }
+ for (auto &level : getTilingLevels())
+ tileSizes.push_back(SmallVector<int64_t>(level.getSizes()));
return tileSizes;
}
SmallVector<int64_t> LoweringConfigAttr::getTileSizeVals(unsigned level) {
- ArrayAttr tileSizesAttr = getTileSizes();
- if (!tileSizesAttr || tileSizesAttr.size() <= level)
+ auto levels = getTilingLevels();
+ if (level >= levels.size())
return {};
- return getIntegerVals(llvm::cast<ArrayAttr>(tileSizesAttr[level]));
+ return SmallVector<int64_t>(levels[level].getSizes());
}
SmallVector<int64_t>
LoweringConfigAttr::getTileInterchangeVals(unsigned level) {
- ArrayAttr tileInterchangeAttr = getTileInterchange();
- if (!tileInterchangeAttr || tileInterchangeAttr.size() <= level)
+ auto levels = getTilingLevels();
+ if (level >= levels.size())
return {};
- return getIntegerVals(llvm::cast<ArrayAttr>(tileInterchangeAttr[level]));
-}
-
-SmallVector<int64_t> LoweringConfigAttr::getNativeVectorSizeVals() {
- ArrayAttr nativeVectorSizeAttr = getNativeVectorSize();
- if (!nativeVectorSizeAttr)
- return {};
- return getIntegerVals(nativeVectorSizeAttr);
+ return SmallVector<int64_t>(levels[level].getInterchange());
}
LogicalResult
LoweringConfigAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- ArrayAttr tileSizes, ArrayAttr tileInterchange,
- ArrayAttr nativeVectorSize) {
- if (!tileSizes) {
- return emitError() << "expected tile_sizes to be specified (even is "
- "specified as empty)";
- }
- auto hasNonIntElems = [](ArrayAttr sizes) -> bool {
- return llvm::any_of(sizes, [](Attribute attr) {
- auto arrayAttr = llvm::dyn_cast<ArrayAttr>(attr);
- return !arrayAttr || !checkIntegerArrayAttr(arrayAttr);
- });
- };
- if (hasNonIntElems(tileSizes)) {
- return emitError()
- << "expected all elements of tile_sizes to be a list of integers";
- }
- if (tileInterchange && hasNonIntElems(tileInterchange)) {
- return emitError() << "expected all elements of tile_interchange to be a "
- "list of integers";
- }
- if (nativeVectorSize) {
- if (!checkIntegerArrayAttr(nativeVectorSize)) {
- return emitError()
- << "expected native_vector_size to be a list of integer values";
- }
- }
+ LoweringConfigTilingLevelsAttr levels,
+ ArrayRef<int64_t> nativeVectorSizes) {
+ (void)nativeVectorSizes;
+ if (!levels)
+ return emitError() << "missing lowering config levels";
return success();
}
@@ -216,27 +208,16 @@
// iree.compilation_info
//===----------------------------------------------------------------------===//
-CompilationInfoAttr
-CompilationInfoAttr::get(MLIRContext *context, LoweringConfigAttr configAttr,
- TranslationInfoAttr translationInfo,
- ArrayRef<int64_t> workgroupSize,
- std::optional<int64_t> subgroupSize) {
- ArrayAttr workgroupSizeAttr = getI64IntegerArrayAttr(context, workgroupSize);
- return get(context, configAttr, translationInfo, workgroupSizeAttr,
- subgroupSize);
-}
-
LogicalResult CompilationInfoAttr::verify(
function_ref<InFlightDiagnostic()> emitError,
LoweringConfigAttr loweringConfig, TranslationInfoAttr translationInfo,
- ArrayAttr workgroupSize, std::optional<int64_t> subgroupSize) {
+ ArrayRef<int64_t> workgroupSize, std::optional<int64_t> subgroupSize) {
if (!loweringConfig) {
return emitError() << "missing lowering config";
}
- if (failed(
- LoweringConfigAttr::verify(emitError, loweringConfig.getTileSizes(),
- loweringConfig.getTileInterchange(),
- loweringConfig.getNativeVectorSize()))) {
+ if (failed(LoweringConfigAttr::verify(
+ emitError, loweringConfig.getTilingLevels(),
+ loweringConfig.getNativeVectorSize()))) {
return failure();
}
if (!translationInfo) {
@@ -248,21 +229,9 @@
translationInfo.getSoftwarePipelineStoreStage()))) {
return failure();
}
- if (workgroupSize) {
- if (!checkIntegerArrayAttr(workgroupSize)) {
- return emitError() << "expected workgroup_size to be a list of integers";
- }
- }
return success();
}
-SmallVector<int64_t> CompilationInfoAttr::getWorkgroupSizeVals() {
- ArrayAttr workgroupSizeAttr = getWorkgroupSize();
- if (!workgroupSizeAttr)
- return {};
- return getIntegerVals(workgroupSizeAttr);
-}
-
//===----------------------------------------------------------------------===//
// Initialize attributes
//===----------------------------------------------------------------------===//
@@ -291,7 +260,9 @@
SmallVector<int64_t> getWorkgroupSize(IREE::HAL::ExecutableExportOp exportOp) {
if (std::optional<ArrayAttr> workgroupSizeAttrList =
exportOp.getWorkgroupSize()) {
- return getIntegerVals(*workgroupSizeAttrList);
+ return llvm::map_to_vector(*workgroupSizeAttrList, [](auto attr) {
+ return llvm::cast<IntegerAttr>(attr).getInt();
+ });
}
return {};
}
@@ -311,8 +282,7 @@
return failure();
MLIRContext *context = exportOp->getContext();
if (!workgroupSize.empty()) {
- auto attr = getIndexIntegerArrayAttr(context, workgroupSize);
- exportOp->setWorkgroupSizeAttr(attr);
+ exportOp->setWorkgroupSizeAttr(getIndexArrayAttr(context, workgroupSize));
}
if (subgroupSize) {
exportOp->setSubgroupSizeAttr(Builder(context).getIndexAttr(*subgroupSize));
@@ -372,7 +342,7 @@
IREE::Codegen::LoweringConfigAttr configAttr = getLoweringConfig(op);
if (!configAttr)
return 0;
- return configAttr.getTileSizes().size();
+ return configAttr.getTilingLevels().size();
}
void setLoweringConfig(Operation *op,
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.td b/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.td
index d1aaea1..f444c0b 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/IREECodegenAttrs.td
@@ -155,6 +155,24 @@
let genVerifyDecl = 1;
}
+def IREECodegen_LoweringConfigTilingLevelAttr :
+ AttrDef<IREECodegen_Dialect, "LoweringConfigTilingLevel", []>
+{
+ let mnemonic = "lowering_config_level";
+ let parameters = (ins
+ ArrayRefParameter<"int64_t",
+ "The tile sizes to use for this level of tiling">:$sizes,
+ OptionalArrayRefParameter<"int64_t",
+ "The tile interchange to use for this level of tiling">:$interchange);
+ let hasCustomAssemblyFormat = 1;
+}
+
+def IREECodegen_LoweringConfigTilingLevelsAttr :
+ ArrayOfAttr<IREECodegen_Dialect, "LoweringConfigTilingLevels",
+ "lowering_config_levels", "LoweringConfigTilingLevelAttr", []>
+{
+}
+
def IREECodegen_LoweringConfigAttr :
AttrDef<IREECodegen_Dialect, "LoweringConfig", []> {
let mnemonic = "lowering_config";
@@ -171,21 +189,17 @@
implementation. In future, each pass pipeline could verify that
the lowering configuration has all the necessary attributes for
the pipeline.
-
}];
let assemblyFormat = [{
- `<` `tile_sizes` `=` $tileSizes
- (`,` `tile_interchange` `=` $tileInterchange^)?
- (`,` `native_vector_size` `=` $nativeVectorSize^)? `>`
+ `<` `tile_sizes` `=` $tilingLevels
+ (`,` `native_vector_size` `=` `[` $nativeVectorSize^ `]`)? `>`
}];
let parameters = (ins
- AttrParameter<"ArrayAttr",
- "The tile sizes to use for different levels of tiling">:$tileSizes,
- DefaultValuedParameter<"ArrayAttr", "ArrayAttr::get($_ctxt, {})",
- "The tile interchange to use for different levels of tiling">:$tileInterchange,
- DefaultValuedParameter<"ArrayAttr", "ArrayAttr::get($_ctxt, {})",
+ AttrParameter<"LoweringConfigTilingLevelsAttr",
+ "The lowering config at different levels">:$tilingLevels,
+ OptionalArrayRefParameter<"int64_t",
"The native vector size to use for the given operation">:$nativeVectorSize
);
let builders = [
@@ -204,8 +218,11 @@
SmallVector<int64_t> getTileInterchangeVals(unsigned level);
// Returns the native vector size to use.
- SmallVector<int64_t> getNativeVectorSizeVals();
+ SmallVector<int64_t> getNativeVectorSizeVals() {
+ return SmallVector<int64_t>(getNativeVectorSize());
+ }
}];
+
let genVerifyDecl = 1;
}
@@ -233,28 +250,21 @@
let parameters = (ins
AttrParameter<"LoweringConfigAttr", "">:$loweringConfig,
AttrParameter<"TranslationInfoAttr", "">:$translationInfo,
- DefaultValuedParameter<"ArrayAttr", "ArrayAttr::get($_ctxt, {})",
- "The workgroup size to use during translation.">:$workgroupSize,
+ OptionalArrayRefParameter<"int64_t", "The workgroup size to use during translation.">:$workgroupSize,
OptionalParameter<"std::optional<int64_t>",
"The subgroup size to use during translation.">:$subgroupSize
);
let assemblyFormat = [{
`<` `lowering_config` `=` $loweringConfig `,` `translation_info` `=` $translationInfo
- (`,` `workgroup_size` `=` $workgroupSize^)?
+ (`,` `workgroup_size` `=` `[` $workgroupSize^ `]`)?
(`,` `subgroup_size` `=` $subgroupSize^)? `>`
}];
- // The builder is externally for auto-tuner to generate the attributes.
- let builders = [
- AttrBuilder<(ins "LoweringConfigAttr":$configAttr,
- "TranslationInfoAttr":$translationInfo,
- "ArrayRef<int64_t>":$workgroupSize,
- "std::optional<int64_t>":$subgroupSize
- )>,
- ];
let extraClassDeclaration = [{
- SmallVector<int64_t> getWorkgroupSizeVals();
+ SmallVector<int64_t> getWorkgroupSizeVals() {
+ return SmallVector<int64_t>(getWorkgroupSize());
+ }
}];
let genVerifyDecl = 1;
}
@@ -266,10 +276,13 @@
Allows setting workgroup size for pre-formed dispatches.
}];
let parameters = (ins
- AttrParameter<"ArrayAttr", "Workgroup Size to use">:$workgroup_size
+ ArrayRefParameter<"int64_t", "Workgroup Size to use">:$workgroup_size
);
let assemblyFormat = [{
- `<` `workgroup_size` `=` $workgroup_size `>`
+ `<` `workgroup_size` `=` `[` $workgroup_size `]` `>`
+ }];
+ let extraClassDeclaration = [{
+ ArrayAttr getWorkgroupSizeIndexArray();
}];
let genVerifyDecl = 1;
}
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/test/invalid.mlir b/compiler/src/iree/compiler/Codegen/Dialect/test/invalid.mlir
index 2917002..34e6fda 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/test/invalid.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/test/invalid.mlir
@@ -1,20 +1,11 @@
// RUN: iree-opt --split-input-file --verify-diagnostics %s
-module {
- func.func @export_config_invalid_type() attributes {
- // expected-error @+1 {{expected workgroup size to contain values of index type}}
- export_config = #iree_codegen.export_config<workgroup_size = [4, 1]>
- } {
- return
- }
-}
-
// -----
module {
func.func @export_config_invalid_type() attributes {
// expected-error @+1 {{expected workgroup size to have atmost 3 entries}}
- export_config = #iree_codegen.export_config<workgroup_size = [4: index, 1: index, 1: index, 1: index]>
+ export_config = #iree_codegen.export_config<workgroup_size = [4, 1, 1, 1]>
} {
return
}
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir b/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
index fd2c094..206fcde 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
@@ -22,11 +22,11 @@
module {
func.func @test() attributes {
- lowering_config = #iree_codegen.lowering_config<tile_sizes = [[], [10]], tile_interchange = [[], []], native_vector_size = [32, 32]>} {
+ lowering_config = #iree_codegen.lowering_config<tile_sizes = [[], [10]], native_vector_size = [32, 32]>} {
return
}
}
-// CHECK: #config = #iree_codegen.lowering_config<tile_sizes = {{\[}}[], [10]{{\]}}, tile_interchange = {{\[}}[], []], native_vector_size = [32, 32]>
+// CHECK: #config = #iree_codegen.lowering_config<tile_sizes = {{\[}}[], [10]{{\]}}, native_vector_size = [32, 32]>
// -----
@@ -34,8 +34,7 @@
func.func @test() attributes {
compilation_info = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = []>,
- translation_info = <CPUDefault>,
- workgroup_size = []>} {
+ translation_info = <CPUDefault>>} {
return
}
}
@@ -64,9 +63,9 @@
module {
func.func @test() attributes {
- export_config = #iree_codegen.export_config<workgroup_size = [4: index, 1: index]>
+ export_config = #iree_codegen.export_config<workgroup_size = [4, 1]>
} {
return
}
}
-// CHECK: #iree_codegen.export_config<workgroup_size = [4 : index, 1 : index]
+// CHECK: #iree_codegen.export_config<workgroup_size = [4, 1]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
index d05fc84..5095561 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
@@ -223,17 +223,14 @@
}
// Verify interchange
- if (!tilingConfig.getTileInterchange().empty()) {
- for (auto level : llvm::seq<unsigned>(
- 0,
- static_cast<unsigned>(tilingConfig.getTileInterchange().size()))) {
- auto tileSizes = tilingConfig.getTileSizes()[level];
- auto interchange = tilingConfig.getTileInterchangeSizes(level);
- if (!isValidInterchange(interchange, tileSizes.size())) {
- return op->emitOpError("expected [0, ")
- << tileSizes.size()
- << ") to be set exactly once in interchange #" << level;
- }
+ auto tileSizesForLevel = tilingConfig.getTileSizes();
+ for (int level = 0; level < tilingConfig.getNumTilingLevels(); level++) {
+ auto interchange = tilingConfig.getTileInterchangeSizes(level);
+ auto &tileSizes = tileSizesForLevel[level];
+ if (!isValidInterchange(interchange, tileSizes.size())) {
+ return op->emitOpError("expected [0, ")
+ << tileSizes.size() << ") to be set exactly once in interchange #"
+ << level;
}
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.cpp
index bd262c8..c738f05 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.cpp
@@ -28,7 +28,7 @@
// [vector-inner-parallel]]
// 4. [[distribution], [cache-parallel], [cache-reduction],
// [vector-parallel], [vector-reduction]]
- int numTileLevels = loweringConfig.getTileSizes().size();
+ int numTileLevels = loweringConfig.getTilingLevels().size();
switch (numTileLevels) {
case 4:
tilingLevelToActualLevelMap[VectorInnerParallelTiles] = 3;
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.h b/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.h
index 5376999..443ed04 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.h
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/TileSizeSelection.h
@@ -30,7 +30,7 @@
/// Returns the number of tiling levels of the configuration.
unsigned getNumTilingLevels() {
- return loweringConfig.getTileSizes().size();
+ return loweringConfig.getTilingLevels().size();
};
/// Returns the number of dimensions of the configuration. All the tiling
@@ -104,7 +104,6 @@
SmallVector<int64_t> getFusableLevels();
// TODO(dcaballe): Revisit if these features are ever used.
- ArrayAttr getTileInterchange() { return loweringConfig.getTileInterchange(); }
SmallVector<int64_t> getTileInterchangeSizes(unsigned level) {
return loweringConfig.getTileInterchangeVals(level);
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/illegal_configuration.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/illegal_configuration.mlir
index 1c485e0..14555d7 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/illegal_configuration.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/illegal_configuration.mlir
@@ -29,7 +29,7 @@
// -----
-#config = #iree_codegen.lowering_config<tile_sizes = [[4, 8], [8, 8, 0], [0, 0, 8], [0, 0, 0]], tile_interchange = [[], [], [], []], native_vector_size = [0, 0, 4]>
+#config = #iree_codegen.lowering_config<tile_sizes = [[4, 8], [8, 8, 0], [0, 0, 8], [0, 0, 0]], native_vector_size = [0, 0, 4]>
#translation = #iree_codegen.translation_info<CPUDoubleTilingExpert>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
@@ -116,7 +116,7 @@
// -----
-#config = #iree_codegen.lowering_config<tile_sizes = [[4, 8], [8, 8, 0], [0, 0, 8], [0, 0, 0]], tile_interchange = [[1], [], [], []]>
+#config = #iree_codegen.lowering_config<tile_sizes = [{sizes=[4, 8], interchange=[1]}, [8, 8, 0], [0, 0, 8], [0, 0, 0]]>
#translation = #iree_codegen.translation_info<CPUDoubleTilingExpert>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
index 47c415e..1bd9eef 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
@@ -328,8 +328,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[64, 64, 0], [32, 32, 0], [0, 0, 32], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPadExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -1902,7 +1901,7 @@
}
}
}
-// CHECK: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault>
+// CHECK: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault>
// CHECK: hal.executable private @no_compute_ops
// CHECK: hal.executable.export public @test
// CHECK-SAME: translation_info = #[[TRANSLATION]]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/peel_and_vectorize.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/peel_and_vectorize.mlir
index 97c1abb..e863b6b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/peel_and_vectorize.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/peel_and_vectorize.mlir
@@ -4,8 +4,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[64, 64, 0], [8, 32, 0], [0, 0, 16], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPeelingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPeelingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -60,8 +59,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[65, 65, 0], [8, 32, 0], [0, 0, 16], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPeelingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPeelingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -128,8 +126,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[64, 64, 0], [8, 32, 0], [0, 0, 16], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPeelingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPeelingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_tests.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_tests.mlir
index e4d2abd..68f1078 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_tests.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_tests.mlir
@@ -70,8 +70,7 @@
// vector ops will be generated.
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[65, 65], [8, 32, 0], [0, 0, 16], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPadExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -116,8 +115,7 @@
// vector ops will be generated.
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[192, 128, 0], [8, 32, 0], [0, 0, 16], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPadExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/vector_masking.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/vector_masking.mlir
index bae6d8c..02eec62 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/vector_masking.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/vector_masking.mlir
@@ -3,8 +3,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[127, 255], [8, 32], [0, 0]]>,
- translation_info = <CPUDoubleTilingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -73,8 +72,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[127, 255], [8, 0], [0, 32]]>,
- translation_info = <CPUDoubleTilingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -131,8 +129,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[127, 255], [8, 32], [0, 0]]>,
- translation_info = <CPUDoubleTilingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -201,8 +198,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[127, 255], [8, 32], [0, 0]]>,
- translation_info = <CPUDoubleTilingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
@@ -323,8 +319,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[127, 255], [8, 32], [0, 0]]>,
- translation_info = <CPUDoubleTilingExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingExpert>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Verifiers.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/Verifiers.cpp
index 6ac3d2d..eb7b7a0 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Verifiers.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Verifiers.cpp
@@ -99,9 +99,9 @@
ArrayRef<int64_t> rhsShape =
llvm::cast<ShapedType>(op->getOperand(1).getType()).getShape();
- if (loweringConfig.getTileSizes().size() != 1) {
+ if (loweringConfig.getTilingLevels().size() != 1) {
return op->emitOpError("expected 1 levels of tiling sizes, got ")
- << loweringConfig.getTileSizes().size();
+ << loweringConfig.getTilingLevels().size();
}
SmallVector<int64_t> tileSizes =
@@ -200,9 +200,9 @@
}
// Verify that there are four level of tile sizes.
- if (loweringConfig.getTileSizes().size() != 4) {
+ if (loweringConfig.getTilingLevels().size() != 4) {
return op->emitOpError("expected 4 levels of tiling sizes, got ")
- << loweringConfig.getTileSizes().size();
+ << loweringConfig.getTilingLevels().size();
}
ArrayRef<int64_t> lhsShape =
@@ -312,7 +312,7 @@
return success();
}
- const int numTileSizeLevels = loweringConfig.getTileSizes().size();
+ const int numTileSizeLevels = loweringConfig.getTilingLevels().size();
SmallVector<int64_t> workgroupTileSizes =
loweringConfig.getTileSizeVals(kWorkgroupTileLevel);
SmallVector<int64_t> threadTileSizes =
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/illegal_configuration.mlir b/compiler/src/iree/compiler/Codegen/SPIRV/test/illegal_configuration.mlir
index 7faec7c..8f794a7 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/test/illegal_configuration.mlir
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/illegal_configuration.mlir
@@ -41,8 +41,7 @@
#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[32, 64], [4, 4], [0, 0, 4]]>,
- translation_info = <SPIRVMatmulPromoteVectorize>,
- workgroup_size = []>
+ translation_info = <SPIRVMatmulPromoteVectorize>>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
index 716625f..97c63a4 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
@@ -350,7 +350,7 @@
if (attr.getValue().isa<IREE::Codegen::ExportConfigAttr>()) {
workgroupSize = attr.getValue()
.cast<IREE::Codegen::ExportConfigAttr>()
- .getWorkgroupSize();
+ .getWorkgroupSizeIndexArray();
if (workgroupSize.size() < 3) {
SmallVector<Attribute> workgroupSizeVals =
llvm::to_vector(workgroupSize);
diff --git a/experimental/dispatch_profiler/matmul.py b/experimental/dispatch_profiler/matmul.py
index 5ad633c..902708d 100644
--- a/experimental/dispatch_profiler/matmul.py
+++ b/experimental/dispatch_profiler/matmul.py
@@ -203,7 +203,7 @@
#${compilation_info_name} = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[${threadblock_shape_m}, ${threadblock_shape_n}, ${threadblock_shape_k}]]>,
translation_info = <${translation_info} pipeline_depth = ${stages}>,
- workgroup_size = [${block_dim_x} : index, ${block_dim_y} : index, ${block_dim_z} : index]
+ workgroup_size = [${block_dim_x}, ${block_dim_y}, ${block_dim_z}]
>
"""
# batch matmul and split-k matmul compilation info template
@@ -212,7 +212,7 @@
#${compilation_info_name} = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[1, ${threadblock_shape_m}, ${threadblock_shape_n}, ${threadblock_shape_k}]]>,
translation_info = <${translation_info} pipeline_depth = ${stages}>,
- workgroup_size = [${block_dim_x} : index, ${block_dim_y} : index, ${block_dim_z} : index]
+ workgroup_size = [${block_dim_x}, ${block_dim_y}, ${block_dim_z}]
>
"""
@@ -262,7 +262,7 @@
# linalg.matmul mlir template
self.linalg_row_row_matmul_template = """
-// Dispatch linalg.matmul row-row layout
+// Dispatch linalg.matmul row-row layout
func.func @${operation_name}_${compilation_info_name}(
%lhs: tensor<${problem_m}x${problem_k}x${datatype_lhs}>,
%rhs: tensor<${problem_k}x${problem_n}x${datatype_rhs}>) -> tensor<${problem_m}x${problem_n}x${datatype_result}>
@@ -270,7 +270,7 @@
%c0 = arith.constant 0.0 : ${datatype_result}
%init = tensor.empty() : tensor<${problem_m}x${problem_n}x${datatype_result}>
%inital_result = linalg.fill ins(%c0 : ${datatype_result}) outs(%init : tensor<${problem_m}x${problem_n}x${datatype_result}>) -> tensor<${problem_m}x${problem_n}x${datatype_result}>
- %result = linalg.matmul ${compilation_info_attribute}
+ %result = linalg.matmul ${compilation_info_attribute}
ins(%lhs, %rhs: tensor<${problem_m}x${problem_k}x${datatype_lhs}>, tensor<${problem_k}x${problem_n}x${datatype_rhs}>)
outs(%inital_result: tensor<${problem_m}x${problem_n}x${datatype_result}>) -> tensor<${problem_m}x${problem_n}x${datatype_result}>
return %result : tensor<${problem_m}x${problem_n}x${datatype_result}>
diff --git a/samples/custom_dispatch/cuda/kernels/ukernel_example.mlir b/samples/custom_dispatch/cuda/kernels/ukernel_example.mlir
index a66feb9..b162791 100644
--- a/samples/custom_dispatch/cuda/kernels/ukernel_example.mlir
+++ b/samples/custom_dispatch/cuda/kernels/ukernel_example.mlir
@@ -15,7 +15,7 @@
// the workgroup size/block size.
// Note: The name "iree_codegen.export_config" is also important for it to be
// propagated through the compiler.
- attributes {iree_codegen.export_config = #iree_codegen.export_config<workgroup_size = [4 : index]>} {
+ attributes {iree_codegen.export_config = #iree_codegen.export_config<workgroup_size = [4]>} {
%id = flow.dispatch.workgroup.id[0] : index
%count = flow.dispatch.workgroup.count[0] : index
diff --git a/tests/e2e/matmul/generate_e2e_matmul_tests.py b/tests/e2e/matmul/generate_e2e_matmul_tests.py
index ebaedb6..bf0df7f 100644
--- a/tests/e2e/matmul/generate_e2e_matmul_tests.py
+++ b/tests/e2e/matmul/generate_e2e_matmul_tests.py
@@ -91,11 +91,9 @@
# Compilation info
workgroup_size: typing.List[int]
- # Prints the workgroup size as 'index' types
+ # Prints the workgroup size
def workgroup_size_str(self):
- return (
- "[" + ", ".join([f"{size} : index" for size in self.workgroup_size]) + "]"
- )
+ return "[" + ", ".join(map(str, self.workgroup_size)) + "]"
# Returns the list of TestShape's to use for the collection of shapes
diff --git a/tests/e2e/regression/lowering_config.mlir b/tests/e2e/regression/lowering_config.mlir
index 36c9003..42a1329 100644
--- a/tests/e2e/regression/lowering_config.mlir
+++ b/tests/e2e/regression/lowering_config.mlir
@@ -1,15 +1,12 @@
#compilation0 = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[32, 32], [8, 8, 0], [0, 0, 8], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPadExpert>>
#compilation1 = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[64, 64], [4, 4, 0], [0, 0, 4], [0, 0, 0]]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ translation_info = <CPUDoubleTilingPadExpert>>
#compilation2 = #iree_codegen.compilation_info<
- lowering_config = <tile_sizes = [[32, 64], [8, 32, 0], [0, 0, 8], [0, 0, 0]], tile_interchange = [[1, 0], [], [], []]>,
- translation_info = <CPUDoubleTilingPadExpert>,
- workgroup_size = []>
+ lowering_config = <tile_sizes = [{sizes=[32, 64], interchange=[1,0]}, [8, 32, 0], [0, 0, 8], [0, 0, 0]]>,
+ translation_info = <CPUDoubleTilingPadExpert>>
func.func @lowering_config_test() {
%a = util.unfoldable_constant dense<1.0> : tensor<128x256xf32>
@@ -28,13 +25,11 @@
// Remove H
#conv_compilation0 = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[0, 7, 7, 64, 0, 0, 0], [6, 1, 7, 32, 0, 0, 0], [0, 0, 0, 0, 1, 3, 4], [0, 0, 0, 0, 0, 0, 0]]>,
- translation_info = <CPUConvTileAndDecomposeExpert>,
- workgroup_size = []>
+ translation_info = <CPUConvTileAndDecomposeExpert>>
// Remove W
#conv_compilation1 = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[0, 7, 7, 64, 0, 0, 0], [6, 7, 1, 32, 0, 0, 0], [0, 0, 0, 0, 3, 1, 4], [0, 0, 0, 0, 0, 0, 0]]>,
- translation_info = <CPUConvTileAndDecomposeExpert>,
- workgroup_size = []>
+ translation_info = <CPUConvTileAndDecomposeExpert>>
func.func @conv() {
%input = util.unfoldable_constant dense<1.0> : tensor<36x7x7x512xf32>
%filter = util.unfoldable_constant dense<1.0> : tensor<3x3x512x512xf32>
