[flang-commits] [flang] [mlir] [MLIR][ODS] Default EnumAttr to angle-bracket assembly (PR #220608)
Mehdi Amini via flang-commits
flang-commits at lists.llvm.org
Thu Sep 3 06:11:08 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/220608
>From 33b7031db6c78efef07dfa68ce762d098bafff0f Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Wed, 2 Sep 2026 07:40:56 -0700
Subject: [PATCH] [MLIR][ODS] Default EnumAttr to angle-bracket assembly
Make angle brackets the default assembly format for EnumAttr.
Use explicit enum directives where operations retain bare symbolic syntax.
Update affected MLIR and Flang assembly.
Assisted-by: Codex
---
.../Dialect/CUF/Attributes/CUFAttr.td | 3 -
.../flang/Optimizer/Dialect/FIRAttr.td | 4 -
.../include/flang/Optimizer/Dialect/FIROps.td | 3 +-
flang/test/Lower/OpenACC/acc-cache.f90 | 6 +-
.../OpenMP/Todo/requires-unnamed-common.f90 | 2 +-
flang/test/Lower/OpenMP/force-usm.f90 | 2 +-
.../Lower/OpenMP/has_device_addr-mapinfo.f90 | 4 +-
flang/test/Lower/OpenMP/interop.f90 | 18 +-
flang/test/Lower/OpenMP/requires-common.f90 | 2 +-
.../test/Lower/OpenMP/requires-force-usm.f90 | 2 +-
flang/test/Lower/OpenMP/requires-usm.f90 | 2 +-
flang/test/Lower/OpenMP/requires.f90 | 2 +-
flang/test/Lower/location.f90 | 3 +-
.../lower-workdistribute-fission-target.mlir | 2 +-
.../omp-map-info-finalization-usm.fir | 2 +-
.../mlir/Dialect/AMDGPU/IR/AMDGPUAttrs.td | 1 -
.../mlir/Dialect/AMDGPU/IR/AMDGPUOps.td | 9 +-
.../mlir/Dialect/Arith/IR/ArithBase.td | 2 -
.../mlir/Dialect/ArmSME/IR/ArmSMEOps.td | 3 -
mlir/include/mlir/Dialect/GPU/IR/GPUBase.td | 2 -
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 41 +++--
.../include/mlir/Dialect/Index/IR/IndexOps.td | 2 +-
mlir/include/mlir/Dialect/LLVMIR/LLVMEnums.td | 2 -
mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td | 2 -
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 168 ++++++++----------
.../include/mlir/Dialect/LLVMIR/ROCDLAttrs.td | 1 -
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 35 ++--
mlir/include/mlir/Dialect/LLVMIR/XeVMOps.td | 4 -
.../mlir/Dialect/Linalg/IR/LinalgBase.td | 6 -
mlir/include/mlir/Dialect/MPI/IR/MPI.td | 1 -
.../mlir/Dialect/OpenACC/OpenACCOps.td | 30 ++--
.../mlir/Dialect/OpenMP/OpenMPClauses.td | 2 +-
.../mlir/Dialect/OpenMP/OpenMPEnums.td | 13 --
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td | 10 +-
.../mlir/Dialect/SPIRV/IR/SPIRVBase.td | 2 -
.../mlir/Dialect/Shard/IR/ShardBase.td | 1 -
.../include/mlir/Dialect/Shard/IR/ShardOps.td | 7 +-
.../SparseTensor/IR/SparseTensorOps.td | 11 +-
.../mlir/Dialect/Tosa/IR/TosaOpBase.td | 1 -
.../Dialect/Vector/IR/VectorAttributes.td | 3 -
.../mlir/Dialect/XeGPU/IR/XeGPUAttrs.td | 45 +++--
.../include/mlir/Dialect/XeGPU/IR/XeGPUOps.td | 7 +-
mlir/include/mlir/IR/EnumAttr.td | 22 ++-
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 20 ---
.../OpenACC/Transforms/ACCImplicitData.cpp | 14 +-
mlir/python/mlir/dialects/irdl.py | 3 +-
mlir/test/CAPI/irdl.c | 8 +-
mlir/test/Dialect/GPU/invalid.mlir | 6 +-
.../Dialect/IRDL/variadics-error.irdl.mlir | 8 +-
mlir/test/Dialect/OpenACC/invalid.mlir | 4 +-
mlir/test/Dialect/OpenACC/ops.mlir | 28 +--
mlir/test/Dialect/OpenMP/invalid.mlir | 12 +-
mlir/test/Dialect/OpenMP/ops.mlir | 36 ++--
mlir/test/IR/array-of-attr.mlir | 4 +-
mlir/test/IR/enum-attr-invalid.mlir | 2 +-
mlir/test/IR/enum-attr-roundtrip.mlir | 18 +-
mlir/test/IR/properties.mlir | 8 +-
...target-atomic-capture-control-options.mlir | 2 +-
...ptarget-atomic-update-control-options.mlir | 2 +-
...mptarget-declare-target-llvm-device-2.mlir | 2 +-
...-declare-target-module-rewrite-device.mlir | 2 +-
.../LLVMIR/omptarget-host-ref-semantics.mlir | 2 +-
.../omptarget-map-pointer-privatization.mlir | 2 +-
.../LLVMIR/omptarget-parallel-llvm-debug.mlir | 2 +-
.../Target/LLVMIR/openmp-interop-depend.mlir | 2 +-
mlir/test/Target/LLVMIR/openmp-interop.mlir | 12 +-
mlir/test/Target/LLVMIR/openmp-llvm.mlir | 2 +-
.../openmp-reduction-task-modifier.mlir | 6 +-
.../openmp-teams-clauses-trunc-ext.mlir | 2 +-
mlir/test/lib/Dialect/Test/TestAttrDefs.td | 23 ++-
mlir/test/lib/Dialect/Test/TestOps.td | 5 +-
.../attr-or-type-format-roundtrip.mlir | 12 +-
mlir/test/mlir-tblgen/attr-or-type-format.td | 2 -
.../test/mlir-tblgen/enums-python-bindings.td | 4 +-
mlir/test/mlir-tblgen/op-format.td | 6 +-
mlir/tools/mlir-tblgen/OpFormatGen.cpp | 24 +--
76 files changed, 339 insertions(+), 436 deletions(-)
diff --git a/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td b/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td
index af21fe3d05fa4..902c8ec41a997 100644
--- a/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td
+++ b/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td
@@ -36,7 +36,6 @@ def cuf_DataAttribute : I32EnumAttr<
def cuf_DataAttributeAttr :
EnumAttr<CUFDialect, cuf_DataAttribute, "cuda"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def cuf_ProcAttribute : I32EnumAttr<
@@ -54,7 +53,6 @@ def cuf_ProcAttribute : I32EnumAttr<
def cuf_ProcAttributeAttr :
EnumAttr<CUFDialect, cuf_ProcAttribute, "cuda_proc"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def cuf_LaunchBoundsAttr : cuf_Attr<"LaunchBounds"> {
@@ -94,7 +92,6 @@ def cuf_DataTransferKind : I32EnumAttr<
def cuf_DataTransferKindAttr :
EnumAttr<CUFDialect, cuf_DataTransferKind, "cuda_transfer"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
#endif // FORTRAN_DIALECT_CUF_CUFATTRS
diff --git a/flang/include/flang/Optimizer/Dialect/FIRAttr.td b/flang/include/flang/Optimizer/Dialect/FIRAttr.td
index cfe8e96e8e4a5..b1d4223f520ff 100644
--- a/flang/include/flang/Optimizer/Dialect/FIRAttr.td
+++ b/flang/include/flang/Optimizer/Dialect/FIRAttr.td
@@ -146,7 +146,6 @@ def fir_FortranProcedureFlagsEnum
def fir_FortranProcedureFlagsAttr :
EnumAttr<FIROpsDialect, fir_FortranProcedureFlagsEnum, "proc_attrs"> {
- let assemblyFormat = "`<` $value `>`";
}
def fir_BoxFieldAttr : I32EnumAttr<
@@ -234,7 +233,6 @@ def fir_PackArrayHeuristics
def fir_PackArrayHeuristicsAttr
: EnumAttr<FIROpsDialect, fir_PackArrayHeuristics,
"pack_array_heuristics"> {
- let assemblyFormat = "`<` $value `>`";
}
def fir_OpenACCSafeTempArrayCopyAttr : fir_Attr<"OpenACCSafeTempArrayCopy"> {
@@ -275,7 +273,6 @@ def LocalitySpecifierType : I32EnumAttr<
def LocalitySpecifierTypeAttr : EnumAttr<FIROpsDialect, LocalitySpecifierType,
"locality_specifier_type"> {
- let assemblyFormat = "`{` `type` `=` $value `}`";
}
/// Fortran inline attribute
@@ -296,7 +293,6 @@ def fir_FortranInlineEnum
def fir_FortranInlineAttr
: EnumAttr<FIROpsDialect, fir_FortranInlineEnum, "inline_attrs"> {
- let assemblyFormat = "`<` $value `>`";
}
// USE statement rename mapping: local_name => use_name
diff --git a/flang/include/flang/Optimizer/Dialect/FIROps.td b/flang/include/flang/Optimizer/Dialect/FIROps.td
index 1a5c29a6e991e..57e23ecc01af8 100644
--- a/flang/include/flang/Optimizer/Dialect/FIROps.td
+++ b/flang/include/flang/Optimizer/Dialect/FIROps.td
@@ -3769,7 +3769,8 @@ def fir_LocalitySpecifierOp : fir_Op<"local", [IsolatedFromAbove,
AnyRegion:$dealloc_region);
let assemblyFormat = [{
- $locality_specifier_type ($sym_visibility^)? $sym_name `:` $type
+ ` ` `{` `type` `=` enum($locality_specifier_type) `}`
+ ($sym_visibility^)? $sym_name `:` $type
(`init` $init_region^)?
(`copy` $copy_region^)?
(`dealloc` $dealloc_region^)?
diff --git a/flang/test/Lower/OpenACC/acc-cache.f90 b/flang/test/Lower/OpenACC/acc-cache.f90
index 2205936b457ce..5eba99f2881e9 100644
--- a/flang/test/Lower/OpenACC/acc-cache.f90
+++ b/flang/test/Lower/OpenACC/acc-cache.f90
@@ -41,7 +41,7 @@ subroutine test_cache_readonly()
end do
! CHECK: acc.loop
-! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%{{.*}} : !fir.ref<!fir.array<10xf32>>) structured(false) name("b") <modifiers = readonly> -> !fir.ref<!fir.array<10xf32>>
+! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%{{.*}} : !fir.ref<!fir.array<10xf32>>) structured(false) name("b") <modifiers = <readonly>> -> !fir.ref<!fir.array<10xf32>>
! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[CACHE]](%{{.*}}) {uniq_name = "_QFtest_cache_readonlyEb"}
! Loop body uses the cached readonly reference
! CHECK: %[[ELEM:.*]] = hlfir.designate %[[DECL]]#0 (%{{.*}}) : (!fir.ref<!fir.array<10xf32>>, i64) -> !fir.ref<f32>
@@ -605,7 +605,7 @@ subroutine test_cache_derived_type_readonly()
! CHECK: acc.loop
! CHECK: %[[ARRAY_COORD:.*]] = hlfir.designate %{{.*}}{"array"} shape %{{.*}} : (!fir.ref<!fir.type<_QFtest_cache_derived_type_readonlyTdt{array:!fir.array<100xf32>}>>, !fir.shape<1>) -> !fir.ref<!fir.array<100xf32>>
! CHECK: %[[BOUND:.*]] = acc.bounds lowerbound(%{{.*}} : index) upperbound(%{{.*}} : index) extent(%{{.*}} : index) stride(%{{.*}} : index) startIdx(%{{.*}} : index)
-! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%[[ARRAY_COORD]] : !fir.ref<!fir.array<100xf32>>) bounds(%[[BOUND]]) structured(false) name("data%array(i-4_4:i+4_4)") <modifiers = readonly> -> !fir.ref<!fir.array<100xf32>>
+! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%[[ARRAY_COORD]] : !fir.ref<!fir.array<100xf32>>) bounds(%[[BOUND]]) structured(false) name("data%array(i-4_4:i+4_4)") <modifiers = <readonly>> -> !fir.ref<!fir.array<100xf32>>
! CHECK: acc.yield
end subroutine
@@ -761,7 +761,7 @@ subroutine test_cache_temp_in_designator(data, a)
! CHECK: %[[ELEMENTAL:.*]] = hlfir.elemental
! CHECK: %[[MAXLOC:.*]] = hlfir.maxloc %[[ELEMENTAL]]
! CHECK: %[[BOUND:.*]] = acc.bounds lowerbound({{.*}}) upperbound({{.*}})
-! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%{{.*}}) bounds(%[[BOUND]]) structured(false) name("data(1:maxloc(a+a,dim=1_4))") <modifiers = readonly> -> !fir.ref<!fir.array<100xf32>>
+! CHECK: %[[CACHE:.*]] = acc.cache varPtr(%{{.*}}) bounds(%[[BOUND]]) structured(false) name("data(1:maxloc(a+a,dim=1_4))") <modifiers = <readonly>> -> !fir.ref<!fir.array<100xf32>>
! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[CACHE]]
! CHECK: hlfir.destroy %[[ELEMENTAL]]
! CHECK: hlfir.designate %[[DECL]]#0
diff --git a/flang/test/Lower/OpenMP/Todo/requires-unnamed-common.f90 b/flang/test/Lower/OpenMP/Todo/requires-unnamed-common.f90
index 578f71a55e43b..6216da04ec87d 100644
--- a/flang/test/Lower/OpenMP/Todo/requires-unnamed-common.f90
+++ b/flang/test/Lower/OpenMP/Todo/requires-unnamed-common.f90
@@ -10,7 +10,7 @@
! XFAIL: *
!CHECK: module attributes {
-!CHECK-SAME: omp.requires = #omp<clause_requires unified_shared_memory>
+!CHECK-SAME: omp.requires = #omp.clause_requires<unified_shared_memory>
block data
!$omp requires unified_shared_memory
integer :: x
diff --git a/flang/test/Lower/OpenMP/force-usm.f90 b/flang/test/Lower/OpenMP/force-usm.f90
index 90bbf3c4d842f..231d28c2cb801 100644
--- a/flang/test/Lower/OpenMP/force-usm.f90
+++ b/flang/test/Lower/OpenMP/force-usm.f90
@@ -7,6 +7,6 @@
! -fopenmp-force-usm is set
!CHECK: module attributes {
-!CHECK-SAME: omp.requires = #omp<clause_requires unified_shared_memory>
+!CHECK-SAME: omp.requires = #omp.clause_requires<unified_shared_memory>
program requires
end program requires
diff --git a/flang/test/Lower/OpenMP/has_device_addr-mapinfo.f90 b/flang/test/Lower/OpenMP/has_device_addr-mapinfo.f90
index 9e0c7a5dd84d6..e85cbba932a49 100644
--- a/flang/test/Lower/OpenMP/has_device_addr-mapinfo.f90
+++ b/flang/test/Lower/OpenMP/has_device_addr-mapinfo.f90
@@ -17,7 +17,7 @@ integer function s(a)
! Check that the map.info for `a` only takes a single parameter.
-!CHECK-DAG: %[[MAP_A:[0-9]+]] = "omp.map.info"(%[[STORAGE_A:[0-9#]+]]) <{map_capture_type = #omp<variable_capture_kind(ByRef)>, map_type = #omp<clause_map_flags to|always|implicit>, name = "a", operandSegmentSizes = array<i32: 1, 0, 0, 0>, partial_map = false, var_ptr_type = !fir.box<!fir.array<?xi32>>}> : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.ref<!fir.array<?xi32>>
-!CHECK-DAG: %[[MAP_T:[0-9]+]] = "omp.map.info"(%[[STORAGE_T:[0-9#]+]]) <{map_capture_type = #omp<variable_capture_kind(ByRef)>, map_type = #omp<clause_map_flags from>, name = "t", operandSegmentSizes = array<i32: 1, 0, 0, 0>, partial_map = false, var_ptr_type = i32}> : (!fir.ref<i32>) -> !fir.ref<i32>
+!CHECK-DAG: %[[MAP_A:[0-9]+]] = "omp.map.info"(%[[STORAGE_A:[0-9#]+]]) <{map_capture_type = #omp.variable_capture_kind<ByRef>, map_type = #omp.clause_map_flags<to|always|implicit>, name = "a", operandSegmentSizes = array<i32: 1, 0, 0, 0>, partial_map = false, var_ptr_type = !fir.box<!fir.array<?xi32>>}> : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.ref<!fir.array<?xi32>>
+!CHECK-DAG: %[[MAP_T:[0-9]+]] = "omp.map.info"(%[[STORAGE_T:[0-9#]+]]) <{map_capture_type = #omp.variable_capture_kind<ByRef>, map_type = #omp.clause_map_flags<from>, name = "t", operandSegmentSizes = array<i32: 1, 0, 0, 0>, partial_map = false, var_ptr_type = i32}> : (!fir.ref<i32>) -> !fir.ref<i32>
!CHECK: "omp.target"(%[[MAP_A]], %[[MAP_T]])
diff --git a/flang/test/Lower/OpenMP/interop.f90 b/flang/test/Lower/OpenMP/interop.f90
index 1fd9b353b535f..219a7ae89c5dd 100644
--- a/flang/test/Lower/OpenMP/interop.f90
+++ b/flang/test/Lower/OpenMP/interop.f90
@@ -7,7 +7,7 @@
!CHECK-LABEL: func.func @_QPtest_interop_init(
!CHECK-SAME: %[[ARG:.*]]: !fir.ref<i64> {fir.bindc_name = "obj"})
!CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG]]
-!CHECK: omp.interop.init %[[DECL]]#0 : !fir.ref<i64> interop_types([#omp<interop_type(target)>])
+!CHECK: omp.interop.init %[[DECL]]#0 : !fir.ref<i64> interop_types([#omp.interop_type<target>])
subroutine test_interop_init(obj)
integer(8) :: obj
!$omp interop init(target: obj)
@@ -18,7 +18,7 @@ subroutine test_interop_init(obj)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_init_targetsync(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(targetsync)>])
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<targetsync>])
subroutine test_interop_init_targetsync(obj)
integer(8) :: obj
!$omp interop init(targetsync: obj)
@@ -29,7 +29,7 @@ subroutine test_interop_init_targetsync(obj)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_init_both(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(targetsync)>, #omp<interop_type(target)>])
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<targetsync>, #omp.interop_type<target>])
subroutine test_interop_init_both(obj)
integer(8) :: obj
!$omp interop init(targetsync, target: obj)
@@ -51,7 +51,7 @@ subroutine test_interop_init_untyped(obj)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_init_nowait(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(target)>]) nowait
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<target>]) nowait
subroutine test_interop_init_nowait(obj)
integer(8) :: obj
!$omp interop init(target: obj) nowait
@@ -63,7 +63,7 @@ subroutine test_interop_init_nowait(obj)
!CHECK-LABEL: func.func @_QPtest_interop_device(
!CHECK: %[[DEV:.*]] = fir.load %{{.*}} : !fir.ref<i32>
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(target)>]) device(%[[DEV]] : i32)
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<target>]) device(%[[DEV]] : i32)
subroutine test_interop_device(obj, dev)
integer(8) :: obj
integer :: dev
@@ -75,7 +75,7 @@ subroutine test_interop_device(obj, dev)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_prefer_str(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(targetsync)>]) prefer_type([1, 6])
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<targetsync>]) prefer_type([1, 6])
subroutine test_interop_prefer_str(obj)
integer(8) :: obj
!$omp interop init(prefer_type("cuda", "level_zero"), targetsync: obj)
@@ -86,7 +86,7 @@ subroutine test_interop_prefer_str(obj)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_prefer_hip(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(target)>]) prefer_type([5])
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<target>]) prefer_type([5])
subroutine test_interop_prefer_hip(obj)
integer(8) :: obj
!$omp interop init(prefer_type("hip"), target: obj)
@@ -97,7 +97,7 @@ subroutine test_interop_prefer_hip(obj)
!===============================================================================
!CHECK-LABEL: func.func @_QPtest_interop_prefer_int(
-!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp<interop_type(targetsync)>]) prefer_type([4])
+!CHECK: omp.interop.init %{{.*}} : !fir.ref<i64> interop_types([#omp.interop_type<targetsync>]) prefer_type([4])
subroutine test_interop_prefer_int(obj)
integer(8) :: obj
integer, parameter :: omp_ifr_sycl = 4
@@ -180,7 +180,7 @@ subroutine test_interop_destroy_device(obj, dev)
!CHECK-LABEL: func.func @_QPtest_interop_init_array_element(
!CHECK: %[[EL:.*]] = hlfir.designate %{{.*}} (%{{.*}}) : (!fir.ref<!fir.array<10xi64>>, index) -> !fir.ref<i64>
-!CHECK: omp.interop.init %[[EL]] : !fir.ref<i64> interop_types([#omp<interop_type(target)>])
+!CHECK: omp.interop.init %[[EL]] : !fir.ref<i64> interop_types([#omp.interop_type<target>])
subroutine test_interop_init_array_element(arr)
integer(8) :: arr(10)
!$omp interop init(target: arr(1))
diff --git a/flang/test/Lower/OpenMP/requires-common.f90 b/flang/test/Lower/OpenMP/requires-common.f90
index b3801a834014f..bfaa2d9de26b6 100644
--- a/flang/test/Lower/OpenMP/requires-common.f90
+++ b/flang/test/Lower/OpenMP/requires-common.f90
@@ -6,7 +6,7 @@
! This test checks the lowering of requires into MLIR
!CHECK: module attributes {
-!CHECK-SAME: omp.requires = #omp<clause_requires unified_shared_memory>
+!CHECK-SAME: omp.requires = #omp.clause_requires<unified_shared_memory>
block data init
!$omp requires unified_shared_memory
integer :: x
diff --git a/flang/test/Lower/OpenMP/requires-force-usm.f90 b/flang/test/Lower/OpenMP/requires-force-usm.f90
index 5f5cf9e64cd70..52ac1bae35162 100644
--- a/flang/test/Lower/OpenMP/requires-force-usm.f90
+++ b/flang/test/Lower/OpenMP/requires-force-usm.f90
@@ -7,7 +7,7 @@
! -fopenmp-force-usm is set, even when other requires directives are present
!CHECK: module attributes {
-!CHECK-SAME: omp.requires = #omp<clause_requires reverse_offload|unified_shared_memory>
+!CHECK-SAME: omp.requires = #omp.clause_requires<reverse_offload|unified_shared_memory>
program requires
!$omp requires reverse_offload
!$omp target
diff --git a/flang/test/Lower/OpenMP/requires-usm.f90 b/flang/test/Lower/OpenMP/requires-usm.f90
index eb5b84cdba78c..9cfaee1686aa6 100644
--- a/flang/test/Lower/OpenMP/requires-usm.f90
+++ b/flang/test/Lower/OpenMP/requires-usm.f90
@@ -17,7 +17,7 @@
! RUN: bbc -fopenmp -fopenmp-is-target-device -emit-hlfir %t/requires-usm-program-after.f90 -o - | FileCheck %s
! CHECK: module attributes {
-! CHECK-SAME: omp.requires = #omp<clause_requires unified_shared_memory>
+! CHECK-SAME: omp.requires = #omp.clause_requires<unified_shared_memory>
!--- requires-usm.f90
module declare_mod
diff --git a/flang/test/Lower/OpenMP/requires.f90 b/flang/test/Lower/OpenMP/requires.f90
index cec836f11a3a0..258ee08962974 100644
--- a/flang/test/Lower/OpenMP/requires.f90
+++ b/flang/test/Lower/OpenMP/requires.f90
@@ -6,7 +6,7 @@
! This test checks the lowering of requires into MLIR
!CHECK: module attributes {
-!CHECK-SAME: omp.requires = #omp<clause_requires reverse_offload|unified_shared_memory>
+!CHECK-SAME: omp.requires = #omp.clause_requires<reverse_offload|unified_shared_memory>
program requires
!$omp requires unified_shared_memory reverse_offload atomic_default_mem_order(seq_cst)
!$omp target
diff --git a/flang/test/Lower/location.f90 b/flang/test/Lower/location.f90
index 744ff7e3bf039..7296ed7e8d828 100644
--- a/flang/test/Lower/location.f90
+++ b/flang/test/Lower/location.f90
@@ -6,8 +6,7 @@ program test
end
! CHECK-LABEL: func.func @_QQmain() attributes {fir.bindc_name = "TEST"} {
-! CHECK: fir.call @_FortranAioOutputAscii(%{{.*}}, %{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<i8>, !fir.ref<i8>, i64) -> i1 loc(fused<#fir<loc_kind_array[base, inclusion, inclusion]>>["{{.*}}location1.inc":1:10, "{{.*}}location0.inc":1:1, "{{.*}}location.f90":4:1])
+! CHECK: fir.call @_FortranAioOutputAscii(%{{.*}}, %{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<i8>, !fir.ref<i8>, i64) -> i1 loc(fused<#fir<loc_kind_array[<base>, <inclusion>, <inclusion>]>>["{{.*}}location1.inc":1:10, "{{.*}}location0.inc":1:1, "{{.*}}location.f90":4:1])
! CHECK: return loc("{{.*}}location.f90":6:1)
! CHECK: } loc("{{.*}}location.f90":3:1)
-
diff --git a/flang/test/Transforms/OpenMP/lower-workdistribute-fission-target.mlir b/flang/test/Transforms/OpenMP/lower-workdistribute-fission-target.mlir
index 224962b949713..beea10f340d04 100644
--- a/flang/test/Transforms/OpenMP/lower-workdistribute-fission-target.mlir
+++ b/flang/test/Transforms/OpenMP/lower-workdistribute-fission-target.mlir
@@ -97,7 +97,7 @@ func.func @x(%lb : index, %ub : index, %step : index, %addr : !fir.ref<index>) {
"omp.target"(%addr, %allocator, %lb_map, %ub_map, %step_map, %addr_map, %addr) <{
allocate_private_indices = array<i64: 0>,
allocate_alignments = array<i64: 64>,
- kernel_type = #omp<kernel_type(generic)>,
+ kernel_type = #omp.kernel_type<generic>,
operandSegmentSizes = array<i32: 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 4, 0, 1, 0>,
private_syms = [@addr_private]
}> ({
diff --git a/flang/test/Transforms/omp-map-info-finalization-usm.fir b/flang/test/Transforms/omp-map-info-finalization-usm.fir
index dd1d477323ca1..370b536af26a4 100644
--- a/flang/test/Transforms/omp-map-info-finalization-usm.fir
+++ b/flang/test/Transforms/omp-map-info-finalization-usm.fir
@@ -5,7 +5,7 @@
// Unified Shared Memory (USM) mode where the parent is in USM (no close) but
// members (like descriptors) might have been initially tagged with close.
-module attributes {omp.requires = #omp<clause_requires unified_shared_memory>} {
+module attributes {omp.requires = #omp.clause_requires<unified_shared_memory>} {
func.func @test_usm_close_flag_cleanup(%arg0: !fir.ref<!fir.type<t{a:!fir.box<!fir.heap<!fir.array<?xf32>>>}>>) {
%map = omp.map.info var_ptr(%arg0 : !fir.ref<!fir.type<t{a:!fir.box<!fir.heap<!fir.array<?xf32>>>}>>, !fir.type<t{a:!fir.box<!fir.heap<!fir.array<?xf32>>>}>) map_clauses(to) capture(ByRef) name("parent") -> !fir.ref<!fir.type<t{a:!fir.box<!fir.heap<!fir.array<?xf32>>>}>>
diff --git a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUAttrs.td b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUAttrs.td
index 65a5f044bf4d8..444ab34a079bc 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUAttrs.td
+++ b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUAttrs.td
@@ -34,7 +34,6 @@ def AMDGPU_AddressSpaceAttr : EnumAttr<AMDGPU_Dialect, AMDGPU_AddressSpace,
structured indexing that is primarily seen in graphics applications. This
is also incompatible with the simple indexing model supported by memref.
}];
- let assemblyFormat = "`<` $value `>`";
}
def AMDGPU_DPPPermAttr : EnumAttr<AMDGPU_Dialect, AMDGPU_DPPPerm,
diff --git a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUOps.td b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUOps.td
index 3a81f2c2126a3..398e754356ece 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUOps.td
+++ b/mlir/include/mlir/Dialect/AMDGPU/IR/AMDGPUOps.td
@@ -784,7 +784,7 @@ def AMDGPU_DPPOp : AMDGPU_Op<"dpp",
}];
let results = (outs AnyType:$result);
let assemblyFormat = [{
- $old $src $kind (`(` $permArgument^ `)`)?
+ $old $src enum($kind) (`(` $permArgument^ `)`)?
oilist (`row_mask` `(` $row_mask `)`
| `bank_mask` `(` $bank_mask `)`
| `bound_ctrl` `(` $bound_ctrl `)`)
@@ -944,7 +944,7 @@ def AMDGPU_SchedBarrierOp :
defined by the ROCDL scheduling group mask enum.
}];
let assemblyFormat = [{
- `allow` `=` $opts attr-dict
+ `allow` `=` enum($opts) attr-dict
}];
}
@@ -1101,7 +1101,7 @@ def AMDGPU_MFMAOp :
(`blocks` `(` $blocks^ `)`)?
custom<MNKDimensionList>($m, $n, $k) $sourceA `*` $sourceB `+` $destC
oilist (`abid` `(` $abid `)`
- | `blgp` `(` $blgp `)`
+ | `blgp` `(` enum($blgp) `)`
| `cbsz` `(` $cbsz `)`
| `reducePrecision` $reducePrecision
| `negateA` $negateA
@@ -2297,7 +2297,8 @@ def AMDGPU_GlobalPrefetchOp :
}];
let assemblyFormat = [{
- $src `[` $indices `]` $temporalHint $cacheScope (`speculative` $speculative^)? attr-dict `:` qualified(type($src))
+ $src `[` $indices `]` enum($temporalHint) enum($cacheScope)
+ (`speculative` $speculative^)? attr-dict `:` qualified(type($src))
}];
let hasVerifier = 1;
diff --git a/mlir/include/mlir/Dialect/Arith/IR/ArithBase.td b/mlir/include/mlir/Dialect/Arith/IR/ArithBase.td
index 760c502ccf1f5..df3d1f483b6fb 100644
--- a/mlir/include/mlir/Dialect/Arith/IR/ArithBase.td
+++ b/mlir/include/mlir/Dialect/Arith/IR/ArithBase.td
@@ -144,7 +144,6 @@ def FastMathFlags : I32BitEnumAttr<
def Arith_FastMathAttr :
EnumAttr<Arith_Dialect, FastMathFlags, "fastmath"> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -167,7 +166,6 @@ def Arith_IntegerOverflowFlags : I32BitEnumAttr<
def Arith_IntegerOverflowAttr :
EnumAttr<Arith_Dialect, Arith_IntegerOverflowFlags, "overflow"> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/ArmSME/IR/ArmSMEOps.td b/mlir/include/mlir/Dialect/ArmSME/IR/ArmSMEOps.td
index 264c3969a1152..5f366ab8e6cbf 100644
--- a/mlir/include/mlir/Dialect/ArmSME/IR/ArmSMEOps.td
+++ b/mlir/include/mlir/Dialect/ArmSME/IR/ArmSMEOps.td
@@ -173,7 +173,6 @@ def TileSliceLayout : I32EnumAttr<"TileSliceLayout", "Layout of a tile slice", [
/// An attribute that specifies the layout of a tile slice in a tile.
def ArmSME_TileSliceLayoutAttr : EnumAttr<ArmSME_Dialect, TileSliceLayout,
"layout"> {
- let assemblyFormat = "`<` $value `>`";
let defaultValue = "TileSliceLayout::Horizontal";
}
@@ -194,7 +193,6 @@ def CombiningKind : I32EnumAttr<"CombiningKind", "Kind of combining function", [
/// FMOPS: https://developer.arm.com/documentation/ddi0602/2022-03/SME-Instructions/FMOPS--non-widening---Floating-point-outer-product-and-subtract-
def ArmSME_CombiningKindAttr : EnumAttr<ArmSME_Dialect, CombiningKind,
"kind"> {
- let assemblyFormat = "`<` $value `>`";
let defaultValue = "CombiningKind::Add";
}
@@ -210,7 +208,6 @@ def TypeSize : I32EnumAttr<"TypeSize", "Size of a vector element type", [
def ArmSME_TypeSizeAttr : EnumAttr<ArmSME_Dialect, TypeSize,
"type_size"> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
index 639a34b665878..55c2bdd88cfe0 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
@@ -96,7 +96,6 @@ class GPU_I32Enum<string name, string description, list<I32EnumAttrCase> cases>
}
class GPU_I32EnumAttr<string mnemonic, GPU_I32Enum enumInfo> :
EnumAttr<GPU_Dialect, enumInfo, mnemonic> {
- let assemblyFormat = "`<` $value `>`";
}
def GPU_AddressSpaceGlobal : I32EnumAttrCase<"Global", 1, "global">;
@@ -131,7 +130,6 @@ def GPU_BarrierScopeEnum : I32Enum<
def GPU_BarrierScopeAttr :
EnumAttr<GPU_Dialect, GPU_BarrierScopeEnum, "barrier_scope"> {
- let assemblyFormat = "`<` $value `>`";
}
def GPU_Dimension : GPU_I32Enum<"Dimension",
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index e6573c94395d9..2aa5219e90404 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -45,7 +45,8 @@ class GPU_IndexOp<string mnemonic, list<Trait> traits = []> :
DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>])>,
Arguments<(ins GPU_DimensionAttr:$dimension,
OptionalAttr<IndexAttr>:$upper_bound)>, Results<(outs Index)> {
- let assemblyFormat = "$dimension (`upper_bound` $upper_bound^)? attr-dict";
+ let assemblyFormat =
+ "enum($dimension) (`upper_bound` $upper_bound^)? attr-dict";
let extraClassDefinition = [{
void $cppClass::getAsmResultNames(
llvm::function_ref<void(mlir::Value, mlir::StringRef)> setNameFn) {
@@ -1215,7 +1216,7 @@ def GPU_AllReduceOp : GPU_Op<"all_reduce",
let results = (outs AnyIntegerOrFloat:$result);
let regions = (region AnyRegion:$body);
- let assemblyFormat = [{ custom<AllReduceOperation>($op) $value
+ let assemblyFormat = [{ (enum($op)^)? $value
(`uniform` $uniform^)? $body attr-dict
`:` functional-type(operands, results) }];
@@ -1294,7 +1295,7 @@ def GPU_SubgroupReduceOp : GPU_Op<"subgroup_reduce", [SameOperandsAndResultType,
}]>
];
- let assemblyFormat = [{ custom<AllReduceOperation>($op) $value
+ let assemblyFormat = [{ enum($op) $value
(`uniform` $uniform^)?
(`cluster` `(` `size` `=` $cluster_size^ (`,` `stride` `=` $cluster_stride^)? `)`)? attr-dict
`:` functional-type(operands, results) }];
@@ -1378,7 +1379,7 @@ def GPU_ShuffleOp : GPU_Op<
}];
let assemblyFormat = [{
- $mode $value `,` $offset `,` $width attr-dict `:` type($value)
+ enum($mode) $value `,` $offset `,` $width attr-dict `:` type($value)
}];
let builders = [
@@ -2262,7 +2263,7 @@ def GPU_SubgroupMmaElementwiseOp : GPU_Op<"subgroup_mma_elementwise",
}];
let assemblyFormat = [{
- $opType $args attr-dict `:` functional-type($args, $res)
+ enum($opType) $args attr-dict `:` functional-type($args, $res)
}];
}
@@ -2582,7 +2583,7 @@ def GPU_Create2To4SpMatOp : GPU_Op<"create_2to4_spmat", [GPU_AsyncOpInterface]>
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- `{` $pruneFlag `}` $rows `,` $cols `,` $memref attr-dict `:` type($memref)
+ `{` enum($pruneFlag) `}` $rows `,` $cols `,` $memref attr-dict `:` type($memref)
}];
}
@@ -2685,7 +2686,7 @@ def GPU_SpMVBufferSizeOp : GPU_Op<"spmv_buffer_size", [GPU_AsyncOpInterface]> {
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $spmatA (`{` $modeA^ `}`)? `,` $dnX `,` $dnY attr-dict `into` $computeType
+ $spmatA (`{` enum($modeA)^ `}`)? `,` $dnX `,` $dnY attr-dict `into` $computeType
}];
}
@@ -2735,7 +2736,7 @@ def GPU_SpMVOp : GPU_Op<"spmv", [GPU_AsyncOpInterface]> {
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $spmatA (`{` $modeA^ `}`)? `,` $dnX `,` $dnY `,` $buffer attr-dict `:` type($buffer) `into` $computeType
+ $spmatA (`{` enum($modeA)^ `}`)? `,` $dnX `,` $dnY `,` $buffer attr-dict `:` type($buffer) `into` $computeType
}];
}
@@ -2788,7 +2789,9 @@ def GPU_SpMMBufferSizeOp : GPU_Op<"spmm_buffer_size", [GPU_AsyncOpInterface, Att
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $spmatA (`{` $modeA^ `}`)? `,` $dnmatB (`{` $modeB^ `}`)? `,` $dnmatC attr-dict `:` type($bufferSzs) `into` $computeType
+ $spmatA (`{` enum($modeA)^ `}`)? `,` $dnmatB
+ (`{` enum($modeB)^ `}`)? `,` $dnmatC attr-dict `:` type($bufferSzs)
+ `into` $computeType
}];
}
@@ -2841,7 +2844,9 @@ def GPU_SpMMOp : GPU_Op<"spmm", [GPU_AsyncOpInterface, AttrSizedOperandSegments]
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $spmatA (`{` $modeA^ `}`)? `,` $dnmatB (`{` $modeB^ `}`)? `,` $dnmatC `,` $buffers attr-dict `:` type($buffers) `into` $computeType
+ $spmatA (`{` enum($modeA)^ `}`)? `,` $dnmatB
+ (`{` enum($modeB)^ `}`)? `,` $dnmatC `,` $buffers attr-dict `:`
+ type($buffers) `into` $computeType
}];
}
@@ -2893,7 +2898,8 @@ def GPU_SDDMMBufferSizeOp : GPU_Op<"sddmm_buffer_size", [GPU_AsyncOpInterface]>
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $dnmatA (`{` $modeA^ `}`)? `,` $dnmatB (`{` $modeB^ `}`)? `,` $spmatC attr-dict `into` $computeType
+ $dnmatA (`{` enum($modeA)^ `}`)? `,` $dnmatB
+ (`{` enum($modeB)^ `}`)? `,` $spmatC attr-dict `into` $computeType
}];
}
@@ -2946,7 +2952,9 @@ def GPU_SDDMMOp : GPU_Op<"sddmm", [GPU_AsyncOpInterface]> {
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $dnmatA (`{` $modeA^ `}`)? `,` $dnmatB (`{` $modeB^ `}`)? `,` $spmatC `,` $buffer attr-dict `:` type($buffer) `into` $computeType
+ $dnmatA (`{` enum($modeA)^ `}`)? `,` $dnmatB
+ (`{` enum($modeB)^ `}`)? `,` $spmatC `,` $buffer attr-dict `:`
+ type($buffer) `into` $computeType
}];
}
@@ -3079,7 +3087,9 @@ def GPU_SpGEMMWorkEstimationOrComputeOp : GPU_Op<"spgemm_work_estimation_or_comp
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- `{` $kind `}` $spmatA (`{` $modeA^ `}`)? `,` $spmatB (`{` $modeB^ `}`)? `,` $spmatC `,` $desc `,` $bufferSz `,` $buffer attr-dict `:` $computeType `into` type($buffer)
+ `{` enum($kind) `}` $spmatA (`{` enum($modeA)^ `}`)? `,` $spmatB
+ (`{` enum($modeB)^ `}`)? `,` $spmatC `,` $desc `,` $bufferSz `,`
+ $buffer attr-dict `:` $computeType `into` type($buffer)
}];
}
@@ -3130,7 +3140,8 @@ def GPU_SpGEMMCopyOp : GPU_Op<"spgemm_copy", [GPU_AsyncOpInterface]> {
let assemblyFormat = [{
custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- $spmatA (`{` $modeA^ `}`)? `,` $spmatB (`{` $modeB^ `}`)? `,` $spmatC `,` $desc attr-dict `:` $computeType
+ $spmatA (`{` enum($modeA)^ `}`)? `,` $spmatB
+ (`{` enum($modeB)^ `}`)? `,` $spmatC `,` $desc attr-dict `:` $computeType
}];
}
@@ -3370,7 +3381,7 @@ def GPU_SubgroupBroadcastOp : GPU_Op<"subgroup_broadcast",
}];
let results = (outs AnyType:$result);
let assemblyFormat = [{
- $src `,` $broadcast_type ($lane^)? attr-dict `:` type($result)
+ $src `,` enum($broadcast_type) ($lane^)? attr-dict `:` type($result)
}];
let hasFolder = 1;
let hasVerifier = 1;
diff --git a/mlir/include/mlir/Dialect/Index/IR/IndexOps.td b/mlir/include/mlir/Dialect/Index/IR/IndexOps.td
index d97d5be698034..cfac59d0db659 100644
--- a/mlir/include/mlir/Dialect/Index/IR/IndexOps.td
+++ b/mlir/include/mlir/Dialect/Index/IR/IndexOps.td
@@ -560,7 +560,7 @@ def Index_CmpOp : IndexOp<"cmp", [Pure]> {
let arguments = (ins IndexCmpPredicateAttr:$pred, Index:$lhs, Index:$rhs);
let results = (outs I1:$result);
- let assemblyFormat = "$pred `(` $lhs `,` $rhs `)` attr-dict";
+ let assemblyFormat = "enum($pred) `(` $lhs `,` $rhs `)` attr-dict";
let hasFolder = 1;
let hasCanonicalizeMethod = 1;
}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMEnums.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMEnums.td
index 4304ec207acf6..4bfa486c16e06 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMEnums.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMEnums.td
@@ -521,7 +521,6 @@ def IntegerOverflowFlags : I32BitEnum<
def LLVM_IntegerOverflowFlagsAttr :
EnumAttr<LLVM_Dialect, IntegerOverflowFlags, "overflow"> {
- let assemblyFormat = "`<` $value `>`";
}
def LLVM_IntegerOverflowFlagsProp :
@@ -557,7 +556,6 @@ def FastmathFlags : I32BitEnumAttr<
def LLVM_FastmathFlagsAttr :
EnumAttr<LLVM_Dialect, FastmathFlags, "fastmath"> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td
index 42d196c5662d1..b6ebb8041e736 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td
@@ -33,7 +33,6 @@ def FPRoundingMode : I32EnumAttr<"FPRoundingMode", "NVVM FPRoundingMode kind",
let cppNamespace = "::mlir::NVVM";
}
def FPRoundingModeAttr : EnumAttr<NVVM_Dialect, FPRoundingMode, "fp_rnd_mode"> {
- let assemblyFormat = "`<` $value `>`";
}
def SaturationModeNone : I32EnumAttrCase<"NONE", 0, "none">;
@@ -66,7 +65,6 @@ def SaturationModeAttr : EnumAttr<NVVM_Dialect, SaturationMode, "sat_mode"> {
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt)
}];
- let assemblyFormat = "`<` $value `>`";
}
#endif // NVVMIR_ENUMS
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9671dca32bd20..7061c33197ee9 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -75,7 +75,6 @@ def CacheEvictionPriority : I32Enum<"CacheEvictionPriority",
def CacheEvictionPriorityAttr : EnumAttr<NVVM_Dialect, CacheEvictionPriority,
"cache_eviction_priority"> {
- let assemblyFormat = "$value";
}
// Memory Space enum definitions
@@ -109,7 +108,6 @@ def NVVMMemorySpaceAttr :
DeclareAttrInterfaceMethods<LLVM_LLVMAddrSpaceAttrInterface>,
DeclareAttrInterfaceMethods<MemorySpaceAttrInterface>
]> {
- let assemblyFormat = "`<` $value `>`";
}
// Attrs describing the scope of the Memory Operation
@@ -124,7 +122,6 @@ def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
let cppNamespace = "::mlir::NVVM";
}
def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
- let assemblyFormat = "$value";
}
// Attrs to disambiguate the cta or cluster space within shared memory
@@ -136,7 +133,6 @@ def SharedSpace : I32EnumAttr<"SharedSpace", "Shared memory space",
let cppNamespace = "::mlir::NVVM";
}
def SharedSpaceAttr : EnumAttr<NVVM_Dialect, SharedSpace, "shared_space"> {
- let assemblyFormat = "`<` $value `>`";
}
// Attrs describing the Memory Ordering Semantics
@@ -157,7 +153,6 @@ def MemOrderKind : I32EnumAttr<"MemOrderKind", "NVVM Memory Ordering kind",
let cppNamespace = "::mlir::NVVM";
}
def MemOrderKindAttr : EnumAttr<NVVM_Dialect, MemOrderKind, "mem_order"> {
- let assemblyFormat = "`<` $value `>`";
}
// Attrs for supported Reduction operations
@@ -575,7 +570,7 @@ def NVVM_ReduxOp :
$res = createIntrinsicCall(builder, intId, {$val, $mask_and_clamp});
}];
let assemblyFormat = [{
- $kind $val `,` $mask_and_clamp oilist(
+ enum($kind) $val `,` $mask_and_clamp oilist(
`abs` `=` $abs | `nan` `=` $nan
) attr-dict `:` type($val) `->` type($res)
}];
@@ -723,7 +718,7 @@ def NVVM_MBarrierExpectTxOp : NVVM_VoidIntrinsicOp<"mbarrier.expect_tx"> {
DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope);
let assemblyFormat = [{
- $addr `,` $txcount oilist(`scope` `=` $scope) attr-dict `:` type(operands)
+ $addr `,` $txcount oilist(`scope` `=` enum($scope)) attr-dict `:` type(operands)
}];
let hasVerifier = 1;
@@ -748,7 +743,7 @@ def NVVM_MBarrierCompleteTxOp : NVVM_VoidIntrinsicOp<"mbarrier.complete_tx"> {
DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope);
let assemblyFormat = [{
- $addr `,` $txcount oilist(`scope` `=` $scope) attr-dict `:` type(operands)
+ $addr `,` $txcount oilist(`scope` `=` enum($scope)) attr-dict `:` type(operands)
}];
let hasVerifier = 1;
@@ -799,7 +794,7 @@ def NVVM_MBarrierArriveOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive",
DefaultValuedAttr<BoolAttr, "false">:$relaxed);
let assemblyFormat = [{
- $addr (`,` $count^)? oilist(`scope` `=` $scope | `relaxed` `=` $relaxed)
+ $addr (`,` $count^)? oilist(`scope` `=` enum($scope) | `relaxed` `=` $relaxed)
attr-dict `:` type($addr)
(`->` type($res)^)?
}];
@@ -829,7 +824,7 @@ def NVVM_MBarrierArriveDropOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_dr
DefaultValuedAttr<BoolAttr, "false">:$relaxed);
let assemblyFormat = [{
- $addr (`,` $count^)? oilist(`scope` `=` $scope | `relaxed` `=` $relaxed)
+ $addr (`,` $count^)? oilist(`scope` `=` enum($scope) | `relaxed` `=` $relaxed)
attr-dict `:` type($addr)
(`->` type($res)^)?
}];
@@ -932,7 +927,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t
let assemblyFormat = [{
$addr `,` $txcount (`,` `predicate` `=` $predicate^)?
- oilist(`scope` `=` $scope | `relaxed` `=` $relaxed)
+ oilist(`scope` `=` enum($scope) | `relaxed` `=` $relaxed)
attr-dict `:` type(operands) (`->` type($res)^)?
}];
let hasVerifier = 1;
@@ -978,7 +973,7 @@ def NVVM_MBarrierArriveDropExpectTxOp : NVVM_SingleResultIntrinsicOp<"mbarrier.a
DefaultValuedAttr<BoolAttr, "false">:$relaxed);
let assemblyFormat = [{
- $addr `,` $txcount oilist(`scope` `=` $scope | `relaxed` `=` $relaxed)
+ $addr `,` $txcount oilist(`scope` `=` enum($scope) | `relaxed` `=` $relaxed)
attr-dict `:` type(operands)
(`->` type($res)^)?
}];
@@ -1107,7 +1102,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.test.wait">
let assemblyFormat = [{
$addr `,` $stateOrPhase oilist(
- `scope` `=` $scope | `relaxed` `=` $relaxed
+ `scope` `=` enum($scope) | `relaxed` `=` $relaxed
) attr-dict `:` type(operands) `->` type($res)
}];
let hasVerifier = 1;
@@ -1140,7 +1135,7 @@ def NVVM_MBarrierTryWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.try_wait"> {
let assemblyFormat = [{
$addr `,` $stateOrPhase (`,` $ticks^)? oilist(
- `scope` `=` $scope | `relaxed` `=` $relaxed
+ `scope` `=` enum($scope) | `relaxed` `=` $relaxed
) attr-dict `:` type(operands) `->` type($res)
}];
let hasVerifier = 1;
@@ -1164,7 +1159,6 @@ def BarrierReduction
}
def BarrierReductionAttr
: EnumAttr<NVVM_Dialect, BarrierReduction, "reduction"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_BarrierOp : NVVM_VoidIntrinsicOp<"barrier",
@@ -1375,7 +1369,7 @@ def NVVM_MembarOp : NVVM_Op<"memory.barrier">,
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar)
}];
- let assemblyFormat = "$scope attr-dict";
+ let assemblyFormat = "enum($scope) attr-dict";
let llvmBuilder = [{
createIntrinsicCall(builder, getMembarIntrinsicID($scope));
}];
@@ -1441,7 +1435,6 @@ def ProxyKindAttr : EnumAttr<NVVM_Dialect, ProxyKind, "proxy_kind"> {
applied to a method of memory access. When two memory operations use distinct
methods of memory access, they are said to be different proxies.
}];
- let assemblyFormat = "`<` $value `>`";
}
defvar ProxyKindNotTensormapOrGeneric =
@@ -1492,7 +1485,7 @@ def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
}];
let assemblyFormat = [{
- $scope $addr `,` $size oilist(
+ enum($scope) $addr `,` $size oilist(
`from_proxy` `=` $fromProxy | `to_proxy` `=` $toProxy
) attr-dict
}];
@@ -1524,7 +1517,7 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
}];
let assemblyFormat = [{
- $scope oilist(
+ enum($scope) oilist(
`from_proxy` `=` $fromProxy | `to_proxy` `=` $toProxy
) attr-dict
}];
@@ -1573,7 +1566,7 @@ def SetMaxRegisterActionAttr : EnumAttr<NVVM_Dialect, SetMaxRegisterAction, "act
def NVVM_SetMaxRegisterOp : NVVM_Op<"setmaxregister"> {
let arguments = (ins I32Attr:$regCount, SetMaxRegisterActionAttr:$action);
- let assemblyFormat = "$action $regCount attr-dict";
+ let assemblyFormat = "enum($action) $regCount attr-dict";
let hasVerifier = 1;
string llvmBuilder = [{
auto intId = (op.getAction() == NVVM::SetMaxRegisterAction::increase) ?
@@ -1631,7 +1624,7 @@ def NVVM_ShflOp :
intId, {$thread_mask, $val, $offset, $mask_and_clamp});
}];
let assemblyFormat = [{
- $kind $thread_mask `,` $val `,` $offset `,` $mask_and_clamp
+ enum($kind) $thread_mask `,` $val `,` $offset `,` $mask_and_clamp
(`return_value_and_is_valid` $return_value_and_is_valid^)? attr-dict
`:` type($val) `->` type($res)
}];
@@ -1679,7 +1672,7 @@ def NVVM_VoteSyncOp
auto intId = getVoteSyncIntrinsicId($kind);
$res = createIntrinsicCall(builder, intId, {$mask, $pred});
}];
- let assemblyFormat = "$kind $mask `,` $pred attr-dict `->` type($res)";
+ let assemblyFormat = "enum($kind) $mask `,` $pred attr-dict `->` type($res)";
let hasVerifier = 1;
}
@@ -1763,7 +1756,6 @@ def PermuteMode : I32EnumAttr<"PermuteMode", "NVVM permute mode",
}
def PermuteModeAttr : EnumAttr<NVVM_Dialect, PermuteMode, "permute_mode"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_PermuteOp : NVVM_SingleResultIntrinsicOp<"prmt", [Pure]>,
@@ -1884,7 +1876,9 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
I32Attr:$size,
LoadCacheModifierAttr:$modifier,
Optional<I32>:$cpSize)> {
- let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` $modifier (`,` $cpSize^)? attr-dict `:` type(operands)";
+ let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` "
+ "enum($modifier) (`,` $cpSize^)? attr-dict `:` "
+ "type(operands)";
let hasVerifier = 1;
let extraClassDeclaration = [{
static llvm::Intrinsic::ID
@@ -2962,7 +2956,6 @@ def MMAB1Op : I32EnumAttr<"MMAB1Op", "MMA binary operations",
let cppNamespace = "::mlir::NVVM";
}
def MMAB1OpAttr : EnumAttr<NVVM_Dialect, MMAB1Op, "mma_b1op"> {
- let assemblyFormat = "`<` $value `>`";
}
/// Enum attribute type for the overflow behavior of MMA integer operations
@@ -2974,7 +2967,6 @@ def MMAIntOverflow : I32EnumAttr<"MMAIntOverflow", "MMA overflow options",
let cppNamespace = "::mlir::NVVM";
}
def MMAIntOverflowAttr : EnumAttr<NVVM_Dialect, MMAIntOverflow, "mma_int_overflow"> {
- let assemblyFormat = "`<` $value `>`";
}
/// MMA kind types (for mixed-precision FP8 operations)
def MMAKindF8F6F4 : I32EnumAttrCase<"f8f6f4", 0>;
@@ -2984,7 +2976,6 @@ def MMAKind : I32EnumAttr<"MMAKind", "MMA operation kind",
let cppNamespace = "::mlir::NVVM";
}
def MMAKindAttr : EnumAttr<NVVM_Dialect, MMAKind, "mma_kind"> {
- let assemblyFormat = "`<` $value `>`";
}
/// Attribute to hold the MMA shape
@@ -3115,7 +3106,6 @@ def MMALayout : I32EnumAttr<"MMALayout", "NVVM MMA layout",
let cppNamespace = "::mlir::NVVM";
}
def MMALayoutAttr : EnumAttr<NVVM_Dialect, MMALayout, "mma_layout"> {
- let assemblyFormat = "`<` $value `>`";
}
/// Enum attribute of the different PTX element types used for MMA operands.
@@ -3146,7 +3136,6 @@ def MMATypes : I32EnumAttr<"MMATypes", "NVVM MMA types",
let cppNamespace = "::mlir::NVVM";
}
def MMATypesAttr : EnumAttr<NVVM_Dialect, MMATypes, "mma_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def MMAFragA : I32EnumAttrCase<"a", 0>;
@@ -3160,7 +3149,6 @@ def MMAFrag: I32EnumAttr<"MMAFrag", "NVVM MMA frag type",
let cppNamespace = "::mlir::NVVM";
}
def MMAFragAttr : EnumAttr<NVVM_Dialect, MMAFrag, "mma_frag"> {
- let assemblyFormat = "$value";
}
def NVVM_WMMALoadOp: NVVM_Op<"wmma.load">,
@@ -3242,7 +3230,7 @@ def NVVM_WMMALoadOp: NVVM_Op<"wmma.load">,
$ptr `,` $stride `,` `m` `=` $m `,` `n` `=` $n `,` `k` `=` $k
`,` `layout` `=` $layout
`,` `element_type` `=` $eltype
- `,` `fragment` `=` $frag attr-dict `:` functional-type($ptr, $res)
+ `,` `fragment` `=` enum($frag) attr-dict `:` functional-type($ptr, $res)
}];
let hasVerifier = 1;
}
@@ -3387,7 +3375,6 @@ def LdStMatrixEltType : I32EnumAttr<"LdStMatrixEltType", "Element type for ldmat
let cppNamespace = "::mlir::NVVM";
}
def LdStMatrixEltTypeAttr : EnumAttr<NVVM_Dialect, LdStMatrixEltType, "ld_st_matrix_elt_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_StMatrixOp: NVVM_Op<"stmatrix">,
@@ -3870,7 +3857,6 @@ def ScaleVecSize : I32EnumAttr<
}
def ScaleVecSizeAttr : EnumAttr<NVVM_Dialect, ScaleVecSize, "scale_vec_size"> {
- let assemblyFormat = "`<` $value `>`";
}
def UE8M0 : I32EnumAttrCase<"UE8M0", 0, "ue8m0">;
@@ -3886,7 +3872,6 @@ def BlockScaleFormat : I32EnumAttr<
}
def BlockScaleFormatAttr : EnumAttr<NVVM_Dialect, BlockScaleFormat, "block_scale_format"> {
- let assemblyFormat = "`<` $value `>`";
}
def MMABlockScaleKindMXF8F6F4 : I32EnumAttrCase<"MXF8F6F4", 0, "mxf8f6f4">;
@@ -3915,7 +3900,6 @@ def MMABlockScaleKindAttr : EnumAttr<NVVM_Dialect, MMABlockScaleKind, "block_sca
+--------------+-------------------------------------------+
```
}];
- let assemblyFormat = "`<` $value `>`";
}
/// Generate enum value of the mma.block_scale intrinsic.
@@ -4339,7 +4323,6 @@ def TMALoadModeAttr : EnumAttr<NVVM_Dialect, TMALoadMode, "tma_load_mode"> {
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes)
}];
- let assemblyFormat = "$value";
}
// List of modes supported for TMA Store and Reduction Ops
@@ -4355,7 +4338,6 @@ def TMAStoreMode : I32EnumAttr<"TMAStoreMode", "NVVM TMA Store Mode",
let cppNamespace = "::mlir::NVVM";
}
def TMAStoreModeAttr : EnumAttr<NVVM_Dialect, TMAStoreMode, "tma_store_mode"> {
- let assemblyFormat = "$value";
}
// Num CTAs in a group participating in the TMA/MMA operations.
@@ -4372,7 +4354,6 @@ def CTAGroupKind : I32EnumAttr<"CTAGroupKind",
}
def CTAGroupKindAttr :
EnumAttr<NVVM_Dialect, CTAGroupKind, "cta_group"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_CpAsyncBulkCommitGroupOp : NVVM_Op<"cp.async.bulk.commit.group">,
@@ -4462,7 +4443,7 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
`im2col` `[` $im2colOffsets `]`
| `multicast_mask` `=` $multicastMask
| `l2_cache_hint` `=` $l2CacheHint | `predicate` `=` $predicate
- | `mode` `=` $mode | `cta_only` `=` $isCTAOnly
+ | `mode` `=` enum($mode) | `cta_only` `=` $isCTAOnly
| `group` `=` custom<CTAGroup>($group)
)
attr-dict `:` type($dstMem) `,` type($tmaDescriptor)
@@ -4549,7 +4530,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
`box` `[`$coordinates `]`
oilist(
`l2_cache_hint` `=` $l2CacheHint | `predicate` `=` $predicate
- | `mode` `=` $mode
+ | `mode` `=` enum($mode)
)
attr-dict `:` type($tmaDescriptor) `,` type($srcMem)
}];
@@ -4672,7 +4653,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
| `lower_stride` `[` $lowerStride `]`
| `upper_stride` `[` $upperStride `]`
| `l2_cache_hint` `=` $l2CacheHint
- | `mode` `=` $mode
+ | `mode` `=` enum($mode)
)
attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
@@ -4694,7 +4675,6 @@ def PrefetchCacheLevel : I32Enum<"PrefetchCacheLevel",
}
def PrefetchCacheLevelAttr : EnumAttr<NVVM_Dialect, PrefetchCacheLevel, "prefetch_cache_level"> {
- let assemblyFormat = "$value";
}
def NVVM_PrefetchOp : NVVM_Op<"prefetch",
@@ -4736,7 +4716,13 @@ def NVVM_PrefetchOp : NVVM_Op<"prefetch",
UnitAttr:$tensormap,
UnitAttr:$uniform,
UnitAttr:$in_param_space);
- let assemblyFormat = "(`level` `=` $cacheLevel^ (`uniform` $uniform^)? `,`)? (`tensormap` $tensormap^ (`in_param_space` $in_param_space^)? `,`)? (`evict_priority` `=` $evictPriority^ `,`)? $addr (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
+ let assemblyFormat = "(`level` `=` enum($cacheLevel)^ "
+ "(`uniform` $uniform^)? `,`)? "
+ "(`tensormap` $tensormap^ "
+ "(`in_param_space` $in_param_space^)? `,`)? "
+ "(`evict_priority` `=` enum($evictPriority)^ `,`)? "
+ "$addr (`,` `predicate` `=` $predicate^)? attr-dict "
+ "`:` type(operands)";
let hasVerifier = 1;
let extraClassDeclaration = [{
@@ -4814,7 +4800,7 @@ def NVVM_CpAsyncBulkTensorPrefetchOp :
`box` `[`$coordinates `]`
oilist(
`im2col` `[` $im2colOffsets `]`
- | `l2_cache_hint` `=` $l2CacheHint | `mode` `=` $mode
+ | `l2_cache_hint` `=` $l2CacheHint | `mode` `=` enum($mode)
)
attr-dict `:` type($tmaDescriptor)
}];
@@ -4844,7 +4830,6 @@ def TMAReduxKind : I32EnumAttr<"TMAReduxKind", "NVVM TMA redux kind",
let cppNamespace = "::mlir::NVVM";
}
def TMAReduxKindAttr : EnumAttr<NVVM_Dialect, TMAReduxKind, "tma_redux_kind"> {
- let assemblyFormat = "$value";
}
def NVVM_CpAsyncBulkTensorReduceOp :
@@ -4877,7 +4862,7 @@ def NVVM_CpAsyncBulkTensorReduceOp :
$srcMem `,`
`box` `[`$coordinates `]`
(`l2_cache_hint` `=` $l2CacheHint^ )?
- `,` `reduction` `=` $redKind oilist(`mode` `=` $mode)
+ `,` `reduction` `=` enum($redKind) oilist(`mode` `=` enum($mode))
attr-dict `:` type($tmaDescriptor) `,` type($srcMem)
}];
@@ -4965,7 +4950,7 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
(`lower_stride` `[`$lowerStride^ `]`)?
(`upper_stride` `[`$upperStride^ `]`)?
(`l2_cache_hint` `=` $l2CacheHint^)?
- `,` `reduction` `=` $redKind oilist(`mode` `=` $mode)
+ `,` `reduction` `=` enum($redKind) oilist(`mode` `=` enum($mode))
attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
@@ -5162,7 +5147,6 @@ def WGMMAScaleIn : I32EnumAttr<"WGMMAScaleIn", "WGMMA overflow options",
let cppNamespace = "::mlir::NVVM";
}
def WGMMAScaleInAttr : EnumAttr<NVVM_Dialect, WGMMAScaleIn, "wgmma_scale_in"> {
- let assemblyFormat = "`<` $value `>`";
}
/// Enum attribute type for the output operand
@@ -5174,7 +5158,6 @@ def WGMMAScaleOut : I32EnumAttr<"WGMMAScaleOut", "WGMMA input predicate",
let cppNamespace = "::mlir::NVVM";
}
def WGMMAScaleOutAttr : EnumAttr<NVVM_Dialect, WGMMAScaleOut, "wgmma_scale_out"> {
- let assemblyFormat = "`<` $value `>`";
}
/// Enum attribute of the different PTX element types used for WGMMA operands.
@@ -5198,7 +5181,6 @@ def WGMMATypes : I32EnumAttr<"WGMMATypes", "NVVM WGMMA types",
let cppNamespace = "::mlir::NVVM";
}
def WGMMATypesAttr : EnumAttr<NVVM_Dialect, WGMMATypes, "wgmma_type"> {
- let assemblyFormat = "`<` $value `>`";
}
@@ -5329,7 +5311,7 @@ def NVVM_GriddepcontrolOp : NVVM_Op<"griddepcontrol", []> {
let arguments = (ins GridDepActionAttr:$kind);
- let assemblyFormat = "$kind attr-dict";
+ let assemblyFormat = "enum($kind) attr-dict";
string llvmBuilder = [{
llvm::Intrinsic::ID id;
@@ -5410,7 +5392,8 @@ def NVVM_MatchSyncOp : NVVM_Op<"match.sync", [InferTypeOpAdaptor]>,
$res = createIntrinsicCall(builder,
intId, {$thread_mask, $val});
}];
- let assemblyFormat = "$kind $thread_mask `,` $val attr-dict `:` type($val) `->` type($res)";
+ let assemblyFormat = "enum($kind) $thread_mask `,` $val attr-dict `:` "
+ "type($val) `->` type($res)";
let hasVerifier = 1;
}
@@ -5472,7 +5455,7 @@ def NVVM_AsyncStoreGlobalOp: NVVM_VoidIntrinsicOp<"store.async.global",
DefaultValuedAttr<BoolAttr, "false">:$mmio);
let results = (outs );
let assemblyFormat = [{
- $addr `,` $value `scope` `=` $scope
+ $addr `,` $value `scope` `=` enum($scope)
oilist(`multimem` `=` $multimem | `mmio` `=` $mmio)
attr-dict `:` type($addr) `,` type($value)
}];
@@ -5542,7 +5525,6 @@ def Tcgen05FenceKind : I32EnumAttr<"Tcgen05FenceKind", "NVVM Tcgen05 fence kind"
}
def Tcgen05FenceKindAttr :
EnumAttr<NVVM_Dialect, Tcgen05FenceKind, "tcgen05_fence"> {
- let assemblyFormat = "`<` $value `>`";
}
def Tcgen05WaitLoad : I32EnumAttrCase<"LOAD", 0, "load">;
@@ -5554,7 +5536,6 @@ def Tcgen05WaitKind : I32EnumAttr<"Tcgen05WaitKind", "NVVM Tcgen05 wait kind",
}
def Tcgen05WaitKindAttr :
EnumAttr<NVVM_Dialect, Tcgen05WaitKind, "tcgen05_wait"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMf<[100, 101, 110]>]> {
@@ -5769,7 +5750,6 @@ def Tcgen05CpShape : I32EnumAttr<"Tcgen05CpShape", "tcgen05 cp shapes",
let genSpecializedAttr = 0;
}
def Tcgen05CpShapeAttr : EnumAttr<NVVM_Dialect, Tcgen05CpShape, "tcgen05_cp_shape"> {
- let assemblyFormat = "$value";
}
def Tcgen05CpMulticastNone: I32EnumAttrCase<"NONE", 0, "none">;
@@ -5784,7 +5764,6 @@ def Tcgen05CpMulticast : I32EnumAttr<"Tcgen05CpMulticast", "tcgen05 cp multicast
let genSpecializedAttr = 0;
}
def Tcgen05CpMulticastAttr : EnumAttr<NVVM_Dialect, Tcgen05CpMulticast, "tcgen05_cp_multicast"> {
- let assemblyFormat = "$value";
}
def FormatB6x16_P32: I32EnumAttrCase<"B6x16_P32", 0, "b6x16_p32">;
@@ -5796,7 +5775,6 @@ def Tcgen05CpSrcFormat : I32EnumAttr<"Tcgen05CpSrcFormat", "tcgen05 cp source fo
let genSpecializedAttr = 0;
}
def Tcgen05CpSrcFormatAttr : EnumAttr<NVVM_Dialect, Tcgen05CpSrcFormat, "tcgen05_cp_src_fmt"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMf<[100, 101, 110]>]> {
@@ -5826,9 +5804,9 @@ def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMf<[100, 101, 110]>]>
I64:$smem_desc);
let assemblyFormat = [{
- $taddr `,` $smem_desc `,` `shape` `=` $shape
+ $taddr `,` $smem_desc `,` `shape` `=` enum($shape)
oilist(
- `group` `=` $group | `multicast` `=` $multicast
+ `group` `=` $group | `multicast` `=` enum($multicast)
| `source_format` `=` $srcFormat
) attr-dict
}];
@@ -5931,7 +5909,6 @@ def Tcgen05LdStShape: I32EnumAttr<
}
def Tcgen05LdStShapeAttr: EnumAttr<NVVM_Dialect, Tcgen05LdStShape, "tcgen05_ldst_shape"> {
- let assemblyFormat = "$value";
}
//===----------------------------------------------------------------------===//
@@ -5954,7 +5931,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMf<[100, 101, 110]>]>
let assemblyFormat = [{
$tmemAddr (`,` $offset^)? (`pack` $pack^)?
- `shape` `=` $shape attr-dict `:` type($res)
+ `shape` `=` enum($shape) attr-dict `:` type($res)
}];
let description = [{
@@ -6045,7 +6022,7 @@ def NVVM_Tcgen05LdRedOp : NVVM_Op<"tcgen05.ld.red",
AnyTypeOf<[I32, F32]>:$redVal);
let assemblyFormat = [{
- $op $addr (`,` $offset^)? `shape` `=` $shape
+ enum($op) $addr (`,` $offset^)? `shape` `=` enum($shape)
oilist(`abs` $abs | `nan` $nan) attr-dict
`:` type($data) `,` type($redVal)
}];
@@ -6132,7 +6109,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMf<[100, 101, 110]>]>
let assemblyFormat = [{
$tmemAddr `,` $val (`,` $offset^)? (`unpack` $unpack^)?
- `shape` `=` $shape attr-dict `:` type($val)
+ `shape` `=` enum($shape) attr-dict `:` type($val)
}];
let description = [{
@@ -6217,7 +6194,6 @@ def DotAccumulateType : I32EnumAttr<"DotAccumulateType",
}
def DotAccumulateTypeAttr : EnumAttr<NVVM_Dialect, DotAccumulateType, "dot_accumulate_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def NVVM_DotAccumulate4WayOp : NVVM_SingleResultIntrinsicOp<"dot.accumulate.4way"> {
@@ -6365,7 +6341,6 @@ def ClusterLaunchControlQueryType
def ClusterLaunchControlQueryTypeAttr
: EnumAttr<NVVM_Dialect,
ClusterLaunchControlQueryType, "cluster_launch_control_query_type"> {
- let assemblyFormat = "$value";
}
def NVVM_ClusterLaunchControlQueryCancelOp
@@ -6392,7 +6367,8 @@ def NVVM_ClusterLaunchControlQueryCancelOp
I128:$try_cancel_response);
let results = (outs AnyTypeOf<[I1, I32]>:$res);
- let assemblyFormat = "`query` `=` $query_type `,` $try_cancel_response attr-dict `:` type($res)";
+ let assemblyFormat = "`query` `=` enum($query_type) `,` "
+ "$try_cancel_response attr-dict `:` type($res)";
let hasVerifier = 1;
@@ -6453,7 +6429,6 @@ def Tcgen05MMAKindAttr : EnumAttr<NVVM_Dialect, Tcgen05MMAKind, "tcgen05_mma_kin
+-------------+--------------------------------------------+
```
}];
- let assemblyFormat = "$value";
}
def Tcgen05MMACollectorOpDiscard : I32EnumAttrCase<"DISCARD", 0, "discard">;
@@ -6481,7 +6456,6 @@ def Tcgen05MMACollectorOpAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorOp, "t
* fill : Fill buffer
* use : Use buffer without modification
}];
- let assemblyFormat = "$value";
}
defvar Tcgen05MMANonBlockScaleKindList =
@@ -6569,10 +6543,11 @@ def NVVM_Tcgen05MMAOp : NVVM_Op<"tcgen05.mma",
let assemblyFormat = [{
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD
oilist(`scale` `=` $scaleInputD | `mask` `=` $disableOutputLane)
- `,` `kind` `=` $kind
+ `,` `kind` `=` enum($kind)
`,` `cta_group` `=` $ctaGroup
oilist(
- `collector_a` `=` $collectorOp | `collector_b` `=` $collectorOpB
+ `collector_a` `=` enum($collectorOp)
+ | `collector_b` `=` enum($collectorOpB)
| `a_shift` $aShift
) attr-dict `:` `(` type(operands) `)`
}];
@@ -6642,10 +6617,11 @@ def NVVM_Tcgen05MMASparseOp : NVVM_Op<"tcgen05.mma.sp",
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD `,`
$sparseMetadata
oilist(`scale` `=` $scaleInputD | `mask` `=` $disableOutputLane)
- `,` `kind` `=` $kind
+ `,` `kind` `=` enum($kind)
`,` `cta_group` `=` $ctaGroup
oilist(
- `collector_a` `=` $collectorOp | `collector_b` `=` $collectorOpB
+ `collector_a` `=` enum($collectorOp)
+ | `collector_b` `=` enum($collectorOpB)
| `a_shift` $aShift
) attr-dict `:` `(` type(operands) `)`
}];
@@ -6680,7 +6656,6 @@ def Tcgen05MMABlockScale
def Tcgen05MMABlockScaleAttr : EnumAttr<NVVM_Dialect, Tcgen05MMABlockScale,
"tcgen05_mma_block_scale"> {
- let assemblyFormat = "$value";
}
def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
@@ -6741,11 +6716,12 @@ def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
let assemblyFormat = [{
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD `,`
- $scaleA `,` $scaleB `,` `kind` `=` $kind
+ $scaleA `,` $scaleB `,` `kind` `=` enum($kind)
`,` `cta_group` `=` $ctaGroup
oilist(
- `block_scale` `=` $blockScale | `collector_a` `=` $collectorOp
- | `collector_b` `=` $collectorOpB
+ `block_scale` `=` enum($blockScale)
+ | `collector_a` `=` enum($collectorOp)
+ | `collector_b` `=` enum($collectorOpB)
)
attr-dict `:` `(` type(operands) `)`
}];
@@ -6813,11 +6789,12 @@ def NVVM_Tcgen05MMASparseBlockScaleOp : NVVM_Op<"tcgen05.mma.sp.block_scale",
let assemblyFormat = [{
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD `,`
- $sparseMetadata `,` $scaleA `,` $scaleB `,` `kind` `=` $kind
+ $sparseMetadata `,` $scaleA `,` $scaleB `,` `kind` `=` enum($kind)
`,` `cta_group` `=` $ctaGroup
oilist(
- `block_scale` `=` $blockScale | `collector_a` `=` $collectorOp
- | `collector_b` `=` $collectorOpB
+ `block_scale` `=` enum($blockScale)
+ | `collector_a` `=` enum($collectorOp)
+ | `collector_b` `=` enum($collectorOpB)
)
attr-dict `:` `(` type(operands) `)`
}];
@@ -6852,7 +6829,6 @@ def Tcgen05MMACollectorBBuffer : I32EnumAttr<
}
def Tcgen05MMACollectorBBufferAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorBBuffer, "tcgen05_mma_collectorb"> {
- let assemblyFormat = "$value";
}
def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
@@ -6909,10 +6885,10 @@ def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
let assemblyFormat = [{
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD (`,` $zeroColMask^)?
- `kind` `=` $kind
+ `kind` `=` enum($kind)
oilist(
- `collector_b_buffer` `=` $collectorBBuffer
- | `collector_b` `=` $collectorOp
+ `collector_b_buffer` `=` enum($collectorBBuffer)
+ | `collector_b` `=` enum($collectorOp)
)
attr-dict `:` `(` type(operands) `)`
}];
@@ -6975,10 +6951,10 @@ def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
let assemblyFormat = [{
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD `,`
- $sparseMetadata (`,` $zeroColMask^)? `kind` `=` $kind
+ $sparseMetadata (`,` $zeroColMask^)? `kind` `=` enum($kind)
oilist(
- `collector_b_buffer` `=` $collectorBBuffer
- | `collector_b` `=` $collectorOp
+ `collector_b_buffer` `=` enum($collectorBBuffer)
+ | `collector_b` `=` enum($collectorOp)
)
attr-dict `:` `(` type(operands) `)`
}];
@@ -7060,7 +7036,8 @@ def NVVM_Tcgen05MMADecompressBOp :
$matrixD `,` $matrixA `,` $matrixB `,` $idesc `,` $enableInputD
`,` $decompressBMetadata (`,` `mask` `=` $disableOutputLane^)?
`cta_group` `=` $ctaGroup
- oilist(`collector_a` `=` $collectorOpA | `collector_b` `=` $collectorOpB)
+ oilist(`collector_a` `=` enum($collectorOpA)
+ | `collector_b` `=` enum($collectorOpB))
attr-dict `:` `(` type(operands) `)`
}];
@@ -7133,9 +7110,9 @@ def NVVM_Tcgen05MMABlockScaleDecompressBOp :
`,` $scaleA `,` $scaleB `,` $decompressBMetadata
`cta_group` `=` $ctaGroup
oilist(
- `block_scale` `=` $blockScale
- | `collector_a` `=` $collectorOpA
- | `collector_b` `=` $collectorOpB
+ `block_scale` `=` enum($blockScale)
+ | `collector_a` `=` enum($collectorOpA)
+ | `collector_b` `=` enum($collectorOpB)
) attr-dict `:` `(` type(operands) `)`
}];
}
@@ -7366,7 +7343,6 @@ def TensormapField : I32Enum<"TensormapField", "NVVM Tensormap Field Kind",
}
def TensormapFieldAttr : EnumAttr<NVVM_Dialect, TensormapField, "tensormap_field"> {
- let assemblyFormat = "$value";
}
// Tensormap Elemtype Enum
@@ -7400,7 +7376,6 @@ def TensormapElemtype : I32Enum<"TensormapElemtype", "NVVM Tensormap Elemtype",
}
def TensormapElemtypeAttr : EnumAttr<NVVM_Dialect, TensormapElemtype, "tensormap_elemtype"> {
- let assemblyFormat = "`<` $value `>`";
}
// Tensormap Interleave Layout Enum
@@ -7419,7 +7394,6 @@ def TensormapInterleaveLayout :
def TensormapInterleaveLayoutAttr : EnumAttr<NVVM_Dialect,
TensormapInterleaveLayout, "tensormap_interleave_layout"> {
- let assemblyFormat = "`<` $value `>`";
}
// Tensormap Swizzle Mode Enum
@@ -7440,7 +7414,6 @@ def TensormapSwizzleMode : I32Enum<"TensormapSwizzleMode", "NVVM Tensormap Swizz
def TensormapSwizzleModeAttr : EnumAttr<NVVM_Dialect,
TensormapSwizzleMode, "tensormap_swizzle_mode"> {
- let assemblyFormat = "`<` $value `>`";
}
// Tensormap Swizzle Atomicity Enum
@@ -7459,7 +7432,6 @@ def TensormapSwizzleAtomicity :
def TensormapSwizzleAtomicityAttr : EnumAttr<NVVM_Dialect,
TensormapSwizzleAtomicity, "tensormap_swizzle_atomicity"> {
- let assemblyFormat = "`<` $value `>`";
}
// Tensormap Fill Mode Enum
@@ -7473,7 +7445,6 @@ def TensormapFillMode : I32Enum<"TensormapFillMode", "NVVM Tensormap Fill Mode",
def TensormapFillModeAttr : EnumAttr<NVVM_Dialect,
TensormapFillMode, "tensormap_fill_mode"> {
- let assemblyFormat = "`<` $value `>`";
}
def TensormapFieldValueAttr :
@@ -7525,7 +7496,8 @@ def NVVM_TensormapReplaceOp : NVVM_VoidIntrinsicOp<"tensormap.replace",
);
let assemblyFormat = [{
- `field` `=` $field (`[` $ord^ `]`)? `,` `new_value` `=` ($new_value_attr^):($new_value)? `in` $addr attr-dict `:` type(operands)
+ `field` `=` enum($field) (`[` $ord^ `]`)? `,` `new_value` `=`
+ ($new_value_attr^):($new_value)? `in` $addr attr-dict `:` type(operands)
}];
}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
index 4f84513d23576..075fa75ab5ad8 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
@@ -34,7 +34,6 @@ def ROCDL_BufferOOBModeAttr :
Specifies the AMDGPU buffer out-of-bounds mode encoded in
`amdgpu.buffer.oob.mode` or `amdgpu.tbuffer.oob.mode` LLVM module flags.
}];
- let assemblyFormat = "`<` $value `>`";
}
class ROCDL_BufferOOBModeModuleFlagAttrBase<string attrName, string mnemonic>
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 4e17e5e229cbb..439fb0f32a409 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -757,7 +757,7 @@ def ROCDL_SetPrioOp : ROCDL_ConcreteNonMemIntrOp<"s.setprio", [], 0, [0], ["prio
def ROCDL_SchedBarrier : ROCDL_ConcreteNonMemIntrOp<"sched.barrier", [], 0, [0],["mask"]>,
Arguments<(ins ROCDL_SchedGroupMaskAttr:$mask)> {
- let assemblyFormat = "$mask attr-dict";
+ let assemblyFormat = "enum($mask) attr-dict";
let description = [{
Insert a scheduling barrier with the given mask. The mask is a
bitfield that controls which instruction types may be scheduled
@@ -785,7 +785,7 @@ def ROCDL_SchedGroupBarrier
[0, 1, 2], ["mask", "size", "groupId"]>,
Arguments<(ins ROCDL_SchedGroupMaskAttr:$mask, I32Attr:$size,
I32Attr:$groupId)> {
- let assemblyFormat = "$mask `,` $size `,` $groupId attr-dict";
+ let assemblyFormat = "enum($mask) `,` $size `,` $groupId attr-dict";
let description = [{
Insert a scheduling group barrier. The first parameter uses the same
scheduling group mask values as `rocdl.sched.barrier`.
@@ -827,7 +827,8 @@ class ROCDL_Mfma_IntrOpBase<string mnemonic, ROCDL_NamedType ABType,
blgpAttr:$blgp)> {
let results = (outs CDType:$res);
let assemblyFormat = [{
- $a `,` $b `,` $c `,` $cbsz `,` $abid `,` $blgp attr-dict `:` functional-type(operands, $res)
+ $a `,` $b `,` $c `,` $cbsz `,` $abid `,` enum($blgp)
+ attr-dict `:` functional-type(operands, $res)
}];
}
@@ -877,7 +878,9 @@ class ROCDL_Mfma_Scale_IntrOp<string mnemonic, Type AB, Type CD> :
I32:$scaleB)> {
let results = (outs LLVM_ScalarOrVectorOf<CD>:$res);
let assemblyFormat = [{
- $a `,` $b `,` $c `,` $cbsz `,` $blgp `,` $opselA `,` $scaleA `,` $opselB `,` $scaleB attr-dict `:` functional-type(operands, $res)
+ $a `,` $b `,` $c `,` enum($cbsz) `,` enum($blgp) `,` $opselA `,`
+ $scaleA `,` $opselB `,` $scaleB attr-dict `:`
+ functional-type(operands, $res)
}];
let description = [{
Scaled matrix fused multiply-add (MFMA) intrinsic with per-operand scaling.
@@ -1115,7 +1118,7 @@ class ROCDL_WMMA_ModsAll_Reuse_IntrOp<string mnemonic, Type AB, Type CD> : ROCDL
DefaultValuedAttr<I1Attr, "0">:$reuseB)> {
let results = (outs LLVM_ScalarOrVectorOf<CD>:$res);
let assemblyFormat = [{
- $a `,` $b `,` $c `,` `modC` `=` $modC prop-dict attr-dict `:`
+ $a `,` $b `,` $c `,` `modC` `=` enum($modC) prop-dict attr-dict `:`
functional-type(operands, $res)
}];
let description = [{
@@ -1142,7 +1145,7 @@ class ROCDL_WMMA_ModsC_IntrOp<string mnemonic, Type AB, Type CD> : ROCDL_IntrOp<
DefaultValuedAttr<I1Attr, "0">:$reuseB)> {
let results = (outs LLVM_ScalarOrVectorOf<CD>:$res);
let assemblyFormat = [{
- $a `,` $b `,` $c `,` `modC` `=` $modC prop-dict attr-dict `:`
+ $a `,` $b `,` $c `,` `modC` `=` enum($modC) prop-dict attr-dict `:`
functional-type(operands, $res)
}];
let description = [{
@@ -1169,7 +1172,7 @@ class ROCDL_WMMA_ModsC_Diff_IntrOp<string mnemonic, Type AB, Type C, Type D> : R
DefaultValuedAttr<I1Attr, "0">:$reuseB)> {
let results = (outs LLVM_ScalarOrVectorOf<D>:$res);
let assemblyFormat = [{
- $a `,` $b `,` $c `,` `modC` `=` $modC prop-dict attr-dict `:`
+ $a `,` $b `,` $c `,` `modC` `=` enum($modC) prop-dict attr-dict `:`
functional-type(operands, $res)
}];
let description = [{
@@ -1244,10 +1247,12 @@ class ROCDL_WMMA_Scale_IntrOp<string mnemonic, Type AB, Type CD, Type ScaleExpTy
let results = (outs LLVM_ScalarOrVectorOf<CD>:$res);
let assemblyFormat = [{
$a `,` $b `,` $c `,` $scaleA `,` $scaleB `,`
- `fmtA` `=` $fmtA `,` `fmtB` `=` $fmtB `,`
- `modC` `=` $modC `,`
- `scaleAType` `=` $scaleAType `,` `fmtScaleA` `=` $fmtScaleA `,`
- `scaleBType` `=` $scaleBType `,` `fmtScaleB` `=` $fmtScaleB
+ `fmtA` `=` enum($fmtA) `,` `fmtB` `=` enum($fmtB) `,`
+ `modC` `=` enum($modC) `,`
+ `scaleAType` `=` enum($scaleAType) `,`
+ `fmtScaleA` `=` enum($fmtScaleA) `,`
+ `scaleBType` `=` enum($scaleBType) `,`
+ `fmtScaleB` `=` enum($fmtScaleB)
prop-dict attr-dict `:` functional-type(operands, $res)
}];
let description = [{
@@ -1289,9 +1294,11 @@ class ROCDL_WMMA_Scale_F4_IntrOp<string mnemonic, Type AB, Type CD, Type ScaleEx
let results = (outs LLVM_ScalarOrVectorOf<CD>:$res);
let assemblyFormat = [{
$a `,` $b `,` $c `,` $scaleA `,` $scaleB `,`
- `modC` `=` $modC `,`
- `scaleAType` `=` $scaleAType `,` `fmtScaleA` `=` $fmtScaleA `,`
- `scaleBType` `=` $scaleBType `,` `fmtScaleB` `=` $fmtScaleB
+ `modC` `=` enum($modC) `,`
+ `scaleAType` `=` enum($scaleAType) `,`
+ `fmtScaleA` `=` enum($fmtScaleA) `,`
+ `scaleBType` `=` enum($scaleBType) `,`
+ `fmtScaleB` `=` enum($fmtScaleB)
prop-dict attr-dict `:` functional-type(operands, $res)
}];
let description = [{
diff --git a/mlir/include/mlir/Dialect/LLVMIR/XeVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/XeVMOps.td
index 65e35aa6163d1..2cd7bbb07244c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/XeVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/XeVMOps.td
@@ -130,7 +130,6 @@ def XeVM_LoadCacheControl
def XeVM_LoadCacheControlAttr
: EnumAttr<XeVM_Dialect, XeVM_LoadCacheControl, "load_cache_control"> {
let summary = [{Describe the cache settings for load operators}];
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -195,7 +194,6 @@ def XeVM_StoreCacheControl
def XeVM_StoreCacheControlAttr
: EnumAttr<XeVM_Dialect, XeVM_StoreCacheControl, "store_cache_control"> {
let summary = [{Describe the cache settings for store operators}];
- let assemblyFormat = "`<` $value `>`";
}
def XeVM_BlockLoadOp
@@ -409,7 +407,6 @@ def XeVM_MemScope
}
def XeVM_MemScopeAttr : EnumAttr<XeVM_Dialect, XeVM_MemScope, "mem_scope"> {
let summary = [{Describe memory scopes}];
- let assemblyFormat = "`<` $value `>`";
}
def AddrSpacePrivate : I32EnumAttrCase<"PRIVATE", 0, "private">;
@@ -427,7 +424,6 @@ def XeVM_AddrSpace
}
def XeVM_AddrSpaceAttr : EnumAttr<XeVM_Dialect, XeVM_AddrSpace, "addr_space"> {
let summary = [{Describe address spaces}];
- let assemblyFormat = "`<` $value `>`";
}
def XeVM_MemfenceOp
diff --git a/mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td b/mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td
index 55b62b643b0e9..e924161165fa1 100644
--- a/mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td
+++ b/mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td
@@ -66,25 +66,19 @@ def Linalg_Dialect : Dialect {
// Define the attribute enums matching elementwise op kind (e.g., add).
def ElementwiseKindAttr : EnumAttr<Linalg_Dialect,
ElementwiseKind, "elementwise_kind"> {
- let assemblyFormat = "`<` $value `>`";
}
// Define the function attribute enums matching the OpDSL functions.
def UnaryFnAttr : EnumAttr<Linalg_Dialect, UnaryFn, "unary_fn"> {
- let assemblyFormat = "`<` $value `>`";
}
def BinaryFnAttr : EnumAttr<Linalg_Dialect, BinaryFn, "binary_fn"> {
- let assemblyFormat = "`<` $value `>`";
}
def TernaryFnAttr : EnumAttr<Linalg_Dialect, TernaryFn, "ternary_fn"> {
- let assemblyFormat = "`<` $value `>`";
}
def TypeFnAttr : EnumAttr<Linalg_Dialect, TypeFn, "type_fn"> {
- let assemblyFormat = "`<` $value `>`";
}
def IteratorTypeEnum : EnumAttr<Linalg_Dialect, IteratorType, "iterator_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def IteratorTypeArrayAttr : TypedArrayAttrBase<IteratorTypeEnum,
"Iterator type should be an enum.">;
diff --git a/mlir/include/mlir/Dialect/MPI/IR/MPI.td b/mlir/include/mlir/Dialect/MPI/IR/MPI.td
index 6869df20936d4..3859bf594f591 100644
--- a/mlir/include/mlir/Dialect/MPI/IR/MPI.td
+++ b/mlir/include/mlir/Dialect/MPI/IR/MPI.td
@@ -213,7 +213,6 @@ def MPI_ErrorClassEnum
}
def MPI_ErrorClassAttr : EnumAttr<MPI_Dialect, MPI_ErrorClassEnum, "errclass"> {
- let assemblyFormat = "`<` $value `>`";
}
def MPI_OpNull : I32EnumAttrCase<"MPI_OP_NULL", 0, "MPI_OP_NULL">;
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 1cdd05c6656b5..9a0ccfeeee4af 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -115,7 +115,6 @@ def OpenACC_ReductionOperator
def OpenACC_ReductionOperatorAttr : EnumAttr<OpenACC_Dialect,
OpenACC_ReductionOperator,
"reduction_operator"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
// OpenACC variable type categorization. This is needed because OpenACC
@@ -227,7 +226,6 @@ def OpenACC_ParLevel : I32EnumAttr<"ParLevel",
def OpenACC_ParLevelAttr : EnumAttr<OpenACC_Dialect,
OpenACC_ParLevel,
"par_level"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def OpenACC_PrivateRecipe : I32EnumAttrCase<"private_recipe", 0>;
@@ -247,7 +245,6 @@ def OpenACC_RecipeKind : I32EnumAttr<
def OpenACC_RecipeKindAttr : EnumAttr<OpenACC_Dialect,
OpenACC_RecipeKind,
"recipe_kind"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
// Type used in operation below.
@@ -365,7 +362,7 @@ class OpenACC_Attr<string name, string attrMnemonic,
// easier to find out whether the variable is in a declare clause and what kind
// of clause it is.
def DeclareAttr : OpenACC_Attr<"Declare", "declare"> {
- let parameters = (ins "DataClauseAttr":$dataClause,
+ let parameters = (ins EnumAttrParameter<OpenACC_DataClauseAttr>:$dataClause,
DefaultValuedParameter<"bool", "false">:$implicit);
let assemblyFormat = "`<` struct(params) `>`";
let builders = [AttrBuilder<(ins "DataClauseAttr":$dataClause), [{
@@ -426,7 +423,6 @@ def OpenACC_DeviceType : I32EnumAttr<"DeviceType",
def OpenACC_DeviceTypeAttr : EnumAttr<OpenACC_Dialect,
OpenACC_DeviceType,
"device_type"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def DeviceTypeArrayAttr :
@@ -448,7 +444,6 @@ def OpenACC_GangArgType : I32EnumAttr<"GangArgType",
def OpenACC_GangArgTypeAttr : EnumAttr<OpenACC_Dialect,
OpenACC_GangArgType,
"gang_arg_type"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def GangArgTypeArrayAttr :
TypedArrayAttrBase<OpenACC_GangArgTypeAttr, "gang arg type array attribute"> {
@@ -470,7 +465,6 @@ def OpenACC_CombinedConstructsType : I32EnumAttr<"CombinedConstructsType",
def OpenACC_CombinedConstructsAttr : EnumAttr<OpenACC_Dialect,
OpenACC_CombinedConstructsType,
"combined_constructs"> {
- let assemblyFormat = [{ ```<` $value `>` }];
}
def OpenACC_ParallelConstruct : I64EnumAttrCase<"acc_construct_parallel", 0>;
@@ -817,7 +811,7 @@ class OpenACC_DataEntryOp<string mnemonic, string clause,
| `async` `` custom<DeviceTypeOperandsWithKeywordOnly>($asyncOperands,
type($asyncOperands), $asyncOperandsDeviceType, $asyncOnly)
| `recipe` `(` custom<RecipeSym>($recipe) `)`
- | `dataClause` `(` qualified($dataClause) `)`
+ | `dataClause` `(` enum($dataClause) `)`
| `structured` `(` $structured `)`
| `implicit` `(` $implicit `)`
| `name` `(` $name `)`
@@ -1211,7 +1205,7 @@ class OpenACC_DataExitOpWithVarPtr<string mnemonic, string clause>
type($asyncOperands), $asyncOperandsDeviceType, $asyncOnly)^)?
`to` custom<Var>($var) `:` custom<VarPtrType>(type($var), $varType)
oilist(
- `dataClause` `(` qualified($dataClause) `)`
+ `dataClause` `(` enum($dataClause) `)`
| `structured` `(` $structured `)`
| `implicit` `(` $implicit `)`
| `name` `(` $name `)`
@@ -1289,7 +1283,7 @@ class OpenACC_DataExitOpNoVarPtr<string mnemonic, string clause>
(`async` `` custom<DeviceTypeOperandsWithKeywordOnly>($asyncOperands,
type($asyncOperands), $asyncOperandsDeviceType, $asyncOnly)^)?
oilist(
- `dataClause` `(` qualified($dataClause) `)`
+ `dataClause` `(` enum($dataClause) `)`
| `structured` `(` $structured `)`
| `implicit` `(` $implicit `)`
| `name` `(` $name `)`
@@ -1953,7 +1947,7 @@ def OpenACC_ParallelOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2100,7 +2094,7 @@ def OpenACC_SerialOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2296,7 +2290,7 @@ def OpenACC_KernelsOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2416,7 +2410,7 @@ def OpenACC_DataOp
$waitOnly)
)
$region
- (`defaultAttr` `(` qualified($defaultAttr)^ `)`)?
+ (`defaultAttr` `(` enum($defaultAttr)^ `)`)?
attr-dict-with-keyword
}];
let hasVerifier = 1;
@@ -3214,9 +3208,9 @@ def OpenACC_DeclareExitOp : OpenACC_Op<"declare_exit",
Example showing `acc declare device_resident(a)`:
```mlir
- %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause declare_device_resident>}
+ %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr {dataClause = #acc.data_clause<declare_device_resident>}
acc.declare_exit dataOperands(%0 : !llvm.ptr)
- acc.delete accPtr(%0 : !llvm.ptr) {dataClause = #acc<data_clause declare_device_resident>}
+ acc.delete accPtr(%0 : !llvm.ptr) {dataClause = #acc.data_clause<declare_device_resident>}
```
}];
@@ -3290,9 +3284,9 @@ def OpenACC_GlobalDestructorOp : OpenACC_Op<"global_dtor",
}
acc.global_dtor @acc_destructor {
%0 = llvm.mlir.addressof @globalvar : !llvm.ptr
- %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause create>}
+ %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) -> !llvm.ptr {dataClause = #acc.data_clause<create>}
acc.declare_exit dataOperands(%1 : !llvm.ptr)
- acc.delete accPtr(%1 : !llvm.ptr) {dataClause = #acc<data_clause create>}
+ acc.delete accPtr(%1 : !llvm.ptr) {dataClause = #acc.data_clause<create>}
}
```
}];
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPClauses.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPClauses.td
index ca091cf8729b2..bd9f1321f077b 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPClauses.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPClauses.td
@@ -396,7 +396,7 @@ class OpenMP_DoacrossClauseSkip<
);
let reqAssemblyFormat = [{
- ( `depend_type` `` $doacross_depend_type^ )?
+ ( `depend_type` `(` enum($doacross_depend_type)^ `)` )?
( `depend_vec` `(` $doacross_depend_vars^ `:` type($doacross_depend_vars)
`)` )?
( `num_loops` `(` $doacross_num_loops^ `)` )?
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td
index f39f0d00ad6e1..c865897e462cf 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td
@@ -56,7 +56,6 @@ def DeclareTargetCaptureClause : OpenMP_I32EnumAttr<
def DeclareTargetCaptureClauseAttr : OpenMP_EnumAttr<DeclareTargetCaptureClause,
"capture_clause"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -74,7 +73,6 @@ def ClauseDepend : OpenMP_I32EnumAttr<
]>;
def ClauseDependAttr : OpenMP_EnumAttr<ClauseDepend, "clause_depend"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -94,7 +92,6 @@ def ClauseSeverity : OpenMP_I32EnumAttr<
]>;
def ClauseSeverityAttr : OpenMP_EnumAttr<ClauseSeverity, "clause_severity"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -198,7 +195,6 @@ def ClauseTaskDepend
def ClauseTaskDependAttr : OpenMP_EnumAttr<ClauseTaskDepend,
"clause_task_depend"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -218,7 +214,6 @@ def DataSharingClauseType : OpenMP_I32EnumAttr<
def DataSharingClauseTypeAttr : OpenMP_EnumAttr<DataSharingClauseType,
"data_sharing_type"> {
- let assemblyFormat = "`{` `type` `=` $value `}`";
}
//===----------------------------------------------------------------------===//
@@ -239,7 +234,6 @@ def DeclareTargetDeviceType : OpenMP_I32EnumAttr<
def DeclareTargetDeviceTypeAttr : OpenMP_EnumAttr<DeclareTargetDeviceType,
"device_type"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -270,7 +264,6 @@ def LinearModifier : OpenMP_I32EnumAttr<"LinearModifier", "linear modifier",
LinearModifierUval]>;
def LinearModifierAttr : OpenMP_EnumAttr<LinearModifier, "linear_modifier"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -291,7 +284,6 @@ def ReductionModifier : OpenMP_I32EnumAttr<
def ReductionModifierAttr : OpenMP_EnumAttr<ReductionModifier,
"reduction_modifier"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -335,7 +327,6 @@ def TargetExecMode : OpenMP_I32EnumAttr<
]>;
def TargetExecModeAttr : OpenMP_EnumAttr<TargetExecMode, "kernel_type"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -358,7 +349,6 @@ def VariableCaptureKind : OpenMP_I32EnumAttr<
def VariableCaptureKindAttr : OpenMP_EnumAttr<VariableCaptureKind,
"variable_capture_kind"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -375,7 +365,6 @@ def AccessGroupModifier : OpenMP_I32EnumAttr<
def AccessGroupModifierAttr : OpenMP_EnumAttr<AccessGroupModifier,
"access_group_modifier"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -396,7 +385,6 @@ def FallbackModifier : OpenMP_I32EnumAttr<
def FallbackModifierAttr : OpenMP_EnumAttr<FallbackModifier,
"fallback_modifier"> {
- let assemblyFormat = "`(` $value `)`";
}
//===----------------------------------------------------------------------===//
@@ -414,7 +402,6 @@ def InteropType : OpenMP_I32EnumAttr<
]>;
def InteropTypeAttr : OpenMP_EnumAttr<InteropType, "interop_type"> {
- let assemblyFormat = "`(` $value `)`";
}
#endif // OPENMP_ENUMS
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
index 963aece310af2..0df33d96319f5 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
@@ -115,7 +115,8 @@ def PrivateClauseOp
AnyRegion:$dealloc_region);
let assemblyFormat = [{
- $data_sharing_type ($sym_visibility^)? $sym_name `:` $type
+ ` ` `{` `type` `=` enum($data_sharing_type) `}`
+ ($sym_visibility^)? $sym_name `:` $type
(`init` $init_region^)?
(`copy` $copy_region^)?
(`dealloc` $dealloc_region^)? attr-dict
@@ -1725,7 +1726,7 @@ def TargetOp : OpenMP_Op<"target", traits = [
bool hasHostEvalTripCount();
}] # clausesExtraClassDeclaration;
- let assemblyFormat = "`kernel_type` `` $kernel_type " #
+ let assemblyFormat = "`kernel_type` `(` enum($kernel_type) `)` " #
clausesReqAssemblyFormat # " " #
"oilist(" # clausesOptAssemblyFormat #
" | `map_iterated` `(` $map_iterated `:` type($map_iterated) `)`)" # [{
@@ -2121,7 +2122,7 @@ def AtomicCompareOp : OpenMP_Op<"atomic.compare", traits = [
let assemblyFormat = clausesAssemblyFormat #
"$x `:` type($x) $region "
"oilist(`weak` $weak | "
- "`fail_memory_order` `(` $fail_memory_order `)`) attr-dict";
+ "`fail_memory_order` `(` enum($fail_memory_order) `)`) attr-dict";
let hasVerifier = 1;
let hasRegionVerifier = 1;
@@ -2750,7 +2751,8 @@ def GroupprivateOp : OpenMP_Op<"groupprivate",
);
let results = (outs OpenMP_PointerLikeType:$gp_addr);
let assemblyFormat = [{
- $sym_name (`device_type` $device_type^)? `:` type($gp_addr) attr-dict
+ $sym_name (`device_type` ` ` `(` enum($device_type)^ `)`)?
+ `:` type($gp_addr) attr-dict
}];
}
diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td
index 92653117a451a..e01818dea8278 100644
--- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td
+++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td
@@ -89,7 +89,6 @@ class SPIRV_BitEnum<string name, string description,
class SPIRV_BitEnumAttr<string name, string description, string mnemonic,
list<BitEnumAttrCaseBase> cases> :
EnumAttr<SPIRV_Dialect, SPIRV_BitEnum<name, description, cases>, mnemonic> {
- let assemblyFormat = "`<` $value `>`";
}
// Wrapper over base I32EnumAttr to set common fields.
@@ -102,7 +101,6 @@ class SPIRV_I32Enum<string name, string description,
class SPIRV_I32EnumAttr<string name, string description, string mnemonic,
list<I32EnumAttrCase> cases> :
EnumAttr<SPIRV_Dialect, SPIRV_I32Enum<name, description, cases>, mnemonic> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td b/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
index 9af607f7bca5b..8bdf96fd6b958 100644
--- a/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
+++ b/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
@@ -63,7 +63,6 @@ def Shard_ReductionKind : I32EnumAttr<"ReductionKind",
}
def Shard_ReductionKindAttr : EnumAttr<Shard_Dialect, Shard_ReductionKind, "partial"> {
- let assemblyFormat = "$value";
}
class Shard_Type<string name, string typeMnemonic, list<Trait> traits = [],
diff --git a/mlir/include/mlir/Dialect/Shard/IR/ShardOps.td b/mlir/include/mlir/Dialect/Shard/IR/ShardOps.td
index 95b7f5624a246..2539eb93f7a59 100644
--- a/mlir/include/mlir/Dialect/Shard/IR/ShardOps.td
+++ b/mlir/include/mlir/Dialect/Shard/IR/ShardOps.td
@@ -576,7 +576,8 @@ def Shard_AllReduceOp : Shard_CollectiveCommunicationOpBase<"all_reduce", [
AnyTypeOf<[AnyMemRef, AnyRankedTensor]>:$result
);
let assemblyFormat = [{
- $input `on` $grid (`grid_axes` `=` $grid_axes^)? (`reduction` `=` $reduction^)?
+ $input `on` $grid (`grid_axes` `=` $grid_axes^)?
+ (`reduction` `=` enum($reduction)^)?
attr-dict `:` type($input) `->` type($result)
}];
let hasCanonicalizer = 1;
@@ -888,7 +889,7 @@ def Shard_ReduceOp : Shard_CollectiveCommunicationOpBase<"reduce", [
);
let assemblyFormat = [{
$input `on` $grid (`grid_axes` `=` $grid_axes^)?
- (`reduction` `=` $reduction^)?
+ (`reduction` `=` enum($reduction)^)?
`root` `=` custom<DynamicIndexList>($root_dynamic, $root)
attr-dict `:` functional-type(operands, results)
}];
@@ -952,7 +953,7 @@ def Shard_ReduceScatterOp : Shard_CollectiveCommunicationOpBase<"reduce_scatter"
);
let assemblyFormat = [{
$input `on` $grid (`grid_axes` `=` $grid_axes^)?
- (`reduction` `=` $reduction^)?
+ (`reduction` `=` enum($reduction)^)?
`scatter_dim` `=` $scatter_dim
attr-dict `:` type($input) `->` type($result)
}];
diff --git a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorOps.td b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorOps.td
index 6f235c5dc74bb..6858fd0853fd0 100644
--- a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorOps.td
+++ b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorOps.td
@@ -564,7 +564,7 @@ def SparseTensor_GetStorageSpecifierOp : SparseTensor_Op<"storage_specifier.get"
SparseTensorStorageSpecifierKindAttr:$specifierKind,
OptionalAttr<LevelAttr>:$level);
let results = (outs Index:$result);
- let assemblyFormat = "$specifier $specifierKind (`at` $level^)? attr-dict"
+ let assemblyFormat = "$specifier enum($specifierKind) (`at` $level^)? attr-dict"
"`:` qualified(type($specifier))";
let hasVerifier = 1;
@@ -591,7 +591,7 @@ def SparseTensor_SetStorageSpecifierOp : SparseTensor_Op<"storage_specifier.set"
OptionalAttr<LevelAttr>:$level,
Index:$value);
let results = (outs SparseTensorStorageSpecifier:$result);
- let assemblyFormat = "$specifier $specifierKind (`at` $level^)? `with` $value"
+ let assemblyFormat = "$specifier enum($specifierKind) (`at` $level^)? `with` $value"
" attr-dict `:` qualified(type($result))";
let hasVerifier = 1;
@@ -675,7 +675,8 @@ def SparseTensor_CrdTranslateOp : SparseTensor_Op<"crd_translate", [Pure]> {
SparseTensorCrdTransDirectionAttr:$direction,
SparseTensorEncodingAttr:$encoder);
let results = (outs Variadic<Index>:$out_crds);
- let assemblyFormat = "$direction `[` $in_crds `]` `as` $encoder attr-dict `:` type($out_crds)";
+ let assemblyFormat = "enum($direction) `[` $in_crds `]` `as` $encoder "
+ "attr-dict `:` type($out_crds)";
let hasVerifier = 1;
let hasFolder = 1;
@@ -935,7 +936,7 @@ def SparseTensor_SortOp : SparseTensor_Op<"sort"> {
Variadic<StridedMemRefRankOf<[AnyType], [1]>>:$ys,
AffineMapAttr:$perm_map, OptionalAttr<IndexAttr>:$ny,
SparseTensorSortKindAttr:$algorithm);
- let assemblyFormat = "$algorithm $n"
+ let assemblyFormat = "enum($algorithm) $n"
"`,`$xy (`jointly` $ys^)? `perm_map` `=` $perm_map"
" (`ny` `=` $ny^)? attr-dict"
"`:` type($xy) (`jointly` type($ys)^)?";
@@ -965,7 +966,7 @@ def SparseTensor_ReorderCOOOp : SparseTensor_Op<"reorder_coo", [Pure]> {
let arguments = (ins AnySparseTensor: $input_coo,
SparseTensorSortKindAttr:$algorithm);
let results = (outs AnySparseTensor: $result_coo);
- let assemblyFormat = "$algorithm $input_coo attr-dict"
+ let assemblyFormat = "enum($algorithm) $input_coo attr-dict"
"`:` type($input_coo) `to` type($result_coo)";
let hasFolder = 1;
diff --git a/mlir/include/mlir/Dialect/Tosa/IR/TosaOpBase.td b/mlir/include/mlir/Dialect/Tosa/IR/TosaOpBase.td
index 63acbb09b7a04..d4043d24f6db7 100644
--- a/mlir/include/mlir/Dialect/Tosa/IR/TosaOpBase.td
+++ b/mlir/include/mlir/Dialect/Tosa/IR/TosaOpBase.td
@@ -234,7 +234,6 @@ class Tosa_I32Enum<string name, string description, list<I32EnumAttrCase> cases>
class Tosa_I32EnumAttr<string name, string description, string mnemonic,
list<I32EnumAttrCase> cases>
: EnumAttr<Tosa_Dialect, Tosa_I32Enum<name, description, cases>, mnemonic> {
- let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorAttributes.td b/mlir/include/mlir/Dialect/Vector/IR/VectorAttributes.td
index bcf53dafbce9b..dcf1513a81f9d 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorAttributes.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorAttributes.td
@@ -46,7 +46,6 @@ def CombiningKind : I32EnumAttr<
/// An attribute that specifies the combining function for `vector.contract`,
/// and `vector.reduction`.
def Vector_CombiningKindAttr : EnumAttr<Vector_Dialect, CombiningKind, "kind"> {
- let assemblyFormat = "`<` $value `>`";
}
def Vector_IteratorType : I32EnumAttr<"IteratorType", "Iterator type", [
@@ -59,7 +58,6 @@ def Vector_IteratorType : I32EnumAttr<"IteratorType", "Iterator type", [
def Vector_IteratorTypeEnum
: EnumAttr<Vector_Dialect, Vector_IteratorType, "iterator_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def Vector_IteratorTypeArrayAttr
@@ -79,7 +77,6 @@ def PrintPunctuation : I32EnumAttr<"PrintPunctuation",
}
def Vector_PrintPunctuation : EnumAttr<Vector_Dialect, PrintPunctuation, "punctuation"> {
- let assemblyFormat = "`<` $value `>`";
}
#endif // MLIR_DIALECT_VECTOR_IR_VECTOR_ATTRIBUTES
diff --git a/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
index 2b977fc418553..0a57ee52685dd 100644
--- a/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
+++ b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
@@ -19,6 +19,24 @@ class XeGPUAttr<string name, string attrMnemonic, list<Trait> traits = [],
let mnemonic = attrMnemonic;
}
+//===----------------------------------------------------------------------===//
+// XeGPU Memory Scope Enums.
+//===----------------------------------------------------------------------===//
+def XeGPU_MemorySpaceGlobal: I32EnumAttrCase<"Global", 0, "global">;
+def XeGPU_MemorySpaceShared: I32EnumAttrCase<"SLM", 3, "slm">;
+def XeGPU_MemorySpace: I32EnumAttr<"MemorySpace",
+ "The address space of the memory the tensor descritor is created for",
+ [XeGPU_MemorySpaceGlobal, XeGPU_MemorySpaceShared]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::xegpu";
+}
+
+def XeGPU_MemorySpaceAttr:
+ EnumAttr<XeGPU_Dialect, XeGPU_MemorySpace, "memory_space"> {
+ let summary = [{Describe the location of data described by a `TensorDesc`:
+ Global device memory (`Global`) or Shared local memory (`SLM`).}];
+}
+
class XeGPU_TensorDescAttr<string name, string attrMnemonic, list<Trait> traits = [],
string baseCppClass = "::mlir::Attribute">
: XeGPUAttr<name, attrMnemonic, traits, baseCppClass> {
@@ -42,8 +60,8 @@ def XeGPU_BlockTensorDescAttr: XeGPU_TensorDescAttr<"BlockTensorDesc", "block_td
}];
let parameters = (ins
- DefaultValuedParameter<
- "MemorySpaceAttr",
+ DefaultValuedEnumAttrParameter<
+ XeGPU_MemorySpaceAttr,
"MemorySpaceAttr::get($_ctxt, xegpu::MemorySpace::Global)",
"Data memory location">: $memory_space,
DefaultValuedParameter<
@@ -71,26 +89,6 @@ def XeGPU_BlockTensorDescAttr: XeGPU_TensorDescAttr<"BlockTensorDesc", "block_td
}];
}
-
-//===----------------------------------------------------------------------===//
-// XeGPU Memory Scope Enums.
-//===----------------------------------------------------------------------===//
-def XeGPU_MemorySpaceGlobal: I32EnumAttrCase<"Global", 0, "global">;
-def XeGPU_MemorySpaceShared: I32EnumAttrCase<"SLM", 3, "slm">;
-def XeGPU_MemorySpace: I32EnumAttr<"MemorySpace",
- "The address space of the memory the tensor descritor is created for",
- [XeGPU_MemorySpaceGlobal, XeGPU_MemorySpaceShared]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::xegpu";
-}
-
-def XeGPU_MemorySpaceAttr:
- EnumAttr<XeGPU_Dialect, XeGPU_MemorySpace, "memory_space"> {
- let summary = [{Describe the location of data described by a `TensorDesc`:
- Global device memory (`Global`) or Shared local memory (`SLM`).}];
- let assemblyFormat = "$value";
-}
-
//===----------------------------------------------------------------------===//
// XeGPU Cache Enums.
//===----------------------------------------------------------------------===//
@@ -112,7 +110,6 @@ def XeGPU_CachePolicyEnums : I32EnumAttr<"CachePolicy", "Cache policy",
def XeGPU_CacheHintAttr
: EnumAttr<XeGPU_Dialect, XeGPU_CachePolicyEnums, "cache_hint"> {
let summary = [{Describe the cache settings for prefetch/load/store operators}];
- let assemblyFormat = "`<` $value `>`";
}
def XeGPU_FenceScopeWorkgroup: I32EnumAttrCase<"Workgroup", 0, "workgroup">;
@@ -129,7 +126,6 @@ def XeGPU_FenceScopeAttr:
let summary = [{Describes the scope of fence.
"workgroup" means that the scope is within each work group.
"gpu" means the scope is across work groups within the gpu.}];
- let assemblyFormat = "$value";
}
//===----------------------------------------------------------------------===//
@@ -150,7 +146,6 @@ def XeGPU_LaneShuffleModeAttr:
"pack" gathers a lane's elements into consecutive logical
positions. "unpack" scatters them back out, strided by the
subgroup size.}];
- let assemblyFormat = "$value";
}
def DistributeLayoutAttr: AttrInterface<"DistributeLayoutAttr"> {
diff --git a/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUOps.td b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUOps.td
index ee165bdead276..2821b01e880d5 100644
--- a/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUOps.td
+++ b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUOps.td
@@ -1213,7 +1213,10 @@ def XeGPU_FenceOp: XeGPU_Op<"fence", []> {
}];
let arguments = (ins XeGPU_MemorySpaceAttr: $memory_kind,
XeGPU_FenceScopeAttr: $fence_scope);
- let assemblyFormat = [{`memory_kind` `=` $memory_kind `,` `fence_scope` `=` $fence_scope attr-dict}];
+ let assemblyFormat = [{
+ `memory_kind` `=` enum($memory_kind) `,`
+ `fence_scope` `=` enum($fence_scope) attr-dict
+ }];
let extraClassDeclaration = extraBaseClassDeclaration;
}
@@ -1571,7 +1574,7 @@ def XeGPU_LaneShuffleOp
}];
let assemblyFormat = [{
- $source $mode attr-dict `:` type($source)
+ $source enum($mode) attr-dict `:` type($source)
}];
let hasVerifier = 1;
diff --git a/mlir/include/mlir/IR/EnumAttr.td b/mlir/include/mlir/IR/EnumAttr.td
index b0704c81805c9..1c3a223dbb849 100644
--- a/mlir/include/mlir/IR/EnumAttr.td
+++ b/mlir/include/mlir/IR/EnumAttr.td
@@ -525,19 +525,19 @@ class EnumParameter<EnumInfo enumInfo>
// def MyEnumAttr : EnumAttr<MyDialect, MyEnum, "enum">;
// ```
//
-// By default, the assembly format of the attribute works best with operation
-// assembly formats. For example:
+// By default, the assembly format of the attribute wraps the symbolic value in
+// angle brackets. Use the `enum` directive to print only the symbolic value in
+// an operation assembly format. For example:
//
// ```
// def MyOp : Op<MyDialect, "my_op"> {
// let arguments = (ins MyEnumAttr:$enum);
-// let assemblyFormat = "$enum attr-dict";
+// let assemblyFormat = "enum($enum) attr-dict";
// }
// ```
//
// The op will appear in the IR as `my_dialect.my_op first`. However, the
-// generic format of the attribute will be `#my_dialect<"enum first">`. Override
-// the attribute's assembly format as required.
+// generic format of the attribute will be `#my_dialect.enum<first>`.
class EnumAttr<Dialect dialect, EnumInfo enumInfo, string name = "",
list <Trait> traits = []>
: AttrDef<dialect, enumInfo.className, traits> {
@@ -566,9 +566,8 @@ class EnumAttr<Dialect dialect, EnumInfo enumInfo, string name = "",
// If a mnemonic was provided, use it to generate a custom assembly format.
let mnemonic = name;
- // The default assembly format for enum attributes. Selected to best work with
- // operation assembly formats.
- let assemblyFormat = "$value";
+ // The default assembly format for enum attributes.
+ let assemblyFormat = "`<` $value `>`";
}
// An EnumAttr used as a parameter of another attribute or type. This retains
@@ -587,6 +586,13 @@ class OptionalEnumAttrParameter<EnumAttr enumAttr, string desc = "">
let defaultValue = cppStorageType # "()";
}
+// An EnumAttr parameter with a default value.
+class DefaultValuedEnumAttrParameter<EnumAttr enumAttr, string value,
+ string desc = "">
+ : EnumAttrParameter<enumAttr, desc> {
+ let defaultValue = value;
+}
+
// A property wrapping by a C++ enum. This class will automatically create bytecode
// serialization logic for the given enum, as well as arranging for parser and
// printer calls.
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index f4145967e94dc..a39b43c9803eb 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -611,26 +611,6 @@ OpFoldResult gpu::AllReduceOp::fold(FoldAdaptor /*adaptor*/) {
return nullptr;
}
-// TODO: Support optional custom attributes (without dialect prefix).
-static ParseResult parseAllReduceOperation(AsmParser &parser,
- AllReduceOperationAttr &attr) {
- StringRef enumStr;
- if (!parser.parseOptionalKeyword(&enumStr)) {
- std::optional<AllReduceOperation> op =
- gpu::symbolizeAllReduceOperation(enumStr);
- if (!op)
- return parser.emitError(parser.getCurrentLocation(), "invalid op kind");
- attr = AllReduceOperationAttr::get(parser.getContext(), *op);
- }
- return success();
-}
-
-static void printAllReduceOperation(AsmPrinter &printer, Operation *op,
- AllReduceOperationAttr attr) {
- if (attr)
- attr.print(printer);
-}
-
//===----------------------------------------------------------------------===//
// SubgroupReduceOp
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
index 8d7d5d3282327..28217741208d9 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
@@ -124,7 +124,7 @@
// func.func @test() {
// %scalar = memref.alloca() {acc.var_name = "n"} : memref<i32>
// %copyin = acc.copyin varPtr(%scalar : memref<i32>) -> memref<i32>
-// {dataClause = #acc<data_clause acc_copy>,
+// {dataClause = #acc.data_clause<acc_copy>,
// implicit = true, name = "n"}
// acc.kernels dataOperands(%copyin : memref<i32>) {
// %val = memref.load %copyin[] : memref<i32>
@@ -132,7 +132,7 @@
// }
// acc.copyout accPtr(%copyin : memref<i32>)
// to varPtr(%scalar : memref<i32>)
-// {dataClause = #acc<data_clause acc_copy>,
+// {dataClause = #acc.data_clause<acc_copy>,
// implicit = true, name = "n"}
// }
//
@@ -153,7 +153,7 @@
// %array = memref.alloca() {acc.var_name = "arr"} : memref<100xf32>
// %copyin = acc.copyin varPtr(%array : memref<100xf32>)
// -> memref<100xf32>
-// {dataClause = #acc<data_clause acc_copy>,
+// {dataClause = #acc.data_clause<acc_copy>,
// implicit = true, name = "arr"}
// acc.parallel dataOperands(%copyin : memref<100xf32>) {
// %c0 = arith.constant 0 : index
@@ -162,7 +162,7 @@
// }
// acc.copyout accPtr(%copyin : memref<100xf32>)
// to varPtr(%array : memref<100xf32>)
-// {dataClause = #acc<data_clause acc_copy>,
+// {dataClause = #acc.data_clause<acc_copy>,
// implicit = true, name = "arr"}
// }
//
@@ -175,7 +175,7 @@
// %c0 = arith.constant 0 : index
// %val = memref.load %array[%c0] : memref<100xf32>
// acc.yield
-// } attributes {defaultAttr = #acc<defaultvalue present>}
+// } attributes {defaultAttr = #acc.defaultvalue<present>}
// }
//
// After:
@@ -185,13 +185,13 @@
// -> memref<100xf32>
// {implicit = true, name = "arr"}
// acc.parallel dataOperands(%present : memref<100xf32>)
-// attributes {defaultAttr = #acc<defaultvalue present>} {
+// attributes {defaultAttr = #acc.defaultvalue<present>} {
// %c0 = arith.constant 0 : index
// %val = memref.load %present[%c0] : memref<100xf32>
// acc.yield
// }
// acc.delete accPtr(%present : memref<100xf32>)
-// {dataClause = #acc<data_clause acc_present>,
+// {dataClause = #acc.data_clause<acc_present>,
// implicit = true, name = "arr"}
// }
//
diff --git a/mlir/python/mlir/dialects/irdl.py b/mlir/python/mlir/dialects/irdl.py
index 1ec951b69b646..1cbb8c742359d 100644
--- a/mlir/python/mlir/dialects/irdl.py
+++ b/mlir/python/mlir/dialects/irdl.py
@@ -87,6 +87,7 @@ def attribute(
@register_attribute_builder("VariadicityArrayAttr")
def _variadicity_array_attr(x: Sequence[Variadicity], context) -> _ods_ir.Attribute:
+ variadicities = ", ".join(f"<{i}>" for i in x)
return _ods_ir.Attribute.parse(
- f"#irdl<variadicity_array [{', '.join(str(i) for i in x)}]>", context
+ f"#irdl<variadicity_array [{variadicities}]>", context
)
diff --git a/mlir/test/CAPI/irdl.c b/mlir/test/CAPI/irdl.c
index 4ee4bc0cc35b6..96177b21ac1e9 100644
--- a/mlir/test/CAPI/irdl.c
+++ b/mlir/test/CAPI/irdl.c
@@ -42,19 +42,19 @@ void testVariadicityAttributes(MlirContext ctx) {
MlirAttribute variadicitySingle =
mlirIRDLVariadicityAttrGet(ctx, mlirStringRefCreateFromCString("single"));
- // CHECK: #irdl<variadicity single>
+ // CHECK: #irdl.variadicity<single>
mlirAttributeDump(variadicitySingle);
MlirAttribute variadicityOptional = mlirIRDLVariadicityAttrGet(
ctx, mlirStringRefCreateFromCString("optional"));
- // CHECK: #irdl<variadicity optional>
+ // CHECK: #irdl.variadicity<optional>
mlirAttributeDump(variadicityOptional);
MlirAttribute variadicityVariadic = mlirIRDLVariadicityAttrGet(
ctx, mlirStringRefCreateFromCString("variadic"));
- // CHECK: #irdl<variadicity variadic>
+ // CHECK: #irdl.variadicity<variadic>
mlirAttributeDump(variadicityVariadic);
MlirAttribute variadicities[] = {variadicitySingle, variadicityOptional,
@@ -62,7 +62,7 @@ void testVariadicityAttributes(MlirContext ctx) {
MlirAttribute variadicityArray =
mlirIRDLVariadicityArrayAttrGet(ctx, 3, variadicities);
- // CHECK: #irdl<variadicity_array[single, optional, variadic]>
+ // CHECK: #irdl<variadicity_array[<single>, <optional>, <variadic>]>
mlirAttributeDump(variadicityArray);
}
diff --git a/mlir/test/Dialect/GPU/invalid.mlir b/mlir/test/Dialect/GPU/invalid.mlir
index 376957fe34109..1b6d3a41d2bd0 100644
--- a/mlir/test/Dialect/GPU/invalid.mlir
+++ b/mlir/test/Dialect/GPU/invalid.mlir
@@ -379,14 +379,14 @@ func.func @reduce_op_and_body(%arg0 : f32) {
%res = "gpu.all_reduce"(%arg0) ({
^bb(%lhs : f32, %rhs : f32):
"gpu.yield"(%lhs) : (f32) -> ()
- }) {op = #gpu<all_reduce_op add>} : (f32) -> (f32)
+ }) {op = #gpu.all_reduce_op<add>} : (f32) -> (f32)
return
}
// -----
func.func @reduce_invalid_op(%arg0 : f32) {
- // expected-error at +1 {{invalid op kind}}
+ // expected-error at +1 {{expected SSA operand}}
%res = gpu.all_reduce foo %arg0 {} : (f32) -> (f32)
return
}
@@ -604,7 +604,7 @@ func.func @reduce_incorrect_yield(%arg0 : f32) {
func.func @shuffle_mismatching_type(%arg0 : f32, %arg1 : i32, %arg2 : i32) {
// expected-error at +1 {{op failed to verify that all of {value, shuffleResult} have same type}}
- %shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = #gpu<shuffle_mode xor> } : (f32, i32, i32) -> (i32, i1)
+ %shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = #gpu.shuffle_mode<xor> } : (f32, i32, i32) -> (i32, i1)
return
}
diff --git a/mlir/test/Dialect/IRDL/variadics-error.irdl.mlir b/mlir/test/Dialect/IRDL/variadics-error.irdl.mlir
index 3e29dd837916b..a400ea0458612 100644
--- a/mlir/test/Dialect/IRDL/variadics-error.irdl.mlir
+++ b/mlir/test/Dialect/IRDL/variadics-error.irdl.mlir
@@ -5,7 +5,7 @@ irdl.dialect @errors {
%0 = irdl.is i32
// expected-error at +1 {{'irdl.operands' op the number of operands and their variadicities must be the same, but got 2 and 1 respectively}}
- "irdl.operands"(%0, %0) <{names = ["foo", "bar"], variadicity = #irdl<variadicity_array[single]>}> : (!irdl.attribute, !irdl.attribute) -> ()
+ "irdl.operands"(%0, %0) <{names = ["foo", "bar"], variadicity = #irdl<variadicity_array[<single>]>}> : (!irdl.attribute, !irdl.attribute) -> ()
}
}
@@ -16,7 +16,7 @@ irdl.dialect @errors {
%0 = irdl.is i32
// expected-error at +1 {{'irdl.operands' op the number of operands and their variadicities must be the same, but got 1 and 2 respectively}}
- "irdl.operands"(%0) <{names = ["foo"], variadicity = #irdl<variadicity_array[single, single]>}> : (!irdl.attribute) -> ()
+ "irdl.operands"(%0) <{names = ["foo"], variadicity = #irdl<variadicity_array[<single>, <single>]>}> : (!irdl.attribute) -> ()
}
}
@@ -27,7 +27,7 @@ irdl.dialect @errors {
%0 = irdl.is i32
// expected-error at +1 {{'irdl.results' op the number of results and their variadicities must be the same, but got 2 and 1 respectively}}
- "irdl.results"(%0, %0) <{names = ["foo", "bar"], variadicity = #irdl<variadicity_array[single]>}> : (!irdl.attribute, !irdl.attribute) -> ()
+ "irdl.results"(%0, %0) <{names = ["foo", "bar"], variadicity = #irdl<variadicity_array[<single>]>}> : (!irdl.attribute, !irdl.attribute) -> ()
}
}
@@ -38,6 +38,6 @@ irdl.dialect @errors {
%0 = irdl.is i32
// expected-error at +1 {{'irdl.results' op the number of results and their variadicities must be the same, but got 1 and 2 respectively}}
- "irdl.results"(%0) <{names = ["foo"], variadicity = #irdl<variadicity_array[single, single]>}> : (!irdl.attribute) -> ()
+ "irdl.results"(%0) <{names = ["foo"], variadicity = #irdl<variadicity_array[<single>, <single>]>}> : (!irdl.attribute) -> ()
}
}
diff --git a/mlir/test/Dialect/OpenACC/invalid.mlir b/mlir/test/Dialect/OpenACC/invalid.mlir
index 10ab1c453ddbd..eeb8e75c5c039 100644
--- a/mlir/test/Dialect/OpenACC/invalid.mlir
+++ b/mlir/test/Dialect/OpenACC/invalid.mlir
@@ -870,13 +870,13 @@ func.func @acc_loop_container() {
%value = memref.alloc() : memref<f32>
// expected-error @below {{no data clause modifiers are allowed}}
-%0 = acc.private varPtr(%value : memref<f32>) <{modifiers = #acc<data_clause_modifier zero>}> -> memref<f32>
+%0 = acc.private varPtr(%value : memref<f32>) <{modifiers = #acc.data_clause_modifier<zero>}> -> memref<f32>
// -----
%value = memref.alloc() : memref<f32>
// expected-error @below {{invalid data clause modifiers: readonly}}
-%0 = acc.create varPtr(%value : memref<f32>) <{modifiers = #acc<data_clause_modifier readonly,zero,capture,always>}> -> memref<f32>
+%0 = acc.create varPtr(%value : memref<f32>) <{modifiers = #acc.data_clause_modifier<readonly,zero,capture,always>}> -> memref<f32>
// -----
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir b/mlir/test/Dialect/OpenACC/ops.mlir
index 24668cf18c6d5..d4f49673646bb 100644
--- a/mlir/test/Dialect/OpenACC/ops.mlir
+++ b/mlir/test/Dialect/OpenACC/ops.mlir
@@ -1021,27 +1021,27 @@ func.func @testdataop(%a: memref<f32>, %b: memref<f32>, %c: memref<f32>) -> () {
// -----
func.func @testdataopmodifiers(%a: memref<f32>, %b: memref<f32>, %c: memref<f32>) -> () {
- %0 = acc.create varPtr(%a : memref<f32>) <{modifiers = #acc<data_clause_modifier capture,zero>}> -> memref<f32>
- %1 = acc.copyin varPtr(%b : memref<f32>) <{modifiers = #acc<data_clause_modifier readonly,capture,always>}> -> memref<f32>
- %2 = acc.copyin varPtr(%c : memref<f32>) <{modifiers = #acc<data_clause_modifier always>}> -> memref<f32>
- %3 = acc.create varPtr(%c : memref<f32>) <{modifiers = #acc<data_clause_modifier always>}> -> memref<f32>
+ %0 = acc.create varPtr(%a : memref<f32>) <{modifiers = #acc.data_clause_modifier<capture,zero>}> -> memref<f32>
+ %1 = acc.copyin varPtr(%b : memref<f32>) <{modifiers = #acc.data_clause_modifier<readonly,capture,always>}> -> memref<f32>
+ %2 = acc.copyin varPtr(%c : memref<f32>) <{modifiers = #acc.data_clause_modifier<always>}> -> memref<f32>
+ %3 = acc.create varPtr(%c : memref<f32>) <{modifiers = #acc.data_clause_modifier<always>}> -> memref<f32>
acc.data dataOperands(%0, %1, %2, %3 : memref<f32>, memref<f32>, memref<f32>, memref<f32>) {
}
- acc.copyout accPtr(%0 : memref<f32>) to varPtr(%a : memref<f32>) <{modifiers = #acc<data_clause_modifier zero,capture,always>}>
- acc.delete accPtr(%2 : memref<f32>) <{modifiers = #acc<data_clause_modifier always>}>
- acc.copyout accPtr(%3 : memref<f32>) to varPtr(%c : memref<f32>) <{modifiers = #acc<data_clause_modifier always>}>
+ acc.copyout accPtr(%0 : memref<f32>) to varPtr(%a : memref<f32>) <{modifiers = #acc.data_clause_modifier<zero,capture,always>}>
+ acc.delete accPtr(%2 : memref<f32>) <{modifiers = #acc.data_clause_modifier<always>}>
+ acc.copyout accPtr(%3 : memref<f32>) to varPtr(%c : memref<f32>) <{modifiers = #acc.data_clause_modifier<always>}>
func.return
}
// CHECK: func @testdataopmodifiers(%[[ARGA:.*]]: memref<f32>, %[[ARGB:.*]]: memref<f32>, %[[ARGC:.*]]: memref<f32>) {
-// CHECK: %[[CREATEA:.*]] = acc.create varPtr(%[[ARGA]] : memref<f32>) <modifiers = "zero,capture"> -> memref<f32>
-// CHECK: %[[COPYINB:.*]] = acc.copyin varPtr(%[[ARGB]] : memref<f32>) <modifiers = "always,readonly,capture"> -> memref<f32>
-// CHECK: %[[COPYINC:.*]] = acc.copyin varPtr(%[[ARGC]] : memref<f32>) <modifiers = "always"> -> memref<f32>
-// CHECK: %[[CREATEC:.*]] = acc.create varPtr(%[[ARGC]] : memref<f32>) <modifiers = "always"> -> memref<f32>
-// CHECK: acc.copyout accPtr(%[[CREATEA]] : memref<f32>) to varPtr(%[[ARGA]] : memref<f32>) <modifiers = "always,zero,capture">
-// CHECK: acc.delete accPtr(%[[COPYINC]] : memref<f32>) <modifiers = "always">
-// CHECK: acc.copyout accPtr(%[[CREATEC]] : memref<f32>) to varPtr(%[[ARGC]] : memref<f32>) <modifiers = "always">
+// CHECK: %[[CREATEA:.*]] = acc.create varPtr(%[[ARGA]] : memref<f32>) <modifiers = <zero,capture>> -> memref<f32>
+// CHECK: %[[COPYINB:.*]] = acc.copyin varPtr(%[[ARGB]] : memref<f32>) <modifiers = <always,readonly,capture>> -> memref<f32>
+// CHECK: %[[COPYINC:.*]] = acc.copyin varPtr(%[[ARGC]] : memref<f32>) <modifiers = <always>> -> memref<f32>
+// CHECK: %[[CREATEC:.*]] = acc.create varPtr(%[[ARGC]] : memref<f32>) <modifiers = <always>> -> memref<f32>
+// CHECK: acc.copyout accPtr(%[[CREATEA]] : memref<f32>) to varPtr(%[[ARGA]] : memref<f32>) <modifiers = <always,zero,capture>>
+// CHECK: acc.delete accPtr(%[[COPYINC]] : memref<f32>) <modifiers = <always>>
+// CHECK: acc.copyout accPtr(%[[CREATEC]] : memref<f32>) to varPtr(%[[ARGC]] : memref<f32>) <modifiers = <always>>
// -----
diff --git a/mlir/test/Dialect/OpenMP/invalid.mlir b/mlir/test/Dialect/OpenMP/invalid.mlir
index 5d47b78dce48e..73f63e13f90f6 100644
--- a/mlir/test/Dialect/OpenMP/invalid.mlir
+++ b/mlir/test/Dialect/OpenMP/invalid.mlir
@@ -2339,7 +2339,7 @@ func.func @omp_task_depend_iterated_no_vars(%data_var: memref<i32>) {
// expected-error @below {{op unexpected depend iterated values}}
"omp.task"() ({
"omp.terminator"() : () -> ()
- }) {depend_iterated_kinds = [#omp<clause_task_depend(taskdependin)>], operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>} : () -> ()
+ }) {depend_iterated_kinds = [#omp.clause_task_depend<taskdependin>], operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>} : () -> ()
"func.return"() : () -> ()
}
@@ -3277,7 +3277,7 @@ func.func @omp_target_depend(%data_var: memref<i32>) {
// expected-error @below {{op expected as many depend values as depend variables}}
"omp.target"(%data_var) ({
"omp.terminator"() : () -> ()
- }) {kernel_type = #omp<kernel_type(generic)>, depend_kinds = [], operandSegmentSizes = array<i32: 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>} : (memref<i32>) -> ()
+ }) {kernel_type = #omp.kernel_type<generic>, depend_kinds = [], operandSegmentSizes = array<i32: 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>} : (memref<i32>) -> ()
"func.return"() : () -> ()
}
@@ -3923,7 +3923,7 @@ func.func @target_private_count_mismatch(%arg0: !llvm.ptr) {
// expected-error @below {{inconsistent number of private variables and privatizer op symbols, private vars: 1 vs. privatizer op symbols: 2}}
"omp.target"(%arg0) <{operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0>,
private_syms = [@x.privatizer, @y.privatizer],
- kernel_type = #omp<kernel_type(generic)>}> ({
+ kernel_type = #omp.kernel_type<generic>}> ({
^bb0(%arg1 : !llvm.ptr):
omp.terminator
}) : (!llvm.ptr) -> ()
@@ -4774,7 +4774,7 @@ func.func @omp_wsloop_linear_modifiers_mismatch(%lb : index, %ub : index, %step
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {linear_modifiers = [#omp<linear_modifier(val)>, #omp<linear_modifier(val)>],
+ }) {linear_modifiers = [#omp.linear_modifier<val>, #omp.linear_modifier<val>],
operandSegmentSizes = array<i32: 0, 0, 1, 1, 0, 0, 0>} : (memref<i32>, i32) -> ()
return
}
@@ -4788,7 +4788,7 @@ func.func @omp_simd_linear_modifiers_mismatch(%lb : index, %ub : index, %step :
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {linear_modifiers = [#omp<linear_modifier(val)>, #omp<linear_modifier(val)>],
+ }) {linear_modifiers = [#omp.linear_modifier<val>, #omp.linear_modifier<val>],
operandSegmentSizes = array<i32: 0, 0, 1, 1, 0, 0, 0>} : (memref<i32>, i32) -> ()
return
}
@@ -4797,7 +4797,7 @@ func.func @omp_simd_linear_modifiers_mismatch(%lb : index, %ub : index, %step :
func.func @omp_declare_simd_linear_modifiers_mismatch(%iv : i32, %step : i32) {
// expected-error @below {{'omp.declare_simd' op expected as many linear modifiers as linear variables}}
- "omp.declare_simd"(%iv, %step) <{linear_modifiers = [#omp<linear_modifier(val)>, #omp<linear_modifier(ref)>], operandSegmentSizes = array<i32: 0, 1, 1, 0>}> : (i32, i32) -> ()
+ "omp.declare_simd"(%iv, %step) <{linear_modifiers = [#omp.linear_modifier<val>, #omp.linear_modifier<ref>], operandSegmentSizes = array<i32: 0, 1, 1, 0>}> : (i32, i32) -> ()
return
}
diff --git a/mlir/test/Dialect/OpenMP/ops.mlir b/mlir/test/Dialect/OpenMP/ops.mlir
index a0aff2cdfda0e..92cc73f59ae58 100644
--- a/mlir/test/Dialect/OpenMP/ops.mlir
+++ b/mlir/test/Dialect/OpenMP/ops.mlir
@@ -116,7 +116,7 @@ func.func @omp_parallel(%data_var : memref<i32>, %if_cond : i1, %num_threads : i
}) {operandSegmentSizes = array<i32: 0,0,1,1,0,0>} : (i1, i32) -> ()
omp.terminator
- }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0>, proc_bind_kind = #omp<procbindkind spread>} : (i1, i32) -> ()
+ }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0>, proc_bind_kind = #omp.procbindkind<spread>} : (i1, i32) -> ()
// CHECK: omp.parallel
omp.parallel {
@@ -484,7 +484,7 @@ func.func @omp_wsloop(%lb : index, %ub : index, %step : index, %data_var : memre
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0,0>, schedule_kind = #omp<schedulekind static>,
+ }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0,0>, schedule_kind = #omp.schedulekind<static>,
linear_var_types = [i32]} : (memref<i32>, i32) -> ()
// CHECK: omp.wsloop linear(%{{.*}} : memref<i32> = %{{.*}} : i32, %{{.*}} : memref<i32> = %{{.*}} : i32) linear_var_types([i32, i32]) schedule(static) {
@@ -493,7 +493,7 @@ func.func @omp_wsloop(%lb : index, %ub : index, %step : index, %data_var : memre
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {operandSegmentSizes = array<i32: 0,0,2,2,0,0,0>, schedule_kind = #omp<schedulekind static>,
+ }) {operandSegmentSizes = array<i32: 0,0,2,2,0,0,0>, schedule_kind = #omp.schedulekind<static>,
linear_var_types = [i32,i32]} :
(memref<i32>, memref<i32>, i32, i32) -> ()
@@ -503,7 +503,7 @@ func.func @omp_wsloop(%lb : index, %ub : index, %step : index, %data_var : memre
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0,1>, schedule_kind = #omp<schedulekind dynamic>, ordered = 2,
+ }) {operandSegmentSizes = array<i32: 0,0,1,1,0,0,1>, schedule_kind = #omp.schedulekind<dynamic>, ordered = 2,
linear_var_types = [i32]} : (memref<i32>, i32, i32) -> ()
// CHECK: omp.wsloop nowait schedule(auto) {
@@ -512,7 +512,7 @@ func.func @omp_wsloop(%lb : index, %ub : index, %step : index, %data_var : memre
omp.loop_nest (%iv) : index = (%lb) to (%ub) step (%step) {
omp.yield
}
- }) {operandSegmentSizes = array<i32: 0,0,0,0,0,0,0>, nowait, schedule_kind = #omp<schedulekind auto>} :
+ }) {operandSegmentSizes = array<i32: 0,0,0,0,0,0,0>, nowait, schedule_kind = #omp.schedulekind<auto>} :
() -> ()
// CHECK: omp.wsloop {
@@ -895,7 +895,7 @@ func.func @omp_target(%if_cond : i1, %device : si32, %num_threads : i32, %devic
"omp.target"(%device, %if_cond, %num_threads) ({
// CHECK: omp.terminator
omp.terminator
- }) {kernel_type = #omp<kernel_type(generic)>, nowait, operandSegmentSizes = array<i32: 0,0,0,0,1,0,0,0,1,0,0,0,0,0,1>} : ( si32, i1, i32 ) -> ()
+ }) {kernel_type = #omp.kernel_type<generic>, nowait, operandSegmentSizes = array<i32: 0,0,0,0,1,0,0,0,1,0,0,0,0,0,1>} : ( si32, i1, i32 ) -> ()
// Test with optional map clause.
// CHECK: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[VAL_1:.*]] : memref<?xi32>, tensor<?xi32>) map_clauses(always, to) capture(ByRef) name("") -> memref<?xi32>
@@ -3204,16 +3204,16 @@ func.func @omp_taskloop(%lb: i32, %ub: i32, %step: i32) -> () {
}
// CHECK: func.func @omp_requires_one
-// CHECK-SAME: omp.requires = #omp<clause_requires reverse_offload>
+// CHECK-SAME: omp.requires = #omp.clause_requires<reverse_offload>
func.func @omp_requires_one() -> ()
- attributes {omp.requires = #omp<clause_requires reverse_offload>} {
+ attributes {omp.requires = #omp.clause_requires<reverse_offload>} {
return
}
// CHECK: func.func @omp_requires_multiple
-// CHECK-SAME: omp.requires = #omp<clause_requires unified_address|dynamic_allocators>
+// CHECK-SAME: omp.requires = #omp.clause_requires<unified_address|dynamic_allocators>
func.func @omp_requires_multiple() -> ()
- attributes {omp.requires = #omp<clause_requires unified_address|dynamic_allocators>} {
+ attributes {omp.requires = #omp.clause_requires<unified_address|dynamic_allocators>} {
return
}
@@ -4425,14 +4425,14 @@ func.func @omp_target_map_iterated(%lb : index, %ub : index, %step : index,
// CHECK-LABEL: func.func @omp_interop_init
func.func @omp_interop_init(%obj : !llvm.ptr, %device : i32) -> () {
- // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp<interop_type(target)>])
- omp.interop.init %obj : !llvm.ptr interop_types([#omp<interop_type(target)>])
+ // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp.interop_type<target>])
+ omp.interop.init %obj : !llvm.ptr interop_types([#omp.interop_type<target>])
- // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp<interop_type(targetsync)>])
- omp.interop.init %obj : !llvm.ptr interop_types([#omp<interop_type(targetsync)>])
+ // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp.interop_type<targetsync>])
+ omp.interop.init %obj : !llvm.ptr interop_types([#omp.interop_type<targetsync>])
- // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp<interop_type(targetsync)>, #omp<interop_type(target)>]) prefer_type([1, 6]) device(%{{.*}} : i32) nowait
- omp.interop.init %obj : !llvm.ptr interop_types([#omp<interop_type(targetsync)>, #omp<interop_type(target)>]) prefer_type([1, 6]) device(%device : i32) nowait
+ // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp.interop_type<targetsync>, #omp.interop_type<target>]) prefer_type([1, 6]) device(%{{.*}} : i32) nowait
+ omp.interop.init %obj : !llvm.ptr interop_types([#omp.interop_type<targetsync>, #omp.interop_type<target>]) prefer_type([1, 6]) device(%device : i32) nowait
return
}
@@ -4464,8 +4464,8 @@ func.func @omp_interop_destroy(%obj : !llvm.ptr, %device : i32) -> () {
// CHECK-LABEL: func.func @omp_interop_depend
func.func @omp_interop_depend(%obj : !llvm.ptr, %dep : !llvm.ptr) -> () {
- // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp<interop_type(targetsync)>]) depend(taskdependinout -> %{{.*}} : !llvm.ptr)
- omp.interop.init %obj : !llvm.ptr interop_types([#omp<interop_type(targetsync)>]) depend(taskdependinout -> %dep : !llvm.ptr)
+ // CHECK: omp.interop.init %{{.*}} : !llvm.ptr interop_types([#omp.interop_type<targetsync>]) depend(taskdependinout -> %{{.*}} : !llvm.ptr)
+ omp.interop.init %obj : !llvm.ptr interop_types([#omp.interop_type<targetsync>]) depend(taskdependinout -> %dep : !llvm.ptr)
// CHECK: omp.interop.use %{{.*}} : !llvm.ptr depend(taskdependin -> %{{.*}} : !llvm.ptr)
omp.interop.use %obj : !llvm.ptr depend(taskdependin -> %dep : !llvm.ptr)
diff --git a/mlir/test/IR/array-of-attr.mlir b/mlir/test/IR/array-of-attr.mlir
index 2c7b5009f74a2..bbf2e53616ad6 100644
--- a/mlir/test/IR/array-of-attr.mlir
+++ b/mlir/test/IR/array-of-attr.mlir
@@ -6,8 +6,8 @@ test.array_of_attr_op
a = [begin 0 : index end, begin 2 : index end],
// CHECK-SAME: [0, 1, -42, 42]
b = [0, 1, -42, 42],
- // CHECK-SAME: [a, b, b, a, "+"]
- c = [a, b, b, a, "+"]
+ // CHECK-SAME: [<a>, <b>, <b>, <a>, <"+">]
+ c = [<a>, <b>, <b>, <a>, <"+">]
// CHECK: test.array_of_attr_op
// CHECK-SAME: a = [], b = [], c = []
diff --git a/mlir/test/IR/enum-attr-invalid.mlir b/mlir/test/IR/enum-attr-invalid.mlir
index 0702da66c20b3..fe401060125bf 100644
--- a/mlir/test/IR/enum-attr-invalid.mlir
+++ b/mlir/test/IR/enum-attr-invalid.mlir
@@ -2,7 +2,7 @@
func.func @test_invalid_enum_case() -> () {
// expected-error at +1 {{expected string or keyword containing one of the following enum values for attribute 'value' [first, second, third]}}
- test.op_with_enum #test<enum fourth>
+ test.op_with_enum #test.enum<fourth>
}
// -----
diff --git a/mlir/test/IR/enum-attr-roundtrip.mlir b/mlir/test/IR/enum-attr-roundtrip.mlir
index 2ce5680226765..9b4b8de898942 100644
--- a/mlir/test/IR/enum-attr-roundtrip.mlir
+++ b/mlir/test/IR/enum-attr-roundtrip.mlir
@@ -2,12 +2,12 @@
// CHECK-LABEL: @test_enum_attr_roundtrip
func.func @test_enum_attr_roundtrip() -> () {
- // CHECK: value = #test<enum first>
- "test.op"() {value = #test<enum first>} : () -> ()
- // CHECK: value = #test<enum second>
- "test.op"() {value = #test<enum second>} : () -> ()
- // CHECK: value = #test<enum third>
- "test.op"() {value = #test<enum third>} : () -> ()
+ // CHECK: value = #test.enum<first>
+ "test.op"() {value = #test.enum<first>} : () -> ()
+ // CHECK: value = #test.enum<second>
+ "test.op"() {value = #test.enum<second>} : () -> ()
+ // CHECK: value = #test.enum<third>
+ "test.op"() {value = #test.enum<third>} : () -> ()
return
}
@@ -69,12 +69,12 @@ func.func @test_enum_prop() -> () {
// CHECK: test.op_with_enum_prop_attr_form <value = first>
test.op_with_enum_prop_attr_form <{value = 0 : i32}>
// CHECK: test.op_with_enum_prop_attr_form <value = second>
- test.op_with_enum_prop_attr_form <{value = #test<enum second>}>
+ test.op_with_enum_prop_attr_form <{value = #test.enum<second>}>
// CHECK: test.op_with_enum_prop_attr_form_always <value = first>
- test.op_with_enum_prop_attr_form_always <{value = #test<enum first>}>
+ test.op_with_enum_prop_attr_form_always <{value = #test.enum<first>}>
// CHECK: test.op_with_enum_prop_attr_form_always <value = second>
- test.op_with_enum_prop_attr_form_always <{value = #test<enum second>}>
+ test.op_with_enum_prop_attr_form_always <{value = #test.enum<second>}>
return
}
diff --git a/mlir/test/IR/properties.mlir b/mlir/test/IR/properties.mlir
index 0506ec5fccac2..9c0c8f001fd29 100644
--- a/mlir/test/IR/properties.mlir
+++ b/mlir/test/IR/properties.mlir
@@ -70,10 +70,10 @@ test.with_custom_prop_dict <attr = 11, unit_attr = unit, unit = unit, prop = 12>
// Inherent attributes use their custom assembly printer in the key-value
// spelling. Optional enum attributes compile and are omitted when absent.
-// CHECK: test.with_custom_attr_prop_dict <prop = 9, attr = first>
-test.with_custom_attr_prop_dict <attr = first, prop = 9>
-// CHECK: test.with_custom_attr_prop_dict <prop = 10, attr = first, optionalAttr = second>
-test.with_custom_attr_prop_dict <optionalAttr = second, prop = 10, attr = first>
+// CHECK: test.with_custom_attr_prop_dict <prop = 9, attr = <first>>
+test.with_custom_attr_prop_dict <attr = <first>, prop = 9>
+// CHECK: test.with_custom_attr_prop_dict <prop = 10, attr = <first>, optionalAttr = <second>>
+test.with_custom_attr_prop_dict <optionalAttr = <second>, prop = 10, attr = <first>>
// Properties bound elsewhere in the assembly format are excluded from the
// key-value list.
diff --git a/mlir/test/Target/LLVMIR/omptarget-atomic-capture-control-options.mlir b/mlir/test/Target/LLVMIR/omptarget-atomic-capture-control-options.mlir
index 1145a9800deeb..0d9efc51a2a49 100644
--- a/mlir/test/Target/LLVMIR/omptarget-atomic-capture-control-options.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-atomic-capture-control-options.mlir
@@ -2,7 +2,7 @@
// CHECK: atomicrmw add ptr %loadgep_, i32 1 monotonic, align 4, !amdgpu.no.remote.memory !{{.*}}
-module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<1> = dense<64> : vector<4xi64>, !llvm.ptr<2> = dense<32> : vector<4xi64>, !llvm.ptr<3> = dense<32> : vector<4xi64>, !llvm.ptr<4> = dense<64> : vector<4xi64>, !llvm.ptr<5> = dense<32> : vector<4xi64>, !llvm.ptr<6> = dense<32> : vector<4xi64>, !llvm.ptr<7> = dense<[160, 256, 256, 32]> : vector<4xi64>, !llvm.ptr<8> = dense<[128, 128, 128, 48]> : vector<4xi64>, !llvm.ptr<9> = dense<[192, 256, 256, 32]> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.legal_int_widths" = array<i32: 32, 64>, "dlti.stack_alignment" = 32 : i64, "dlti.alloca_memory_space" = 5 : ui64, "dlti.global_memory_space" = 1 : ui64>, fir.atomic_fine_grained_memory, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "generic-hsa", llvm.data_layout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9", llvm.target_triple = "amdgcn-amd-amdhsa", omp.flags = #omp.flags<openmp_device_version = 31>, omp.is_gpu = true, omp.is_target_device = true, omp.requires = #omp<clause_requires none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
+module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<1> = dense<64> : vector<4xi64>, !llvm.ptr<2> = dense<32> : vector<4xi64>, !llvm.ptr<3> = dense<32> : vector<4xi64>, !llvm.ptr<4> = dense<64> : vector<4xi64>, !llvm.ptr<5> = dense<32> : vector<4xi64>, !llvm.ptr<6> = dense<32> : vector<4xi64>, !llvm.ptr<7> = dense<[160, 256, 256, 32]> : vector<4xi64>, !llvm.ptr<8> = dense<[128, 128, 128, 48]> : vector<4xi64>, !llvm.ptr<9> = dense<[192, 256, 256, 32]> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.legal_int_widths" = array<i32: 32, 64>, "dlti.stack_alignment" = 32 : i64, "dlti.alloca_memory_space" = 5 : ui64, "dlti.global_memory_space" = 1 : ui64>, fir.atomic_fine_grained_memory, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "generic-hsa", llvm.data_layout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9", llvm.target_triple = "amdgcn-amd-amdhsa", omp.flags = #omp.flags<openmp_device_version = 31>, omp.is_gpu = true, omp.is_target_device = true, omp.requires = #omp.clause_requires<none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
llvm.func @_QQmain() attributes {fir.bindc_name = "TEST", omp.declare_target = #omp.declaretarget<device_type = host, capture_clause = to>, target_cpu = "generic-hsa"} {
%0 = llvm.mlir.constant(1 : i64) : i64
%1 = llvm.alloca %0 x i32 {bindc_name = "threads"} : (i64) -> !llvm.ptr<5>
diff --git a/mlir/test/Target/LLVMIR/omptarget-atomic-update-control-options.mlir b/mlir/test/Target/LLVMIR/omptarget-atomic-update-control-options.mlir
index bf6c36cf768a1..ccea511a214c7 100644
--- a/mlir/test/Target/LLVMIR/omptarget-atomic-update-control-options.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-atomic-update-control-options.mlir
@@ -2,7 +2,7 @@
// CHECK: atomicrmw add ptr %loadgep_, i32 1 monotonic, align 4, !amdgpu.ignore.denormal.mode !{{.*}}, !amdgpu.no.fine.grained.memory !{{.*}}, !amdgpu.no.remote.memory !{{.*}}
-module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<1> = dense<64> : vector<4xi64>, !llvm.ptr<2> = dense<32> : vector<4xi64>, !llvm.ptr<3> = dense<32> : vector<4xi64>, !llvm.ptr<4> = dense<64> : vector<4xi64>, !llvm.ptr<5> = dense<32> : vector<4xi64>, !llvm.ptr<6> = dense<32> : vector<4xi64>, !llvm.ptr<7> = dense<[160, 256, 256, 32]> : vector<4xi64>, !llvm.ptr<8> = dense<[128, 128, 128, 48]> : vector<4xi64>, !llvm.ptr<9> = dense<[192, 256, 256, 32]> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.legal_int_widths" = array<i32: 32, 64>, "dlti.stack_alignment" = 32 : i64, "dlti.alloca_memory_space" = 5 : ui64, "dlti.global_memory_space" = 1 : ui64>, fir.atomic_ignore_denormal_mode, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "generic-hsa", llvm.data_layout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9", llvm.target_triple = "amdgcn-amd-amdhsa", omp.flags = #omp.flags<openmp_device_version = 31>, omp.is_gpu = true, omp.is_target_device = true, omp.requires = #omp<clause_requires none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
+module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<1> = dense<64> : vector<4xi64>, !llvm.ptr<2> = dense<32> : vector<4xi64>, !llvm.ptr<3> = dense<32> : vector<4xi64>, !llvm.ptr<4> = dense<64> : vector<4xi64>, !llvm.ptr<5> = dense<32> : vector<4xi64>, !llvm.ptr<6> = dense<32> : vector<4xi64>, !llvm.ptr<7> = dense<[160, 256, 256, 32]> : vector<4xi64>, !llvm.ptr<8> = dense<[128, 128, 128, 48]> : vector<4xi64>, !llvm.ptr<9> = dense<[192, 256, 256, 32]> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.legal_int_widths" = array<i32: 32, 64>, "dlti.stack_alignment" = 32 : i64, "dlti.alloca_memory_space" = 5 : ui64, "dlti.global_memory_space" = 1 : ui64>, fir.atomic_ignore_denormal_mode, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "generic-hsa", llvm.data_layout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9", llvm.target_triple = "amdgcn-amd-amdhsa", omp.flags = #omp.flags<openmp_device_version = 31>, omp.is_gpu = true, omp.is_target_device = true, omp.requires = #omp.clause_requires<none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
llvm.func @_QQmain() attributes {fir.bindc_name = "TEST", omp.declare_target = #omp.declaretarget<device_type = host, capture_clause = to>, target_cpu = "generic-hsa"} {
%0 = llvm.mlir.constant(1 : i64) : i64
%1 = llvm.alloca %0 x i32 {bindc_name = "threads"} : (i64) -> !llvm.ptr<5>
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device-2.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device-2.mlir
index 3b6243c160d47..0ab70b1003f9b 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device-2.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device-2.mlir
@@ -4,7 +4,7 @@
// for link and to clauses. And verifies we continue to make the correct replacement accesses
// within the target region.
-module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_device = true, omp.requires = #omp<clause_requires unified_shared_memory>} {
+module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_device = true, omp.requires = #omp.clause_requires<unified_shared_memory>} {
// CHECK-DAG: @_QMtest_0Evar_to_usm_decl_tgt_ref_ptr = weak global ptr null, align 8
llvm.mlir.global external @_QMtest_0Evar_to_usm() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = any, capture_clause = to>} : i32 {
%0 = llvm.mlir.constant(1 : i32) : i32
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-module-rewrite-device.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-module-rewrite-device.mlir
index 0980f19f3a9c6..ab19b8eabe8d9 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-module-rewrite-device.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-module-rewrite-device.mlir
@@ -219,7 +219,7 @@ module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_devic
// region and indirectly inside of a declare target function invoked from that
// region, and both use-sites must be rewritten.
-module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_device = true, omp.requires = #omp<clause_requires unified_shared_memory>} {
+module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_device = true, omp.requires = #omp.clause_requires<unified_shared_memory>} {
// CHECK-DAG: @_QMtest_0Evar_to_usm_decl_tgt_ref_ptr = weak global ptr null, align 8
llvm.mlir.global external @_QMtest_0Evar_to_usm() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = any, capture_clause = to>} : i32 {
%0 = llvm.mlir.constant(1 : i32) : i32
diff --git a/mlir/test/Target/LLVMIR/omptarget-host-ref-semantics.mlir b/mlir/test/Target/LLVMIR/omptarget-host-ref-semantics.mlir
index 3c005001faa0a..d76fc096791ef 100644
--- a/mlir/test/Target/LLVMIR/omptarget-host-ref-semantics.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-host-ref-semantics.mlir
@@ -3,7 +3,7 @@
// Tests that we correctly lower the different variations of reference pointer
// and attach semantics.
-module attributes {omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp<clause_requires none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 61>} {
+module attributes {omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp.clause_requires<none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 61>} {
llvm.func @attach_always_(%arg0: !llvm.ptr, %arg1: !llvm.ptr) {
%map1 = omp.map.info var_ptr(%arg0 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) var_ptr_ptr(%arg1 : !llvm.ptr, i32) name("") -> !llvm.ptr
%map2 = omp.map.info var_ptr(%arg0 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(always, to) capture(ByRef) members(%map1 : [0] : !llvm.ptr) name("x") -> !llvm.ptr
diff --git a/mlir/test/Target/LLVMIR/omptarget-map-pointer-privatization.mlir b/mlir/test/Target/LLVMIR/omptarget-map-pointer-privatization.mlir
index abbb0efa600aa..7adbabdb7c8ac 100644
--- a/mlir/test/Target/LLVMIR/omptarget-map-pointer-privatization.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-map-pointer-privatization.mlir
@@ -7,7 +7,7 @@
// * the parent which now has the attach map type is still passed as a kernel
// argument (OMP_MAP_TARGET_PARAM), unlike normal attach maps.
-module attributes {omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp<clause_requires none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 52>} {
+module attributes {omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp.clause_requires<none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 52>} {
llvm.func @assumed_shape_array_priv_(%arg0: !llvm.ptr, %arg1: !llvm.ptr) {
%member = omp.map.info var_ptr(%arg0 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) var_ptr_ptr(%arg1 : !llvm.ptr, i32) name("") -> !llvm.ptr
%parent = omp.map.info var_ptr(%arg0 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(target_param, private, attach) capture(ByRef) var_ptr_ptr(%arg1 : !llvm.ptr, i32) members(%member : [0] : !llvm.ptr) name("arr_read_write") -> !llvm.ptr
diff --git a/mlir/test/Target/LLVMIR/omptarget-parallel-llvm-debug.mlir b/mlir/test/Target/LLVMIR/omptarget-parallel-llvm-debug.mlir
index fd4837df64d00..4977b59465289 100644
--- a/mlir/test/Target/LLVMIR/omptarget-parallel-llvm-debug.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-parallel-llvm-debug.mlir
@@ -10,7 +10,7 @@
#int_ty = #llvm.di_basic_type<tag = DW_TAG_base_type, name = "integer", sizeInBits = 32, encoding = DW_ATE_signed>
#var_x = #llvm.di_local_variable<scope = #sp, name = "x", file = #di_file, type = #int_ty>
#var_x1 = #llvm.di_local_variable<scope = #sp1, name = "x", file = #di_file, type = #int_ty>
-module attributes {dlti.dl_spec = #dlti.dl_spec<i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, i64 = dense<64> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<270> = dense<32> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, "dlti.endianness" = "little", "dlti.stack_alignment" = 128 : i64, "dlti.mangling_mode" = "e">, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "x86-64", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.ident = "flang version 21.0.0 (/home/haqadeer/work/src/aomp-llvm-project/flang 793f9220ab32f92fc3b253efec2e332c18090e53)", llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp<clause_requires none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 52>} {
+module attributes {dlti.dl_spec = #dlti.dl_spec<i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, i64 = dense<64> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr = dense<64> : vector<4xi64>, !llvm.ptr<270> = dense<32> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, "dlti.endianness" = "little", "dlti.stack_alignment" = 128 : i64, "dlti.mangling_mode" = "e">, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", fir.target_cpu = "x86-64", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.ident = "flang version 21.0.0 (/home/haqadeer/work/src/aomp-llvm-project/flang 793f9220ab32f92fc3b253efec2e332c18090e53)", llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp.clause_requires<none>, omp.target_triples = ["amdgcn-amd-amdhsa"], omp.version = #omp.version<version = 52>} {
llvm.func @_QQmain() attributes {fir.bindc_name = "test", frame_pointer = #llvm.framePointerKind<all>, target_cpu = "x86-64"} {
%0 = llvm.mlir.constant(1 : i64) : i64
%1 = llvm.alloca %0 x i32 {bindc_name = "x"} : (i64) -> !llvm.ptr
diff --git a/mlir/test/Target/LLVMIR/openmp-interop-depend.mlir b/mlir/test/Target/LLVMIR/openmp-interop-depend.mlir
index d362714add856..59f4ad8685902 100644
--- a/mlir/test/Target/LLVMIR/openmp-interop-depend.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-interop-depend.mlir
@@ -9,7 +9,7 @@
//--- init.mlir
// INIT: error: not yet implemented: Unhandled clause depend in omp.interop.init operation
llvm.func @test_interop_init_depend(%interop: !llvm.ptr, %dep: !llvm.ptr) {
- omp.interop.init %interop : !llvm.ptr interop_types([#omp<interop_type(targetsync)>]) depend(taskdependinout -> %dep : !llvm.ptr)
+ omp.interop.init %interop : !llvm.ptr interop_types([#omp.interop_type<targetsync>]) depend(taskdependinout -> %dep : !llvm.ptr)
llvm.return
}
diff --git a/mlir/test/Target/LLVMIR/openmp-interop.mlir b/mlir/test/Target/LLVMIR/openmp-interop.mlir
index 85cadd241e23d..7ef75ede3f414 100644
--- a/mlir/test/Target/LLVMIR/openmp-interop.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-interop.mlir
@@ -7,14 +7,14 @@
// CHECK: ret void
llvm.func @test_interop_init(%arg0: !llvm.ptr) {
%0 = llvm.mlir.constant(1 : i64) : i64
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(target)>])
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<target>])
llvm.return
}
// CHECK-LABEL: define void @test_interop_init_targetsync(
// CHECK: call void @__tgt_interop_init(ptr @{{.*}}, i32 %{{.*}}, ptr %{{.*}}, i32 2, i32 -1, i32 0, ptr null, i32 0)
llvm.func @test_interop_init_targetsync(%arg0: !llvm.ptr) {
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(targetsync)>])
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<targetsync>])
llvm.return
}
@@ -22,7 +22,7 @@ llvm.func @test_interop_init_targetsync(%arg0: !llvm.ptr) {
// CHECK: call void @__tgt_interop_init(ptr @{{.*}}, i32 %{{.*}}, ptr %{{.*}}, i32 1, i32 -1, i32 0, ptr null, i32 0)
// CHECK-NOT: call void @__tgt_interop_init
llvm.func @test_interop_init_both(%arg0: !llvm.ptr) {
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(targetsync)>, #omp<interop_type(target)>])
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<targetsync>, #omp.interop_type<target>])
llvm.return
}
@@ -47,7 +47,7 @@ llvm.func @test_interop_destroy(%arg0: !llvm.ptr) {
// CHECK-LABEL: define void @test_interop_init_nowait(
// CHECK: call void @__tgt_interop_init(ptr @{{.*}}, i32 %{{.*}}, ptr %{{.*}}, i32 1, i32 -1, i32 0, ptr null, i32 1)
llvm.func @test_interop_init_nowait(%arg0: !llvm.ptr) {
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(target)>]) nowait
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<target>]) nowait
llvm.return
}
@@ -61,7 +61,7 @@ llvm.func @test_interop_destroy_nowait(%arg0: !llvm.ptr) {
// CHECK-LABEL: define void @test_interop_init_device(
// CHECK: call void @__tgt_interop_init(ptr @{{.*}}, i32 %{{.*}}, ptr %{{.*}}, i32 1, i32 %{{.*}}, i32 0, ptr null, i32 0)
llvm.func @test_interop_init_device(%arg0: !llvm.ptr, %arg1: i32) {
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(target)>]) device(%arg1 : i32)
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<target>]) device(%arg1 : i32)
llvm.return
}
@@ -69,7 +69,7 @@ llvm.func @test_interop_init_device(%arg0: !llvm.ptr, %arg1: i32) {
// CHECK: %[[DEVICE:.*]] = trunc i64 %{{.*}} to i32
// CHECK: call void @__tgt_interop_init(ptr @{{.*}}, i32 %{{.*}}, ptr %{{.*}}, i32 1, i32 %[[DEVICE]], i32 0, ptr null, i32 0)
llvm.func @test_interop_init_device_i64(%arg0: !llvm.ptr, %arg1: i64) {
- omp.interop.init %arg0 : !llvm.ptr interop_types([#omp<interop_type(target)>]) device(%arg1 : i64)
+ omp.interop.init %arg0 : !llvm.ptr interop_types([#omp.interop_type<target>]) device(%arg1 : i64)
llvm.return
}
diff --git a/mlir/test/Target/LLVMIR/openmp-llvm.mlir b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
index 794c2ff577c9b..0c004565ab076 100644
--- a/mlir/test/Target/LLVMIR/openmp-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
@@ -4079,7 +4079,7 @@ llvm.func @omp_task_if(%boolexpr: i1) {
// -----
-module attributes {omp.requires = #omp<clause_requires reverse_offload|unified_shared_memory>} {}
+module attributes {omp.requires = #omp.clause_requires<reverse_offload|unified_shared_memory>} {}
// -----
diff --git a/mlir/test/Target/LLVMIR/openmp-reduction-task-modifier.mlir b/mlir/test/Target/LLVMIR/openmp-reduction-task-modifier.mlir
index a15b1ee701a4e..386cc8ead1580 100644
--- a/mlir/test/Target/LLVMIR/openmp-reduction-task-modifier.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-reduction-task-modifier.mlir
@@ -159,7 +159,7 @@ llvm.func @empty_sections_task_reduction(%x: !llvm.ptr) {
// reduction variables must not open a task-reduction scope.
llvm.func @parallel_task_mod_no_reductions() {
- "omp.parallel"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0>, reduction_mod = #omp<reduction_modifier(task)>}> ({
+ "omp.parallel"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0>, reduction_mod = #omp.reduction_modifier<task>}> ({
omp.terminator
}) : () -> ()
llvm.return
@@ -179,7 +179,7 @@ llvm.func @wsloop_task_mod_no_reductions() {
%lb = llvm.mlir.constant(1 : i32) : i32
%ub = llvm.mlir.constant(10 : i32) : i32
%step = llvm.mlir.constant(1 : i32) : i32
- "omp.wsloop"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0, 0>, reduction_mod = #omp<reduction_modifier(task)>}> ({
+ "omp.wsloop"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0, 0, 0, 0>, reduction_mod = #omp.reduction_modifier<task>}> ({
"omp.loop_nest"(%lb, %ub, %step) <{loop_inclusive}> ({
^bb0(%iv: i32):
"omp.yield"() : () -> ()
@@ -201,7 +201,7 @@ llvm.func @wsloop_task_mod_no_reductions() {
// empty-sections early return tested above.
llvm.func @sections_task_mod_no_reductions() {
- "omp.sections"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0>, reduction_mod = #omp<reduction_modifier(task)>}> ({
+ "omp.sections"() <{operandSegmentSizes = array<i32: 0, 0, 0, 0>, reduction_mod = #omp.reduction_modifier<task>}> ({
"omp.section"() ({
"omp.terminator"() : () -> ()
}) : () -> ()
diff --git a/mlir/test/Target/LLVMIR/openmp-teams-clauses-trunc-ext.mlir b/mlir/test/Target/LLVMIR/openmp-teams-clauses-trunc-ext.mlir
index 2524a1e078a1a..f13781960ebeb 100644
--- a/mlir/test/Target/LLVMIR/openmp-teams-clauses-trunc-ext.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-teams-clauses-trunc-ext.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
-module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.ident = "flang version 22.0.0 (git at github.com:mjklemm/llvm-project.git 5d9164c24a474793ab325116c5f782dce0577574)", llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp<clause_requires none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
+module attributes {dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.ident = "flang version 22.0.0 (git at github.com:mjklemm/llvm-project.git 5d9164c24a474793ab325116c5f782dce0577574)", llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_gpu = false, omp.is_target_device = false, omp.requires = #omp.clause_requires<none>, omp.target_triples = [], omp.version = #omp.version<version = 31>} {
omp.private {type = private} @_QFnum_threads_arg_2Ei_private_i32 : i32
omp.private {type = private} @_QFnum_threads_const_2Ei_private_i32 : i32
omp.private {type = private} @_QFnum_threads_arg_8Ei_private_i32 : i32
diff --git a/mlir/test/lib/Dialect/Test/TestAttrDefs.td b/mlir/test/lib/Dialect/Test/TestAttrDefs.td
index e153eeeed74dd..4a9c214234684 100644
--- a/mlir/test/lib/Dialect/Test/TestAttrDefs.td
+++ b/mlir/test/lib/Dialect/Test/TestAttrDefs.td
@@ -352,29 +352,28 @@ def TestArrayOfUglyAttrs : ArrayOfAttr<Test_Dialect, "ArrayOfUglyAttrs",
def TestArrayOfInts : ArrayOfAttr<Test_Dialect, "ArrayOfInts",
"array_of_ints", "int32_t">;
-def TestSimpleEnum8Attr : EnumAttr<Test_Dialect, TestSimpleEnum8, "simple_enum_8"> {
- let assemblyFormat = "`` $value";
-}
+def TestSimpleEnum8Attr
+ : EnumAttr<Test_Dialect, TestSimpleEnum8, "simple_enum_8">;
-def TestSimpleEnum16Attr : EnumAttr<Test_Dialect, TestSimpleEnum16, "simple_enum_16"> {
- let assemblyFormat = "`` $value";
-}
+def TestSimpleEnum16Attr
+ : EnumAttr<Test_Dialect, TestSimpleEnum16, "simple_enum_16">;
// An array of enum attributes.
-def TestSimpleEnumAttr : EnumAttr<Test_Dialect, TestSimpleEnum, "simple_enum"> {
- let assemblyFormat = "`` $value";
-}
+def TestSimpleEnumAttr
+ : EnumAttr<Test_Dialect, TestSimpleEnum, "simple_enum">;
def TestArrayOfEnums : ArrayOfAttr<Test_Dialect, "ArrayOfEnums",
"array_of_enums", "SimpleEnumAttr">;
def TestStructEnumAttr
: EnumAttr<Test_Dialect, TestStructEnum, "struct_enum_value"> {
- let assemblyFormat = "`<` $value `>`";
}
def TestAttrEnumAttrStruct : Test_Attr<"TestAttrEnumAttrStruct"> {
let parameters = (ins
EnumAttrParameter<TestStructEnumAttr>:$required,
- OptionalEnumAttrParameter<TestStructEnumAttr>:$optional
+ OptionalEnumAttrParameter<TestStructEnumAttr>:$optional,
+ DefaultValuedEnumAttrParameter<
+ TestStructEnumAttr,
+ "TestStructEnumAttr::get($_ctxt, TestStructEnum::First)">:$defaulted
);
let mnemonic = "enum_attr_struct";
let assemblyFormat = "`<` struct(params) `>`";
@@ -382,7 +381,6 @@ def TestAttrEnumAttrStruct : Test_Attr<"TestAttrEnumAttrStruct"> {
def TestStructBitEnumAttr
: EnumAttr<Test_Dialect, TestStructBitEnum, "struct_bit_enum_value"> {
- let assemblyFormat = "`<` $value `>`";
}
def TestAttrBitEnumAttrStruct : Test_Attr<"TestAttrBitEnumAttrStruct"> {
@@ -403,7 +401,6 @@ def TestCustomAnchor : Test_Attr<"TestCustomAnchor"> {
def Test_IteratorTypeEnum
: EnumAttr<Test_Dialect, IteratorType, "iterator_type"> {
- let assemblyFormat = "`<` $value `>`";
}
def Test_IteratorTypeArrayAttr
diff --git a/mlir/test/lib/Dialect/Test/TestOps.td b/mlir/test/lib/Dialect/Test/TestOps.td
index 3060de7ac2b4b..4c8d03feb1c64 100644
--- a/mlir/test/lib/Dialect/Test/TestOps.td
+++ b/mlir/test/lib/Dialect/Test/TestOps.td
@@ -491,13 +491,12 @@ def TestEnumAttr : EnumAttr<Test_Dialect, TestEnum, "enum">;
def TestPrettyEnumAttr
: EnumAttr<Test_Dialect, TestPrettyEnum, "pretty_enum"> {
- let assemblyFormat = "`<` $value `>`";
}
// Define an op that contains the enum attribute.
def OpWithEnum : TEST_Op<"op_with_enum"> {
let arguments = (ins TestEnumAttr:$value, OptionalAttr<AnyAttr>:$tag);
- let assemblyFormat = "$value (`tag` $tag^)? attr-dict";
+ let assemblyFormat = "enum($value) (`tag` $tag^)? attr-dict";
}
// Define an op that prints just the symbolic value of an enum attribute whose
@@ -585,7 +584,6 @@ def OpWithSimpleEnum16Prop : TEST_Op<"op_with_simple_enum_16_prop"> {
// Define the enum attribute.
def TestBitEnumAttr : EnumAttr<Test_Dialect, TestBitEnum, "bit_enum"> {
- let assemblyFormat = "`<` $value `>`";
}
// Define an op that contains the enum attribute.
@@ -596,7 +594,6 @@ def OpWithBitEnum : TEST_Op<"op_with_bit_enum"> {
def TestBitEnumVerticalBarAttr
: EnumAttr<Test_Dialect, TestBitEnumVerticalBar, "bit_enum_vbar"> {
- let assemblyFormat = "`<` $value `>`";
}
// Define an op that contains the enum attribute.
diff --git a/mlir/test/mlir-tblgen/attr-or-type-format-roundtrip.mlir b/mlir/test/mlir-tblgen/attr-or-type-format-roundtrip.mlir
index 62abd3a9f68fe..5a9ce63540880 100644
--- a/mlir/test/mlir-tblgen/attr-or-type-format-roundtrip.mlir
+++ b/mlir/test/mlir-tblgen/attr-or-type-format-roundtrip.mlir
@@ -33,10 +33,10 @@ attributes {
attr_12 = #test.attr_with_optional_enum<a>,
// CHECK: #test.attr_with_optional_enum<b>
attr_13 = #test.attr_with_optional_enum<b>,
- // CHECK: #test<simple_enum"+">
- attr_14 = #test<simple_enum "+">,
- // CHECK: #test<simple_enum"dash-separated-sentence">
- attr_15 = #test<simple_enum "dash-separated-sentence">,
+ // CHECK: #test.simple_enum<"+">
+ attr_14 = #test.simple_enum<"+">,
+ // CHECK: #test.simple_enum<"dash-separated-sentence">
+ attr_15 = #test.simple_enum<"dash-separated-sentence">,
// Test that ArrayRefParameter in non-last struct position is wrapped in
// brackets to avoid ambiguity with the struct-level comma (issue #156623).
// CHECK: #test.arr_struct<elements = [1, 2, 3], count = 42>
@@ -58,6 +58,10 @@ attributes {
// CHECK: #test.enum_attr_struct<required = first, optional = second>
attr_enum_attr_struct =
#test.enum_attr_struct<optional = second, required = first>,
+ // Default-valued EnumAttr parameters also use their underlying enum syntax.
+ // CHECK: #test.enum_attr_struct<required = first, defaulted = second>
+ attr_enum_attr_struct_defaulted =
+ #test.enum_attr_struct<defaulted = second, required = first>,
// CHECK: #test.enum_attr_struct<required = third>
attr_enum_attr_struct_optional =
#test.enum_attr_struct<required = third>,
diff --git a/mlir/test/mlir-tblgen/attr-or-type-format.td b/mlir/test/mlir-tblgen/attr-or-type-format.td
index 279b174768d87..6d8198f5f3c9d 100644
--- a/mlir/test/mlir-tblgen/attr-or-type-format.td
+++ b/mlir/test/mlir-tblgen/attr-or-type-format.td
@@ -237,7 +237,6 @@ def TestStructEnum : I32EnumAttr<"TestStructEnum", "TestStructEnumType", [
}
def EnumAttrC : EnumAttr<Test_Dialect, TestStructEnum, "EnumAttrC"> {
- let assemblyFormat = "`<` $value `>`";
}
// EnumAttr parameters in a struct use the underlying enum syntax rather than
@@ -267,7 +266,6 @@ def TestStructBitEnum : I32BitEnum<"TestStructBitEnum", "", [
}
def EnumAttrD : EnumAttr<Test_Dialect, TestStructBitEnum, "EnumAttrD"> {
- let assemblyFormat = "`<` $value `>`";
}
// A non-final comma-separated bit enum is bracketed so its commas are not
diff --git a/mlir/test/mlir-tblgen/enums-python-bindings.td b/mlir/test/mlir-tblgen/enums-python-bindings.td
index 74b9f51b0c2d6..487b60dbf91ce 100644
--- a/mlir/test/mlir-tblgen/enums-python-bindings.td
+++ b/mlir/test/mlir-tblgen/enums-python-bindings.td
@@ -108,8 +108,8 @@ def TestBitEnum_Attr : EnumAttr<Test_Dialect, TestBitEnum, "testbitenum">;
// CHECK: @register_attribute_builder("TestDialect.TestBitEnum_Attr")
// CHECK: def _testbitenum_attr(x, context):
-// CHECK: return _ods_ir.Attribute.parse(f'#TestDialect<testbitenum {str(x)}>', context=context)
+// CHECK: return _ods_ir.Attribute.parse(f'#TestDialect.testbitenum<{str(x)}>', context=context)
// CHECK: @register_attribute_builder("TestDialect.TestMyEnum_Attr")
// CHECK: def _testmyenum_attr(x, context):
-// CHECK: return _ods_ir.Attribute.parse(f'#TestDialect<enum {str(x)}>', context=context)
+// CHECK: return _ods_ir.Attribute.parse(f'#TestDialect.enum<{str(x)}>', context=context)
diff --git a/mlir/test/mlir-tblgen/op-format.td b/mlir/test/mlir-tblgen/op-format.td
index 31e1ef2297688..b198c92ec349b 100644
--- a/mlir/test/mlir-tblgen/op-format.td
+++ b/mlir/test/mlir-tblgen/op-format.td
@@ -134,7 +134,6 @@ def TestEnum : I32Enum<"TestEnum", "a test enum", [TestEnumCase0, TestEnumCase1]
def TestEnumAttr : EnumAttr<TestDialect, TestEnum, "enum">;
def TestPrettyEnumAttr : EnumAttr<TestDialect, TestEnum, "pretty_enum"> {
- let assemblyFormat = "`<` $value `>`";
}
def TestNonKeywordEnumCase : I32EnumCase<"NonKeyword", 0, "non-keyword">;
@@ -145,14 +144,13 @@ def TestNonKeywordEnum
}
def TestNonKeywordEnumAttr
: EnumAttr<TestDialect, TestNonKeywordEnum, "non_keyword_enum"> {
- let assemblyFormat = "`<` $value `>`";
}
// Default-valued optional attributes have a non-optional getter.
// CHECK-LABEL: DefaultOptionalEnumAttrOp::print
// CHECK: auto caseValue = getAttr();
-def DefaultOptionalEnumAttrOp : TestFormat_Op<"($attr^)? attr-dict">,
+def DefaultOptionalEnumAttrOp : TestFormat_Op<"(enum($attr)^)? attr-dict">,
Arguments<(ins DefaultValuedOptionalAttr<
TestEnumAttr, "::test::TestEnum::Case0">:$attr)>;
@@ -160,7 +158,7 @@ def DefaultOptionalEnumAttrOp : TestFormat_Op<"($attr^)? attr-dict">,
// CHECK: symbolizeTestEnum
// CHECK-LABEL: EnumAttrOp::print
// CHECK: stringifyTestEnum
-def EnumAttrOp : TestFormat_Op<"$attr attr-dict">,
+def EnumAttrOp : TestFormat_Op<"enum($attr) attr-dict">,
Arguments<(ins TestEnumAttr:$attr)>;
// Unquoted bit enums use a separator-aware attribute parser instead of the
diff --git a/mlir/tools/mlir-tblgen/OpFormatGen.cpp b/mlir/tools/mlir-tblgen/OpFormatGen.cpp
index e802bd35e914c..e16844d4f8ba6 100644
--- a/mlir/tools/mlir-tblgen/OpFormatGen.cpp
+++ b/mlir/tools/mlir-tblgen/OpFormatGen.cpp
@@ -476,28 +476,14 @@ static bool canFormatEnumAttr(const NamedAttribute *attr) {
if (!baseAttr.isEnumAttr())
return false;
- // For newer EnumAttr-based attributes (which extend AttrDef), only apply
- // enum keyword formatting when the attribute uses the default "$value"
- // assembly format. If it has a custom format (e.g., `<` $value `>`), the
- // attribute's own AttrDef parser/printer handles formatting — using the
- // keyword path here would conflict with that custom format.
- if (baseAttr.isSubClassOf("EnumAttr")) {
- llvm::StringRef asmFmt =
- baseAttr.getDef().getValueAsString("assemblyFormat");
- if (asmFmt != "$value")
- return false;
- }
-
- EnumInfo enumInfo(getEnumInfoRecord(baseAttr));
-
- // Unquoted bit enums may consist of multiple keywords separated by a comma
- // or vertical bar. Implicit formatting defers to the attribute parser;
- // explicit `enum` directives select a separator-aware operation parser.
- if (baseAttr.isSubClassOf("EnumAttr") && enumInfo.isBitEnum() &&
- !enumInfo.printBitEnumQuoted())
+ // New-style EnumAttr-based attributes have a custom AttrDef parser and
+ // printer. Only format their symbolic value directly when requested with an
+ // `enum` directive.
+ if (baseAttr.isSubClassOf("EnumAttr"))
return false;
// The attribute must have a valid underlying type and a constant builder.
+ EnumInfo enumInfo(getEnumInfoRecord(baseAttr));
return !enumInfo.getUnderlyingType().empty() &&
!baseAttr.getConstBuilderTemplate().empty();
}
More information about the flang-commits
mailing list