diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp index 118a781570c6..f80310344e0d 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp @@ -338,15 +338,17 @@ LoweringConfigAttr::verify(function_ref emitError, LogicalResult CompilationInfoAttr::verify(function_ref emitError, - LoweringConfigAttr loweringConfig, + LoweringConfigAttrInterface loweringConfig, TranslationInfoAttr translationInfo) { if (!loweringConfig) { return emitError() << "missing lowering config"; } - if (failed(LoweringConfigAttr::verify( - emitError, loweringConfig.getTilingLevels(), - loweringConfig.getNativeVectorSize()))) { - return failure(); + if (auto defaultConfig = llvm::dyn_cast(loweringConfig)) { + if (failed(LoweringConfigAttr::verify( + emitError, defaultConfig.getTilingLevels(), + defaultConfig.getNativeVectorSize()))) { + return emitError() << "invalid lowering config: " << defaultConfig; + } } if (!translationInfo) { return emitError() << "missing translation info"; @@ -356,7 +358,7 @@ CompilationInfoAttr::verify(function_ref emitError, translationInfo.getCodegenSpec(), translationInfo.getWorkgroupSize(), translationInfo.getSubgroupSize(), translationInfo.getConfiguration()))) { - return failure(); + return emitError() << "invalid translation info: " << translationInfo; } return success(); } diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td index dad9349ee251..cef0735cc2c0 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td +++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td @@ -286,21 +286,19 @@ def IREECodegen_CompilationInfoAttr : Specifies the information that allows controlling the compilation of operations like `linalg.matmul`/`linalg.*conv` within IREE. This information is used to override the defaults used by - the IREE compiler. Currently it is only valid to set this on - `linalg.matmul`/`linalg.*conv*` operations. + the IREE compiler. If set on the input to the compiler, there is no + guarantee that the config survives until codegen. Named operations like + `linalg.matmul`/`linalg.*conv*` are more likely to retain their lowering + configurations. TODO: It is expected that the `TranslationInfoAttr` and the `LoweringConfigAttr` are specified. Currently there is no verification that the values of the `LoweringConfigAttr` fully specifies the behaviour of the compilation path chosen with - `TranslationInfoAttr`. This could be added in the future. Note: - Typically the values used for the first-level tiling in - `LoweringConfigAttr` value in the `TranslationInfoAttr` are the - same since the first-level of tile + distribute is already done - at the `Flow` level. This verification is also a TODO. + `TranslationInfoAttr`. This could be added in the future. }]; let parameters = (ins - AttrParameter<"LoweringConfigAttr", "">:$loweringConfig, + AttrParameter<"LoweringConfigAttrInterface", "">:$loweringConfig, AttrParameter<"TranslationInfoAttr", "">:$translationInfo ); diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir index 8979a72f4411..a19c9532d64d 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir +++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir @@ -33,7 +33,7 @@ module { module { func.func @test() attributes { compilation_info = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = >} { return } @@ -48,7 +48,7 @@ module { module { func.func @test() attributes { compilation_info = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = >} { return } diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/lowering_matmul_promotion.mlir b/compiler/src/iree/compiler/Codegen/SPIRV/test/lowering_matmul_promotion.mlir index 2bb5062260eb..27831cfe2d31 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/test/lowering_matmul_promotion.mlir +++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/lowering_matmul_promotion.mlir @@ -7,7 +7,7 @@ // Verify pipelining + multi-buffering. #compilation = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > #pipeline_layout = #hal.pipeline.layout, + lowering_config = #iree_codegen.lowering_config, translation_info = > #pipeline_layout = #hal.pipeline.layout #compilation = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > hal.executable public @batch_matmul_f16_16x4096x4096x64_truncf_mulf { diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_promotion.mlir b/compiler/src/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_promotion.mlir index bdbad94f0240..854fc57f8079 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_promotion.mlir +++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_promotion.mlir @@ -173,7 +173,7 @@ hal.executable @matmul_f16_128x256x64 { ]> #user_config = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > hal.executable @matmul_f16_32x1280x1280 { diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/split_reduction.mlir b/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/split_reduction.mlir index 6b91b5e6f1cd..b9c2742c4450 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/split_reduction.mlir +++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/split_reduction.mlir @@ -1,7 +1,7 @@ // RUN: iree-opt --pass-pipeline='builtin.module(util.func(iree-flow-split-reduction-ops))' --iree-flow-split-matmul-reduction=4 %s | FileCheck %s #compilation = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > util.func public @matmul(%arg0: tensor<100x200xf32>, %arg1: tensor<200x300xf32>, %arg2: tensor<100x300xf32>) -> tensor<100x300xf32> { %0 = linalg.matmul {compilation_info = #compilation} diff --git a/tests/e2e/matmul/generate_e2e_matmul_tests.py b/tests/e2e/matmul/generate_e2e_matmul_tests.py index 596f61086674..15f0e3320969 100644 --- a/tests/e2e/matmul/generate_e2e_matmul_tests.py +++ b/tests/e2e/matmul/generate_e2e_matmul_tests.py @@ -569,7 +569,7 @@ def generate_function( compilation_info_string = ( f"#compilation{generate_function.compilation_index} = " "#iree_codegen.compilation_info<\n" - f" lowering_config = ,\n" + f" lowering_config = #iree_codegen.lowering_config,\n" f" translation_info = <{compiler_pipeline} {compilation_info.workgroup_size_str()}\n" f" {subgroup_size_str},\n" f" {{ pipeline_depth = {compilation_info.software_pipeline_depth}, " diff --git a/tests/e2e/regression/lowering_config.mlir b/tests/e2e/regression/lowering_config.mlir index c2bfa8ac80fb..bbf63013aa27 100644 --- a/tests/e2e/regression/lowering_config.mlir +++ b/tests/e2e/regression/lowering_config.mlir @@ -1,11 +1,11 @@ #compilation0 = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > #compilation1 = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > #compilation2 = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > func.func @lowering_config_test() { @@ -24,11 +24,11 @@ func.func @lowering_config_test() { // Conv dims: N, OH, OW, OC, KH, KW, (IC) // Remove H #conv_compilation0 = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > // Remove W #conv_compilation1 = #iree_codegen.compilation_info< - lowering_config = , + lowering_config = #iree_codegen.lowering_config, translation_info = > func.func @conv() { %input = util.unfoldable_constant dense<1.0> : tensor<36x7x7x512xf32>