Integrate LLVM to llvm/llvm-project@6cce5bca3b90 (#24912)
* Add prop-dict to IREE operations for strict property assembly formats
introduced by llvm/llvm-project#196269 and enabled for Transform by
llvm/llvm-project#217290.
* Preserve StableHLO's bare enum syntax after llvm/llvm-project#220608
changed EnumAttr's default assembly format to angle brackets.
* Replace deprecated setName() calls with setSymbolName() following
llvm/llvm-project#218921.
* Cherry-pick e45b19d32e7c (llvm/llvm-project#222395) to fix the missing
RemarkStreamer Bazel dependency exposed by llvm/llvm-project#222300.
Assisted by Codex.
---------
Signed-off-by: Juan I. Pisula <pisula@roofline.ai>
diff --git a/compiler/plugins/target/ROCM/test/tuning_spec_mmt_tile_and_fuse.mlir b/compiler/plugins/target/ROCM/test/tuning_spec_mmt_tile_and_fuse.mlir
index ea2e9a9..d11f982 100644
--- a/compiler/plugins/target/ROCM/test/tuning_spec_mmt_tile_and_fuse.mlir
+++ b/compiler/plugins/target/ROCM/test/tuning_spec_mmt_tile_and_fuse.mlir
@@ -24,7 +24,7 @@
transform.named_sequence @main(%variant_op: !transform.any_op {transform.consumed}) -> (!transform.any_op)
attributes { iree_codegen.tuning_spec_entrypoint } {
- transform.print %variant_op {name="Custom spec"} : !transform.any_op
+ transform.print %variant_op name="Custom spec" : !transform.any_op
%res = transform.foreach_match in %variant_op
@match_mmt -> @apply_op_config
: (!transform.any_op) -> !transform.any_op
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_nested_layout_vector_distribution.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_nested_layout_vector_distribution.mlir
index 31bc004..d63740d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_nested_layout_vector_distribution.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_nested_layout_vector_distribution.mlir
@@ -1504,7 +1504,7 @@
builtin.module attributes { transform.with_named_sequence } {
transform.named_sequence @__transform_main(%variant_op: !transform.any_op {transform.readonly}) {
%top_level_func = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- transform.iree.test_gpu_vector_distribution %top_level_func {workgroup_size = array<i64: 256, 1, 1>} : !transform.any_op
+ transform.iree.test_gpu_vector_distribution %top_level_func <{workgroup_size = array<i64: 256, 1, 1>}> : !transform.any_op
transform.yield
}
}
@@ -1538,7 +1538,7 @@
builtin.module attributes { transform.with_named_sequence } {
transform.named_sequence @__transform_main(%variant_op: !transform.any_op {transform.readonly}) {
%top_level_func = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- transform.iree.test_gpu_vector_distribution %top_level_func {workgroup_size = array<i64: 64, 1, 1>} : !transform.any_op
+ transform.iree.test_gpu_vector_distribution %top_level_func <{workgroup_size = array<i64: 64, 1, 1>}> : !transform.any_op
transform.yield
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_distribution.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_distribution.mlir
index 37e825c..7a83986 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_distribution.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_distribution.mlir
@@ -190,7 +190,7 @@
builtin.module attributes { transform.with_named_sequence } {
transform.named_sequence @__transform_main(%variant_op: !transform.any_op {transform.readonly}) {
%top_level_func = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- transform.iree.test_gpu_vector_distribution %top_level_func {workgroup_size = array<i64: 512, 1, 1>}
+ transform.iree.test_gpu_vector_distribution %top_level_func <{workgroup_size = array<i64: 512, 1, 1>}>
: !transform.any_op
transform.yield
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_reduce_bank_conflicts.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_reduce_bank_conflicts.mlir
index 74139e0..b401dbd 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_reduce_bank_conflicts.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_reduce_bank_conflicts.mlir
@@ -24,7 +24,7 @@
transform.named_sequence @__transform_main(
%variant_op: !transform.any_op {transform.readonly}) {
%0 = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- transform.iree.reduce_shared_memory_bank_conflicts %0 { padding_size_bits = 64 } : (!transform.any_op) -> ()
+ transform.iree.reduce_shared_memory_bank_conflicts %0 <{ padding_size_bits = 64 }> : (!transform.any_op) -> ()
transform.yield
}
} // module
diff --git a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensionsOps.td b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensionsOps.td
index bdff088..1a62a26 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensionsOps.td
+++ b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensionsOps.td
@@ -159,7 +159,7 @@
let arguments = (ins DefaultValuedAttr<BoolAttr, "true">:$useNvGpu);
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
- let assemblyFormat = "attr-dict";
+ let assemblyFormat = "prop-dict attr-dict";
}
def ApplyUnrollVectorsGpuMmaSyncPatternsOp : Op<Transform_Dialect,
@@ -435,7 +435,7 @@
);
let results = (outs TransformHandleTypeInterface:$result);
- let assemblyFormat = "attr-dict $target `:` functional-type($target, results)";
+ let assemblyFormat = "prop-dict attr-dict $target `:` functional-type($target, results)";
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
let builders = [
@@ -561,7 +561,7 @@
);
let results = (outs);
- let assemblyFormat = "$target attr-dict `:` functional-type(operands, results)";
+ let assemblyFormat = "$target prop-dict attr-dict `:` functional-type(operands, results)";
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
let extraClassDeclaration = [{
@@ -675,7 +675,7 @@
DefaultValuedOptionalAttr<I64Attr, "64">:$subgroup_size);
let results = (outs);
- let assemblyFormat = [{ $target attr-dict `:` type($target)}];
+ let assemblyFormat = [{ $target prop-dict attr-dict `:` type($target)}];
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
let extraClassDeclaration = [{
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/external_strategy_spec.mlir b/compiler/src/iree/compiler/Codegen/Common/test/external_strategy_spec.mlir
index 7dc54d3..c206027 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/external_strategy_spec.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/external_strategy_spec.mlir
@@ -2,7 +2,7 @@
module @user_spec attributes { transform.with_named_sequence } {
transform.named_sequence @lowering_strategy(%op: !transform.any_op {transform.readonly}) {
- transform.print {name = "I am external", skip_regions}
+ transform.print name = "I am external" skip_regions
transform.yield
}
transform.named_sequence @import_lowering_strategy(%op: !transform.any_op {transform.readonly}) -> !transform.any_op
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/link_tuning_specs.mlir b/compiler/src/iree/compiler/Codegen/Common/test/link_tuning_specs.mlir
index 19fd909..b0be421 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/link_tuning_specs.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/link_tuning_specs.mlir
@@ -17,7 +17,7 @@
module @foo_module attributes { transform.with_named_sequence } {
transform.named_sequence @foo(%arg0: !transform.any_op {transform.readonly}) -> !transform.any_op
attributes { iree_codegen.tuning_spec_entrypoint } {
- transform.print {name = "Foo", skip_regions}
+ transform.print name = "Foo" skip_regions
transform.yield %arg0 : !transform.any_op
}
}
@@ -26,7 +26,7 @@
transform.named_sequence @bar(%arg0: !transform.any_op {transform.readonly}) -> !transform.any_op
attributes { iree_codegen.tuning_spec_entrypoint } {
transform.match.operation_name %arg0 ["func.func"] : !transform.any_op
- transform.print {name = "Bar", skip_regions}
+ transform.print name = "Bar" skip_regions
transform.yield %arg0 : !transform.any_op
}
}
@@ -34,7 +34,7 @@
module @baz_module attributes { transform.with_named_sequence } {
transform.named_sequence @baz(%arg0: !transform.any_op {transform.readonly}) -> !transform.any_op
attributes { iree_codegen.tuning_spec_entrypoint } {
- transform.print {name = "Baz", skip_regions}
+ transform.print name = "Baz" skip_regions
transform.yield %arg0 : !transform.any_op
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/lowering_config_interpreter.mlir b/compiler/src/iree/compiler/Codegen/Common/test/lowering_config_interpreter.mlir
index 02f2944..7b446c7 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/lowering_config_interpreter.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/lowering_config_interpreter.mlir
@@ -31,15 +31,15 @@
}
transform.named_sequence @print_one(%op: !transform.any_op {transform.readonly}) {
- transform.print %op {name = "one"} : !transform.any_op
+ transform.print %op name = "one" : !transform.any_op
transform.yield
}
transform.named_sequence @print_two(%op: !transform.any_op {transform.readonly}) {
- transform.print %op {name = "two"} : !transform.any_op
+ transform.print %op name = "two" : !transform.any_op
transform.yield
}
transform.named_sequence @print_three(%op: !transform.any_op {transform.readonly}) {
- transform.print %op {name = "three"} : !transform.any_op
+ transform.print %op name = "three" : !transform.any_op
transform.yield
}
}
@@ -73,7 +73,7 @@
}
transform.named_sequence @lowering_strategy(%op: !transform.any_op {transform.readonly}) {
- transform.print {name = "I am internal"}
+ transform.print name = "I am internal"
transform.yield
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/materialize_tuning_specs.mlir b/compiler/src/iree/compiler/Codegen/Common/test/materialize_tuning_specs.mlir
index c34d403..26a470a 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/materialize_tuning_specs.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/materialize_tuning_specs.mlir
@@ -31,7 +31,7 @@
// SKIPLINK-LABEL: module @user_spec
// SKIPLINK-SAME: iree_codegen.tuning_spec_with_default_entrypoint
// SKIPLINK-SAME: transform.with_named_sequence
-// SKIPLINK: transform.print {name = "Hello Tuning Spec"}
+// SKIPLINK: transform.print name = "Hello Tuning Spec"
// SKIPLINK-NOT: module @{{.+}}
// SKIPLINK: module attributes
// SKIPLINK-SAME: iree_codegen.tuning_spec_mlirbc = dense<{{.+}}> : vector<{{[0-9]+}}xi8>
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec.mlir b/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec.mlir
index 78fddf7..9e9e86d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec.mlir
@@ -3,7 +3,7 @@
module @user_spec attributes { transform.with_named_sequence } {
transform.named_sequence @hello(%arg0: !transform.any_op {transform.readonly}) -> !transform.any_op
attributes { iree_codegen.tuning_spec_entrypoint } {
- transform.print {name = "Hello Tuning Spec", skip_regions}
+ transform.print name = "Hello Tuning Spec" skip_regions
transform.yield %arg0 : !transform.any_op
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec_default.mlir b/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec_default.mlir
index d5f83b9..671dcfb 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec_default.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/tuning_spec_default.mlir
@@ -2,7 +2,7 @@
module @user_spec attributes { transform.with_named_sequence, iree_codegen.tuning_spec_with_default_entrypoint } {
transform.named_sequence @match(%arg: !transform.any_op {transform.readonly}) -> (!transform.any_op) {
- transform.print {name = "Hello Tuning Spec"}
+ transform.print name = "Hello Tuning Spec"
transform.yield %arg : !transform.any_op
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/verify_tuning_specs.mlir b/compiler/src/iree/compiler/Codegen/Common/test/verify_tuning_specs.mlir
index 610e0ba..24474fa 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/verify_tuning_specs.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/verify_tuning_specs.mlir
@@ -147,7 +147,7 @@
// expected-error @+1{{'__kernel_config' must start with 'ForeachMatchOp' (required by 'iree_codegen.tuning_spec_with_default_entrypoint')}}
transform.named_sequence @__kernel_config(%arg0: !transform.any_op {transform.consumed})
-> (!transform.any_op) attributes { iree_codegen.tuning_spec_entrypoint } {
- transform.print {name = "Hello"}
+ transform.print name = "Hello"
transform.yield %arg0 : !transform.any_op
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
index 724d541..b6cac8a 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
@@ -242,7 +242,9 @@
//===----------------------------------------------------------------------===//
class IREEGPU_MmaEnumAttr<EnumAttrInfo enumInfo, string name = "">
- : EnumAttr<IREEGPU_Dialect, enumInfo, name>;
+ : EnumAttr<IREEGPU_Dialect, enumInfo, name> {
+ let assemblyFormat = "$value";
+}
def IREEGPU_MMAIntrinsicAttr
: IREEGPU_MmaEnumAttr<IREEGPU_MMAIntrinsic, "mma_intrinsic">;
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/IREEGPUExtensionsOps.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/IREEGPUExtensionsOps.td
index 09f39e4..0d18aa4 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/IREEGPUExtensionsOps.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/IREEGPUExtensionsOps.td
@@ -209,7 +209,7 @@
let results = (outs TransformHandleTypeInterface:$result);
let assemblyFormat = [{
- $producer `into` $consumer attr-dict
+ $producer `into` $consumer prop-dict attr-dict
`:` functional-type(operands, results)
}];
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
index 86beadd..b4cf271 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
@@ -873,7 +873,7 @@
LLVM::InlineAsmOp asmOp = LLVM::InlineAsmOp::create(
rewriter, loc, returnType, inputs, code, constraints,
/*has_side_effects=*/false, /*is_align_stack=*/false,
- LLVM::TailCallKind::None, dialectAttr,
+ LLVM::TailCallKind::None, /*convergent=*/false, dialectAttr,
/*operand_attrs=*/ArrayAttr());
// Extract result vectors from the asm op.
SmallVector<Value> resVec;
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
index e677a71..1236653 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
@@ -78,6 +78,7 @@
/*asm_string=*/asmStr, /*constraints=*/"",
/*has_side_effects=*/true,
/*is_align_stack=*/false, LLVM::TailCallKind::None,
+ /*convergent=*/true,
/*asm_dialect=*/asmDialectAttr,
/*operand_attrs=*/ArrayAttr());
} else if (chipset.majorVersion < 12) {
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensionsOps.td b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensionsOps.td
index a84197f..d733d61 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensionsOps.td
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensionsOps.td
@@ -205,7 +205,7 @@
DefaultValuedAttr<I64Attr, "{}">:$warp_size);
let results = (outs TransformHandleTypeInterface:$result);
- let assemblyFormat = "$target attr-dict `:` functional-type($target, $result)";
+ let assemblyFormat = "$target prop-dict attr-dict `:` functional-type($target, $result)";
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
let builders = [
@@ -359,7 +359,7 @@
let assemblyFormat = [{
$target
- attr-dict
+ prop-dict attr-dict
`:` functional-type($target, results)
}];
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
@@ -442,7 +442,7 @@
let assemblyFormat = [{
$for_op
- attr-dict
+ prop-dict attr-dict
`:` functional-type(operands, results)}];
let extraClassDeclaration = [{
@@ -514,7 +514,7 @@
let assemblyFormat = [{
$target
- attr-dict
+ prop-dict attr-dict
`:` functional-type(operands, results)}];
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
@@ -702,7 +702,7 @@
let assemblyFormat = [{
$target (`test_conversion` $test_conversion^)?
- attr-dict `:` functional-type(operands, results)
+ prop-dict attr-dict `:` functional-type(operands, results)
}];
let cppNamespace = "mlir::iree_compiler::IREE::transform_dialect";
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
index 73c2e6b..a15ca43 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -344,4 +344,4 @@
}
// CHECK-LABEL: llvm.func @global_subgroup_barrier
-// CHECK: llvm.inline_asm has_side_effects asm_dialect = att ";;;WARNING: BREAKS DEBUG WATCHES{{.*}}s_barrier"
+// CHECK: llvm.inline_asm has_side_effects convergent asm_dialect = att ";;;WARNING: BREAKS DEBUG WATCHES{{.*}}s_barrier"
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/create_async_groups.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/create_async_groups.mlir
index 5bfb282..ea5d6b2 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/create_async_groups.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/create_async_groups.mlir
@@ -26,7 +26,7 @@
module attributes { transform.with_named_sequence } {
transform.named_sequence @__transform_main(%variant_op: !transform.any_op {transform.readonly}) {
%top_level_func = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- transform.iree.create_async_groups %top_level_func {use_mma_sync} : (!transform.any_op) -> ()
+ transform.iree.create_async_groups %top_level_func <{use_mma_sync}> : (!transform.any_op) -> ()
transform.yield
}
} // module
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_bufferize.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_bufferize.mlir
index c07aa76..6cf7e0b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_bufferize.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_bufferize.mlir
@@ -34,7 +34,7 @@
transform.named_sequence @__transform_main(%arg0: !transform.any_op {transform.readonly}) {
%func = transform.structured.match ops{["func.func"]} in %arg0 : (!transform.any_op) -> !transform.any_op
transform.iree.eliminate_empty_tensors %func : (!transform.any_op) -> ()
- %_ = transform.iree.bufferize { target_gpu } %func: (!transform.any_op) -> !transform.any_op
+ %_ = transform.iree.bufferize <{ target_gpu }> %func: (!transform.any_op) -> !transform.any_op
transform.yield
} // @__transform_main
} // module
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_distribution_spec.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_distribution_spec.mlir
index a24fecf..f67d7eb 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_distribution_spec.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_distribution_spec.mlir
@@ -3,9 +3,9 @@
%if_op = transform.structured.match ops{["scf.if"]} in %variant_op
: (!transform.any_op) -> !transform.any_op
- %warp = transform.iree.vector.to_warp_execute_on_lane_0 %if_op { warp_size = 32 }
+ %warp = transform.iree.vector.to_warp_execute_on_lane_0 %if_op <{ warp_size = 32 }>
: (!transform.any_op) -> !transform.any_op
- %isolated = transform.get_parent_op %warp {isolated_from_above}
+ %isolated = transform.get_parent_op %warp <{isolated_from_above}>
: (!transform.any_op) -> !transform.any_op
transform.iree.vector.warp_distribute %isolated
: (!transform.any_op) -> ()
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir
index ddf7542..1b48151 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir
@@ -3,7 +3,7 @@
%if_op = transform.structured.match ops{["scf.if"]} in %variant_op
: (!transform.any_op) -> !transform.any_op
- transform.iree.vector.to_warp_execute_on_lane_0 %if_op { warp_size = 32 }
+ transform.iree.vector.to_warp_execute_on_lane_0 %if_op <{ warp_size = 32 }>
: (!transform.any_op) -> !transform.any_op
// Late canonicalizations to cleanup and pass the checks.
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_to_nvgpu_mma.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_to_nvgpu_mma.mlir
index 9cf6560..d144dcd 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_to_nvgpu_mma.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_to_nvgpu_mma.mlir
@@ -35,7 +35,7 @@
transform.apply_patterns to %func {
transform.apply_patterns.iree.unroll_vectors_gpu_wmma_sync
} : !transform.any_op
- transform.iree.vector.vector_to_mma_conversion %func { use_wmma } : (!transform.any_op) -> ()
+ transform.iree.vector.vector_to_mma_conversion %func <{ use_wmma }> : (!transform.any_op) -> ()
// Apply canonicalization post-hoc to trigger DCE and pass the test
// (i.e. all vector.contract are dead).
@@ -83,7 +83,7 @@
transform.apply_patterns to %func {
transform.apply_patterns.iree.unroll_vectors_gpu_mma_sync
} : !transform.any_op
- transform.iree.vector.vector_to_mma_conversion %func { use_mma_sync } : (!transform.any_op) -> ()
+ transform.iree.vector.vector_to_mma_conversion %func <{ use_mma_sync }> : (!transform.any_op) -> ()
// Apply canonicalization post-hoc to trigger DCE and pass the test
// (i.e. all vector.contract are dead).
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_gpu_pipelining.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_gpu_pipelining.mlir
index 6a4ea9e..a6a013b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_gpu_pipelining.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_gpu_pipelining.mlir
@@ -60,7 +60,7 @@
transform.named_sequence @__transform_main(%root: !transform.any_op {transform.readonly}) {
%for = transform.structured.match ops{["scf.for"]} in %root : (!transform.any_op) -> !transform.any_op
%1 = transform.cast %for : !transform.any_op to !transform.op<"scf.for">
- %2 = transform.iree.pipeline_shared_memory_copies %1 { depth = 4 } : (!transform.op<"scf.for">) -> !transform.op<"scf.for">
+ %2 = transform.iree.pipeline_shared_memory_copies %1 <{ depth = 4 }> : (!transform.op<"scf.for">) -> !transform.op<"scf.for">
transform.yield
} // @__transform_main
} // module
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_vector_to_mma.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_vector_to_mma.mlir
index 1b6f81c..68f7b6b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_vector_to_mma.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_vector_to_mma.mlir
@@ -56,7 +56,7 @@
transform.apply_patterns to %func {
transform.apply_patterns.iree.unroll_vectors_gpu_wmma_sync
} : !transform.any_op
- transform.iree.vector.vector_to_mma_conversion %func { use_wmma } : (!transform.any_op) -> ()
+ transform.iree.vector.vector_to_mma_conversion %func <{ use_wmma }> : (!transform.any_op) -> ()
// Apply canonicalization post-hoc to trigger DCE and pass the test
// (i.e. all vector.contract are dead).
@@ -140,7 +140,7 @@
transform.apply_patterns to %func {
transform.apply_patterns.iree.unroll_vectors_gpu_wmma_sync
} : !transform.any_op
- transform.iree.vector.vector_to_mma_conversion %func { use_wmma } : (!transform.any_op) -> ()
+ transform.iree.vector.vector_to_mma_conversion %func <{ use_wmma }> : (!transform.any_op) -> ()
transform.apply_patterns to %func {
transform.apply_patterns.canonicalization
} : !transform.any_op
diff --git a/compiler/src/iree/compiler/Dialect/Encoding/IR/EncodingAttrs.td b/compiler/src/iree/compiler/Dialect/Encoding/IR/EncodingAttrs.td
index 875855a..d6ef389 100644
--- a/compiler/src/iree/compiler/Dialect/Encoding/IR/EncodingAttrs.td
+++ b/compiler/src/iree/compiler/Dialect/Encoding/IR/EncodingAttrs.td
@@ -24,7 +24,9 @@
}
class IREEEncoding_EnumAttr<EnumAttrInfo enumInfo, string name = "">
- : EnumAttr<IREEEncoding_Dialect, enumInfo, name>;
+ : EnumAttr<IREEEncoding_Dialect, enumInfo, name> {
+ let assemblyFormat = "$value";
+}
// Enums for tagging operand operation in an EncodingAttr
def MATMUL : I32EnumAttrCase<"matmul", 0>;
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/AnnotateDispatches.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/AnnotateDispatches.cpp
index c377b47..7d6dd50 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/AnnotateDispatches.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/AnnotateDispatches.cpp
@@ -519,7 +519,7 @@
exportOp.setSymName(newName);
exportOp.setFunctionRef(newName);
- funcOp.setName(newName);
+ funcOp.setSymbolName(newName);
auto newSymbolRefAttr =
SymbolRefAttr::get(&getContext(), executableOp.getName(),
diff --git a/compiler/src/iree/compiler/Dialect/LinalgExt/TransformExtensions/LinalgExtExtensionsOps.td b/compiler/src/iree/compiler/Dialect/LinalgExt/TransformExtensions/LinalgExtExtensionsOps.td
index c89de92..6662621 100644
--- a/compiler/src/iree/compiler/Dialect/LinalgExt/TransformExtensions/LinalgExtExtensionsOps.td
+++ b/compiler/src/iree/compiler/Dialect/LinalgExt/TransformExtensions/LinalgExtExtensionsOps.td
@@ -29,11 +29,10 @@
);
let results = (outs Variadic<TransformHandleTypeInterface>:$result);
- let assemblyFormat = "attr-dict $target `:` functional-type(operands, results)";
let cppNamespace = "mlir::iree_compiler::IREE::LinalgExt";
let assemblyFormat = [{
- $target attr-dict `:` functional-type(operands, results)
+ $target prop-dict attr-dict `:` functional-type(operands, results)
}];
let extraClassDeclaration = [{
diff --git a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/test/distribution.mlir b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/test/distribution.mlir
index c348d52..c09cc16 100644
--- a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/test/distribution.mlir
+++ b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/test/distribution.mlir
@@ -17,7 +17,7 @@
module attributes { transform.with_named_sequence } {
transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
%0 = transform.structured.match ops{["iree_linalg_ext.scatter"]} in %module_op : (!transform.any_op) -> !transform.any_op
- %forall, %tiled_op = transform.structured.tile_using_forall %0 tile_sizes [10, 30, 0] { mapping = [#gpu.thread<y>, #gpu.thread<x>] } : (!transform.any_op) -> (!transform.any_op, !transform.any_op)
+ %forall, %tiled_op = transform.structured.tile_using_forall %0 tile_sizes [10, 30, 0] ( mapping = [#gpu.thread<y>, #gpu.thread<x>] ) : (!transform.any_op) -> (!transform.any_op, !transform.any_op)
transform.yield
}
}
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/LinkModules.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/LinkModules.cpp
index 75c9f75..83f13f2 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Transforms/LinkModules.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/LinkModules.cpp
@@ -695,9 +695,9 @@
// Rename the operation.
if (auto funcOp = dyn_cast<FunctionOpInterface>(info->op)) {
- funcOp.setName(rename.newName);
+ funcOp.setSymbolName(rename.newName);
} else if (auto symOp = dyn_cast<SymbolOpInterface>(info->op)) {
- symOp.setName(rename.newName);
+ symOp.setSymbolName(rename.newName);
}
// Update SymbolInfo's names to reflect the rename.
diff --git a/samples/transform_dialect/transform_library.mlir b/samples/transform_dialect/transform_library.mlir
index fda287d..6e6f52c 100644
--- a/samples/transform_dialect/transform_library.mlir
+++ b/samples/transform_dialect/transform_library.mlir
@@ -40,7 +40,7 @@
transform.apply_patterns to %func_1 {
transform.apply_patterns.linalg.erase_unnecessary_inputs
} : !transform.any_op
- %memref_func = transform.iree.bufferize { target_gpu } %func_1 : (!transform.any_op) -> (!transform.any_op)
+ %memref_func = transform.iree.bufferize <{ target_gpu }> %func_1 : (!transform.any_op) -> (!transform.any_op)
// Step 6. Post-bufferization vector distribution
// ===========================================================================
@@ -48,30 +48,30 @@
transform.iree.forall_to_workgroup %func_7 : (!transform.any_op) -> ()
transform.iree.map_nested_forall_to_gpu_threads %func_7
workgroup_dims = [4, 8, 1] : (!transform.any_op) -> ()
- transform.print {name = "Ran custom_transform_strategy"}
+ transform.print name = "Ran custom_transform_strategy"
transform.yield
}
// Send it down a custom transform dialect pipeline.
transform.named_sequence @custom_matmul(%matmul: !transform.any_op {transform.readonly}) {
- %variant_op = transform.get_parent_op %matmul {op_name = "hal.executable.variant"} : (!transform.any_op) -> !transform.any_op
+ %variant_op = transform.get_parent_op %matmul <{op_name = "hal.executable.variant"}> : (!transform.any_op) -> !transform.any_op
%funcs = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
%subgroup_reduce = transform.param.constant #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen
codegen_spec = @custom_transform_strategy> -> !transform.any_param
transform.annotate %funcs "translation_info" = %subgroup_reduce : !transform.any_op, !transform.any_param
- transform.print {name = "Setting matmul strategy to custom_transform_strategy"}
+ transform.print name = "Setting matmul strategy to custom_transform_strategy"
transform.yield
}
// Send it down subgroup reduce with a custom tiling configuration.
transform.named_sequence @use_base_vectorize(%reduce: !transform.any_op {transform.readonly}) {
- %variant_op = transform.get_parent_op %reduce {op_name = "hal.executable.variant"} : (!transform.any_op) -> !transform.any_op
+ %variant_op = transform.get_parent_op %reduce <{op_name = "hal.executable.variant"}> : (!transform.any_op) -> !transform.any_op
%lowering_config = transform.param.constant #iree_codegen.lowering_config<tile_sizes = [[8, 0], [1, 0], [0, 0, 4]]> -> !transform.any_param
transform.annotate %reduce "lowering_config" = %lowering_config : !transform.any_op, !transform.any_param
%funcs = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
%subgroup_reduce = transform.param.constant #iree_codegen.translation_info<pipeline = #iree_gpu.spirv_pipeline<BaseVectorize> workgroup_size = [16, 1, 1]> -> !transform.any_param
transform.annotate %funcs "translation_info" = %subgroup_reduce : !transform.any_op, !transform.any_param
- transform.print {name = "Setting reduce strategy to base vectorize"}
+ transform.print name = "Setting reduce strategy to base vectorize"
transform.yield
}
@@ -105,7 +105,7 @@
%c2 = transform.param.constant 2 : i64 -> !transform.param<i64>
%rank = transform.match.structured.rank %arg1 : (!transform.any_op) -> !transform.param<i64>
transform.match.param.cmpi eq %rank, %c2 : !transform.param<i64>
- transform.match.structured.dim %arg1[-1] {reduction} : !transform.any_op
+ transform.match.structured.dim %arg1[-1] reduction : !transform.any_op
transform.match.structured.yield %arg1 : !transform.any_op
}
transform.yield %matched : !transform.any_op
diff --git a/third_party/llvm-project b/third_party/llvm-project
index d1a9d89..6cce5bc 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit d1a9d89827aaf5badac754cdae6ac9d47b9c306f
+Subproject commit 6cce5bca3b901ec32e250fb43e88eb0fe511af93
diff --git a/third_party/stablehlo b/third_party/stablehlo
index a2727e5..50a9d75 160000
--- a/third_party/stablehlo
+++ b/third_party/stablehlo
@@ -1 +1 @@
-Subproject commit a2727e576ec6f18acc561333dfd70ef49d5f1cdf
+Subproject commit 50a9d75a404653fe5bbf284b5af8b3ffeda42ee0
