[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 05:10:46 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/220608
>From d0999a91c8227375a49ae39c35bfbb1a6d6b8d99 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 | 6 +-
.../flang/Optimizer/Dialect/FIRAttr.td | 2 +-
.../include/flang/Optimizer/Dialect/FIROps.td | 3 +-
.../acc-declare-ctor-dtor-conversion.fir | 12 +-
.../acc-declare-gpu-module-insertion.fir | 20 +--
.../Fir/OpenACC/declare-action-conversion.fir | 4 +-
.../offload-livein-value-canonicalization.fir | 10 +-
.../Fir/OpenACC/propagate-attr-folding.fir | 4 +-
flang/test/Lower/OpenACC/acc-cache.f90 | 6 +-
.../acc-declare-common-in-function.f90 | 4 +-
.../OpenACC/acc-declare-global-component.f90 | 4 +-
.../Lower/OpenACC/acc-declare-globals.f90 | 28 ++--
.../OpenACC/acc-declare-interface-body.f90 | 2 +-
.../Lower/OpenACC/acc-declare-unified.f90 | 2 +-
...acc-declare-use-associated-allocatable.f90 | 2 +-
.../OpenACC/acc-declare-use-associated.f90 | 4 +-
flang/test/Lower/OpenACC/acc-declare.f90 | 44 ++---
.../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 +-
.../Transforms/OpenACC/acc-implicit-data.fir | 4 +-
...cit-declare-type-descriptor-create-box.fir | 2 +-
.../OpenACC/offload-target-verifier.fir | 4 +-
.../lower-workdistribute-fission-target.mlir | 2 +-
.../omp-map-info-finalization-usm.fir | 2 +-
.../mlir/Dialect/AMDGPU/IR/AMDGPUOps.td | 9 +-
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 41 +++--
.../include/mlir/Dialect/Index/IR/IndexOps.td | 2 +-
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 156 ++++++++++--------
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 35 ++--
.../mlir/Dialect/OpenACC/OpenACCOps.td | 34 ++--
.../mlir/Dialect/OpenMP/OpenMPClauses.td | 2 +-
.../mlir/Dialect/OpenMP/OpenMPEnums.td | 26 +--
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td | 10 +-
.../mlir/Dialect/Shard/IR/ShardBase.td | 2 +-
.../include/mlir/Dialect/Shard/IR/ShardOps.td | 7 +-
.../SparseTensor/IR/SparseTensorOps.td | 11 +-
.../mlir/Dialect/XeGPU/IR/XeGPUAttrs.td | 6 +-
.../include/mlir/Dialect/XeGPU/IR/XeGPUOps.td | 7 +-
mlir/include/mlir/IR/EnumAttr.td | 15 +-
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 20 ---
.../ACCDeclareGPUModuleInsertion.cpp | 6 +-
.../OpenACC/Transforms/ACCImplicitData.cpp | 14 +-
.../OpenACC/Transforms/ACCImplicitDeclare.cpp | 6 +-
mlir/python/mlir/dialects/irdl.py | 3 +-
mlir/test/CAPI/irdl.c | 8 +-
.../Conversion/XeGPUToXeVM/loadstore_1d.mlir | 4 +-
.../Conversion/XeGPUToXeVM/loadstore_nd.mlir | 8 +-
.../XeGPUToXeVM/loadstore_nd_sub_byte.mlir | 4 +-
.../Conversion/XeGPUToXeVM/prefetch_nd.mlir | 4 +-
mlir/test/Dialect/GPU/invalid.mlir | 6 +-
.../Dialect/IRDL/variadics-error.irdl.mlir | 8 +-
.../acc-declare-ctor-dtor-conversion.mlir | 16 +-
.../acc-declare-gpu-module-insertion.mlir | 12 +-
.../Dialect/OpenACC/acc-implicit-data.mlir | 4 +-
.../Dialect/OpenACC/acc-implicit-declare.mlir | 10 +-
mlir/test/Dialect/OpenACC/invalid.mlir | 4 +-
...offload-livein-value-canonicalization.mlir | 4 +-
.../OpenACC/offload-target-verifier.mlir | 2 +-
mlir/test/Dialect/OpenACC/ops.mlir | 40 ++---
mlir/test/Dialect/OpenMP/invalid.mlir | 12 +-
mlir/test/Dialect/OpenMP/ops.mlir | 36 ++--
mlir/test/Dialect/XeGPU/invalid.mlir | 2 +-
mlir/test/Dialect/XeGPU/ops.mlir | 4 +-
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 | 15 +-
mlir/test/lib/Dialect/Test/TestOps.td | 2 +-
.../attr-or-type-format-roundtrip.mlir | 8 +-
.../test/mlir-tblgen/enums-python-bindings.td | 4 +-
mlir/test/mlir-tblgen/op-format.td | 4 +-
mlir/tools/mlir-tblgen/OpFormatGen.cpp | 24 +--
92 files changed, 470 insertions(+), 465 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..7a39fc028bb9d 100644
--- a/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td
+++ b/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.td
@@ -36,7 +36,7 @@ def cuf_DataAttribute : I32EnumAttr<
def cuf_DataAttributeAttr :
EnumAttr<CUFDialect, cuf_DataAttribute, "cuda"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def cuf_ProcAttribute : I32EnumAttr<
@@ -54,7 +54,7 @@ def cuf_ProcAttribute : I32EnumAttr<
def cuf_ProcAttributeAttr :
EnumAttr<CUFDialect, cuf_ProcAttribute, "cuda_proc"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def cuf_LaunchBoundsAttr : cuf_Attr<"LaunchBounds"> {
@@ -94,7 +94,7 @@ def cuf_DataTransferKind : I32EnumAttr<
def cuf_DataTransferKindAttr :
EnumAttr<CUFDialect, cuf_DataTransferKind, "cuda_transfer"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ 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..09d1a52d938e9 100644
--- a/flang/include/flang/Optimizer/Dialect/FIRAttr.td
+++ b/flang/include/flang/Optimizer/Dialect/FIRAttr.td
@@ -275,7 +275,7 @@ def LocalitySpecifierType : I32EnumAttr<
def LocalitySpecifierTypeAttr : EnumAttr<FIROpsDialect, LocalitySpecifierType,
"locality_specifier_type"> {
- let assemblyFormat = "`{` `type` `=` $value `}`";
+ let assemblyFormat = "`<` $value `>`";
}
/// Fortran inline attribute
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/Fir/OpenACC/acc-declare-ctor-dtor-conversion.fir b/flang/test/Fir/OpenACC/acc-declare-ctor-dtor-conversion.fir
index 30974ecf9c546..0ab3a4a27cb35 100644
--- a/flang/test/Fir/OpenACC/acc-declare-ctor-dtor-conversion.fir
+++ b/flang/test/Fir/OpenACC/acc-declare-ctor-dtor-conversion.fir
@@ -33,35 +33,35 @@
// NODTOR-NOT: llvm.mlir.global_dtors
module attributes {gpu.container_module} {
- fir.global @_QMmmEarr {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<7xf32> {
+ fir.global @_QMmmEarr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<7xf32> {
%0 = fir.zero_bits !fir.array<7xf32>
fir.has_value %0 : !fir.array<7xf32>
}
- fir.global @_QMmmEotherEarr {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<3xf32> {
+ fir.global @_QMmmEotherEarr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<3xf32> {
%0 = fir.zero_bits !fir.array<3xf32>
fir.has_value %0 : !fir.array<3xf32>
}
acc.global_ctor @_QMmmEarr_acc_ctor {
- %0 = fir.address_of(@_QMmmEarr) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<7xf32>>
+ %0 = fir.address_of(@_QMmmEarr) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<7xf32>>
%1 = acc.create varPtr(%0 : !fir.ref<!fir.array<7xf32>>) structured(false) name("arr") -> !fir.ref<!fir.array<7xf32>>
%2 = acc.declare_enter dataOperands(%1 : !fir.ref<!fir.array<7xf32>>)
acc.terminator
}
acc.global_ctor @_QMmmEotherEarr_acc_ctor {
- %0 = fir.address_of(@_QMmmEotherEarr) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<3xf32>>
+ %0 = fir.address_of(@_QMmmEotherEarr) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<3xf32>>
%1 = acc.create varPtr(%0 : !fir.ref<!fir.array<3xf32>>) structured(false) name("other_arr") -> !fir.ref<!fir.array<3xf32>>
%2 = acc.declare_enter dataOperands(%1 : !fir.ref<!fir.array<3xf32>>)
acc.terminator
}
acc.global_dtor @_QMmmEarr_acc_dtor {
- %0 = fir.address_of(@_QMmmEarr) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<7xf32>>
+ %0 = fir.address_of(@_QMmmEarr) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<7xf32>>
%1 = acc.getdeviceptr varPtr(%0 : !fir.ref<!fir.array<7xf32>>) dataClause(acc_create) structured(false) name("arr") -> !fir.ref<!fir.array<7xf32>>
acc.declare_exit dataOperands(%1 : !fir.ref<!fir.array<7xf32>>)
acc.delete accPtr(%1 : !fir.ref<!fir.array<7xf32>>) dataClause(acc_create) structured(false) name("arr")
acc.terminator
}
acc.global_dtor @_QMmmEotherEarr_acc_dtor {
- %0 = fir.address_of(@_QMmmEotherEarr) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<3xf32>>
+ %0 = fir.address_of(@_QMmmEotherEarr) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<3xf32>>
%1 = acc.getdeviceptr varPtr(%0 : !fir.ref<!fir.array<3xf32>>) dataClause(acc_create) structured(false) name("other_arr") -> !fir.ref<!fir.array<3xf32>>
acc.declare_exit dataOperands(%1 : !fir.ref<!fir.array<3xf32>>)
acc.delete accPtr(%1 : !fir.ref<!fir.array<3xf32>>) dataClause(acc_create) structured(false) name("other_arr")
diff --git a/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
index 5f7edd6970e8d..3e289c202ca74 100644
--- a/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
+++ b/flang/test/Fir/OpenACC/acc-declare-gpu-module-insertion.fir
@@ -3,25 +3,25 @@
// Globals registered as host variables must be declarations in device code.
// device_resident globals remain device definitions.
module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", gpu.container_module} {
- fir.global internal @shared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<7xf32> {
+ fir.global internal @shared {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<7xf32> {
%0 = fir.zero_bits !fir.array<7xf32>
fir.has_value %0 : !fir.array<7xf32>
}
- fir.global @initialized_shared(dense<1.0> : vector<1xf32>) {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.array<1xf32>
- fir.global @resident {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>} : !fir.array<7xf32> {
+ fir.global @initialized_shared(dense<1.0> : vector<1xf32>) {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.array<1xf32>
+ fir.global @resident {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>} : !fir.array<7xf32> {
%0 = fir.zero_bits !fir.array<7xf32>
fir.has_value %0 : !fir.array<7xf32>
}
- fir.global @precloned {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<1xf32> {
+ fir.global @precloned {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<1xf32> {
%0 = fir.zero_bits !fir.array<1xf32>
fir.has_value %0 : !fir.array<1xf32>
}
- fir.global @constant_shared {acc.declare = #acc.declare<dataClause = acc_copyin>} constant : i32 {
+ fir.global @constant_shared {acc.declare = #acc.declare<dataClause = <acc_copyin>>} constant : i32 {
%0 = arith.constant 42 : i32
fir.has_value %0 : i32
}
gpu.module @acc_gpu_module {
- fir.global @precloned {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<1xf32> {
+ fir.global @precloned {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<1xf32> {
%0 = fir.zero_bits !fir.array<1xf32>
fir.has_value %0 : !fir.array<1xf32>
}
@@ -29,13 +29,13 @@ module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", gpu.conta
}
// CHECK-LABEL: gpu.module @acc_gpu_module {
-// CHECK: fir.global @precloned {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<1xf32>
+// CHECK: fir.global @precloned {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<1xf32>
// CHECK-NOT: fir.has_value
-// CHECK: fir.global @shared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<7xf32>
+// CHECK: fir.global @shared {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<7xf32>
// CHECK-NOT: fir.has_value
-// CHECK: fir.global @initialized_shared {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.array<1xf32>
+// CHECK: fir.global @initialized_shared {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.array<1xf32>
// CHECK-NOT: fir.has_value
// CHECK: fir.global @resident {{.*}} : !fir.array<7xf32> {
// CHECK: fir.has_value
-// CHECK: fir.global @constant_shared {acc.declare = #acc.declare<dataClause = acc_copyin>} constant : i32 {
+// CHECK: fir.global @constant_shared {acc.declare = #acc.declare<dataClause = <acc_copyin>>} constant : i32 {
// CHECK: fir.has_value
diff --git a/flang/test/Fir/OpenACC/declare-action-conversion.fir b/flang/test/Fir/OpenACC/declare-action-conversion.fir
index d4f82cf9045b6..2096e2d081a80 100644
--- a/flang/test/Fir/OpenACC/declare-action-conversion.fir
+++ b/flang/test/Fir/OpenACC/declare-action-conversion.fir
@@ -15,7 +15,7 @@
// call sub()
// end program main
module {
- fir.global @_QMmmEarr {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.box<!fir.heap<!fir.array<?xf32>>> {
+ fir.global @_QMmmEarr {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.box<!fir.heap<!fir.array<?xf32>>> {
}
func.func private @_QMmmEarr_acc_declare_update_desc_post_alloc() {
return
@@ -86,7 +86,7 @@ module {
%2 = fir.zero_bits !fir.heap<!fir.array<?xf32>>
%3 = fir.shape %c0 : (index) -> !fir.shape<1>
%4 = fir.embox %2(%3) : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.heap<!fir.array<?xf32>>>
- %5 = fir.declare %1 {acc.declare = #acc.declare<dataClause = acc_copyin>, fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QMmmFsubEarr"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>
+ %5 = fir.declare %1 {acc.declare = #acc.declare<dataClause = <acc_copyin>>, fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QMmmFsubEarr"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>
%8 = fir.allocmem !fir.array<?xf32>, %c100 {fir.must_be_heap = true, uniq_name = "_QMmmFsubEarr.alloc"}
%9 = fir.shape %c100 : (index) -> !fir.shape<1>
%10 = fir.embox %8(%9) : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.heap<!fir.array<?xf32>>>
diff --git a/flang/test/Fir/OpenACC/offload-livein-value-canonicalization.fir b/flang/test/Fir/OpenACC/offload-livein-value-canonicalization.fir
index fb23a361605a1..eb50187c1d7c1 100644
--- a/flang/test/Fir/OpenACC/offload-livein-value-canonicalization.fir
+++ b/flang/test/Fir/OpenACC/offload-livein-value-canonicalization.fir
@@ -96,7 +96,7 @@ func.func @test_fieldindex_sink() {
// -----
// Test fir.address_of with acc.declare sinking
-fir.global @global_with_declare {acc.declare = #acc.declare<dataClause = acc_copyin>} : f32 {
+fir.global @global_with_declare {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : f32 {
%0 = arith.constant 0.0 : f32
fir.has_value %0 : f32
}
@@ -144,7 +144,7 @@ func.func @test_address_of_constant_global_sink() {
// -----
// Test fir.address_of with fir.convert tracing (ViewLikeOpInterface)
-fir.global @global_for_convert {acc.declare = #acc.declare<dataClause = acc_copyin>} : f32 {
+fir.global @global_for_convert {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : f32 {
%0 = arith.constant 0.0 : f32
fir.has_value %0 : f32
}
@@ -169,7 +169,7 @@ func.func @test_address_of_with_convert_sink() {
// -----
// Test fir.declare with PartialEntityAccessOpInterface tracing
-fir.global @global_for_declare {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.array<10xf32> {
+fir.global @global_for_declare {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.array<10xf32> {
%0 = fir.zero_bits !fir.array<10xf32>
fir.has_value %0 : !fir.array<10xf32>
}
@@ -479,7 +479,7 @@ func.func @test_undef_slice_rematerialize(%arg0: !fir.box<!fir.array<?x?xf32>>)
// from the sunk/cloned declare instead of being left as an illegal live-in
// or cloned alongside it.
-fir.global @global_for_dummy_scope {acc.declare = #acc.declare<dataClause = acc_copyin>} : i32 {
+fir.global @global_for_dummy_scope {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : i32 {
%0 = arith.constant 0 : i32
fir.has_value %0 : i32
}
@@ -512,7 +512,7 @@ func.func @test_dummy_scope_sink() {
// the outer declare keeps its dummy_scope untouched, while the clone made
// for the region drops it rather than duplicating the scope.
-fir.global @global_for_dummy_scope_remat {acc.declare = #acc.declare<dataClause = acc_copyin>} : i32 {
+fir.global @global_for_dummy_scope_remat {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : i32 {
%0 = arith.constant 0 : i32
fir.has_value %0 : i32
}
diff --git a/flang/test/Fir/OpenACC/propagate-attr-folding.fir b/flang/test/Fir/OpenACC/propagate-attr-folding.fir
index 5676e9de9e981..d5a890bcd5b97 100644
--- a/flang/test/Fir/OpenACC/propagate-attr-folding.fir
+++ b/flang/test/Fir/OpenACC/propagate-attr-folding.fir
@@ -26,7 +26,7 @@ func.func @_QPsub1(%arg0: !fir.ref<!fir.array<?x?xf32>> {fir.bindc_name = "a"},
%17:3 = fir.box_dims %12, %c1 : (!fir.box<!fir.array<?x?xf32>>, index) -> (index, index, index)
%18 = arith.subi %17#1, %c1 : index
%19 = acc.bounds lowerbound(%c0 : index) upperbound(%18 : index) extent(%17#1 : index) stride(%16 : index) startIdx(%c1 : index) strideInBytes(true)
- %20 = fir.box_addr %12 {acc.declare = #acc.declare<dataClause = acc_present>} : (!fir.box<!fir.array<?x?xf32>>) -> !fir.ref<!fir.array<?x?xf32>>
+ %20 = fir.box_addr %12 {acc.declare = #acc.declare<dataClause = <acc_present>>} : (!fir.box<!fir.array<?x?xf32>>) -> !fir.ref<!fir.array<?x?xf32>>
%21 = acc.present varPtr(%20 : !fir.ref<!fir.array<?x?xf32>>) bounds(%15, %19) name("a") -> !fir.ref<!fir.array<?x?xf32>>
%22 = acc.declare_enter dataOperands(%21 : !fir.ref<!fir.array<?x?xf32>>)
acc.declare_exit token(%22)
@@ -35,6 +35,6 @@ func.func @_QPsub1(%arg0: !fir.ref<!fir.array<?x?xf32>> {fir.bindc_name = "a"},
// CHECK-LABEL: func.func @_QPsub1(
// CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<?x?xf32>> {fir.bindc_name = "a"}
-// CHECK: %[[DECL:.*]] = fir.declare %[[ARG0]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_present>, uniq_name = "_QFsub1Ea"} : (!fir.ref<!fir.array<?x?xf32>>, !fir.shape<2>) -> !fir.ref<!fir.array<?x?xf32>>
+// CHECK: %[[DECL:.*]] = fir.declare %[[ARG0]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_present>>, uniq_name = "_QFsub1Ea"} : (!fir.ref<!fir.array<?x?xf32>>, !fir.shape<2>) -> !fir.ref<!fir.array<?x?xf32>>
// CHECK: %[[PRES:.*]] = acc.present varPtr(%[[DECL]] : !fir.ref<!fir.array<?x?xf32>>) bounds(%{{.*}}, %{{.*}}) name("a") -> !fir.ref<!fir.array<?x?xf32>>
// CHECK: %{{.*}} = acc.declare_enter dataOperands(%[[PRES]] : !fir.ref<!fir.array<?x?xf32>>)
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/OpenACC/acc-declare-common-in-function.f90 b/flang/test/Lower/OpenACC/acc-declare-common-in-function.f90
index ab010d1eab808..a9e06a23cba4e 100644
--- a/flang/test/Lower/OpenACC/acc-declare-common-in-function.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-common-in-function.f90
@@ -13,13 +13,13 @@ program p
data pi/0.0/
! CHECK-DAG: acc.global_ctor @{{.*}}_acc_ctor {
-! CHECK-DAG: %[[ADDR0:.*]] = fir.address_of(@{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copyin>} : {{.*}}
+! CHECK-DAG: %[[ADDR0:.*]] = fir.address_of(@{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : {{.*}}
! CHECK-DAG: acc.declare_enter dataOperands(%{{.*}} : {{.*}})
! CHECK-DAG: acc.terminator
! CHECK-DAG: }
! CHECK-DAG: acc.global_dtor @{{.*}}_acc_dtor {
-! CHECK-DAG: %[[ADDR1:.*]] = fir.address_of(@{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.ref<tuple<f32>>
+! CHECK-DAG: %[[ADDR1:.*]] = fir.address_of(@{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.ref<tuple<f32>>
! CHECK-DAG: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[ADDR1]] : !fir.ref<tuple<f32>>) dataClause(acc_copyin) structured(false) name("com") -> !fir.ref<tuple<f32>>
! CHECK-DAG: acc.declare_exit dataOperands(%[[GDP]] : !fir.ref<tuple<f32>>)
! CHECK-DAG: acc.delete accPtr(%[[GDP]] : !fir.ref<tuple<f32>>) dataClause(acc_copyin) structured(false) name("com")
diff --git a/flang/test/Lower/OpenACC/acc-declare-global-component.f90 b/flang/test/Lower/OpenACC/acc-declare-global-component.f90
index ff74f7f170e72..b246ad0b16708 100644
--- a/flang/test/Lower/OpenACC/acc-declare-global-component.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-global-component.f90
@@ -8,11 +8,11 @@
! RUN: bbc -fopenacc -emit-hlfir %t/mod_same_clause.f90 -o - | FileCheck %s --check-prefixes=SAME
! RUN: bbc -fopenacc -emit-hlfir %t/mod_separate_declare.f90 -o - | FileCheck %s --check-prefixes=SEP
-! SAME-DAG: fir.global @_QMmod_same_clauseEobj {acc.declare = #acc.declare<dataClause = acc_create>}
+! SAME-DAG: fir.global @_QMmod_same_clauseEobj {acc.declare = #acc.declare<dataClause = <acc_create>>}
! SAME-COUNT-1: acc.global_ctor @_QMmod_same_clauseEobj_acc_ctor {
! SAME-COUNT-1: acc.global_dtor @_QMmod_same_clauseEobj_acc_dtor {
-! SEP-DAG: fir.global @_QMmod_separate_declareEobj {acc.declare = #acc.declare<dataClause = acc_create>}
+! SEP-DAG: fir.global @_QMmod_separate_declareEobj {acc.declare = #acc.declare<dataClause = <acc_create>>}
! SEP-COUNT-1: acc.global_ctor @_QMmod_separate_declareEobj_acc_ctor {
! SEP-COUNT-1: acc.global_dtor @_QMmod_separate_declareEobj_acc_dtor {
diff --git a/flang/test/Lower/OpenACC/acc-declare-globals.f90 b/flang/test/Lower/OpenACC/acc-declare-globals.f90
index f0eda1b0f6f76..b332d41fe072f 100644
--- a/flang/test/Lower/OpenACC/acc-declare-globals.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-globals.f90
@@ -4,7 +4,7 @@
! RUN: bbc -fopenacc -emit-hlfir %s -o - | FileCheck %s
module acc_declare_common_test
-! CHECK-LABEL: fir.global @numbers_ {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>, alignment = 4 : i64} : tuple<f32, f32> {
+! CHECK-LABEL: fir.global @numbers_ {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>, alignment = 4 : i64} : tuple<f32, f32> {
! CHECK: acc.global_ctor @numbers__acc_ctor
! CHECK: acc.global_dtor @numbers__acc_dtor
REAL :: one = 1
@@ -12,7 +12,7 @@ module acc_declare_common_test
COMMON /numbers/ one, two
!$acc declare device_resident(/numbers/)
-! CHECK-LABEL: fir.global @numbers_create_ {acc.declare = #acc.declare<dataClause = acc_create>, alignment = 4 : i64} : tuple<f32, f32> {
+! CHECK-LABEL: fir.global @numbers_create_ {acc.declare = #acc.declare<dataClause = <acc_create>>, alignment = 4 : i64} : tuple<f32, f32> {
! CHECK: acc.global_ctor @numbers_create__acc_ctor
! CHECK: acc.global_dtor @numbers_create__acc_dtor
REAL :: one_create = 1
@@ -20,7 +20,7 @@ module acc_declare_common_test
COMMON /numbers_create/ one_create, two_create
!$acc declare create(/numbers_create/)
-! CHECK-LABEL: fir.global @numbers_in_ {acc.declare = #acc.declare<dataClause = acc_copyin>, alignment = 4 : i64} : tuple<f32, f32> {
+! CHECK-LABEL: fir.global @numbers_in_ {acc.declare = #acc.declare<dataClause = <acc_copyin>>, alignment = 4 : i64} : tuple<f32, f32> {
! CHECK: acc.global_ctor @numbers_in__acc_ctor
! CHECK: acc.global_dtor @numbers_in__acc_dtor
REAL :: one_in = 1
@@ -28,7 +28,7 @@ module acc_declare_common_test
COMMON /numbers_in/ one_in, two_in
!$acc declare copyin(/numbers_in/)
-! CHECK-LABEL: fir.global @numbers_link_ {acc.declare = #acc.declare<dataClause = acc_declare_link>, alignment = 4 : i64} : tuple<f32, f32> {
+! CHECK-LABEL: fir.global @numbers_link_ {acc.declare = #acc.declare<dataClause = <acc_declare_link>>, alignment = 4 : i64} : tuple<f32, f32> {
! CHECK: acc.global_ctor @numbers_link__acc_ctor
REAL :: one_link = 1
REAL :: two_link = 2
@@ -42,17 +42,17 @@ module acc_declare_test
!$acc declare create(data1)
end module
-! CHECK-LABEL: fir.global @_QMacc_declare_testEdata1 {acc.declare = #acc.declare<dataClause = acc_create>, alignment = 64 : i64} : !fir.array<100000xf32>
+! CHECK-LABEL: fir.global @_QMacc_declare_testEdata1 {acc.declare = #acc.declare<dataClause = <acc_create>>, alignment = 64 : i64} : !fir.array<100000xf32>
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_testEdata1_acc_ctor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_testEdata1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<100000xf32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<100000xf32>>
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<100000xf32>>) structured(false) name("data1") -> !fir.ref<!fir.array<100000xf32>>
! CHECK: acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<100000xf32>>)
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_testEdata1_acc_dtor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_testEdata1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<100000xf32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<100000xf32>>
! CHECK: %[[DEVICEPTR:.*]] = acc.getdeviceptr varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<100000xf32>>) dataClause(acc_create) structured(false) name("data1") -> !fir.ref<!fir.array<100000xf32>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<100000xf32>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.array<100000xf32>>) dataClause(acc_create) structured(false) name("data1")
@@ -66,14 +66,14 @@ module acc_declare_copyin_test
end module
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_copyin_testEdata1_acc_ctor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_copyin_testEdata1) {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.ref<!fir.array<100000xf32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_copyin_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.ref<!fir.array<100000xf32>>
! CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<100000xf32>>) structured(false) name("data1") -> !fir.ref<!fir.array<100000xf32>>
! CHECK: acc.declare_enter dataOperands(%[[COPYIN]] : !fir.ref<!fir.array<100000xf32>>)
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_copyin_testEdata1_acc_dtor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_copyin_testEdata1) {acc.declare = #acc.declare<dataClause = acc_copyin>} : !fir.ref<!fir.array<100000xf32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_copyin_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : !fir.ref<!fir.array<100000xf32>>
! CHECK: %[[DEVICEPTR:.*]] = acc.getdeviceptr varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<100000xf32>>) dataClause(acc_copyin) structured(false) name("data1") -> !fir.ref<!fir.array<100000xf32>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<100000xf32>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.array<100000xf32>>) dataClause(acc_copyin) structured(false) name("data1")
@@ -86,17 +86,17 @@ module acc_declare_device_resident_test
!$acc declare device_resident(data1)
end module
-! CHECK-LABEL: fir.global @_QMacc_declare_device_resident_testEdata1 {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>, alignment = 64 : i64} : !fir.array<5000xi32>
+! CHECK-LABEL: fir.global @_QMacc_declare_device_resident_testEdata1 {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>, alignment = 64 : i64} : !fir.array<5000xi32>
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_device_resident_testEdata1_acc_ctor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_resident_testEdata1) {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>} : !fir.ref<!fir.array<5000xi32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_resident_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>} : !fir.ref<!fir.array<5000xi32>>
! CHECK: %[[DEVICERESIDENT:.*]] = acc.declare_device_resident varPtr(%0 : !fir.ref<!fir.array<5000xi32>>) structured(false) name("data1") -> !fir.ref<!fir.array<5000xi32>>
! CHECK: acc.declare_enter dataOperands(%[[DEVICERESIDENT]] : !fir.ref<!fir.array<5000xi32>>)
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_device_resident_testEdata1_acc_dtor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_resident_testEdata1) {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>} : !fir.ref<!fir.array<5000xi32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_resident_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>} : !fir.ref<!fir.array<5000xi32>>
! CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<5000xi32>>) dataClause(acc_declare_device_resident) structured(false) name("data1") -> !fir.ref<!fir.array<5000xi32>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<5000xi32>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.array<5000xi32>>) dataClause(acc_declare_device_resident) structured(false) name("data1")
@@ -109,10 +109,10 @@ module acc_declare_device_link_test
!$acc declare link(data1)
end module
-! CHECK-LABEL: fir.global @_QMacc_declare_device_link_testEdata1 {acc.declare = #acc.declare<dataClause = acc_declare_link>, alignment = 64 : i64} : !fir.array<5000xi32> {
+! CHECK-LABEL: fir.global @_QMacc_declare_device_link_testEdata1 {acc.declare = #acc.declare<dataClause = <acc_declare_link>>, alignment = 64 : i64} : !fir.array<5000xi32> {
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_device_link_testEdata1_acc_ctor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_link_testEdata1) {acc.declare = #acc.declare<dataClause = acc_declare_link>} : !fir.ref<!fir.array<5000xi32>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_device_link_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_declare_link>>} : !fir.ref<!fir.array<5000xi32>>
! CHECK: %[[LINK:.*]] = acc.declare_link varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.array<5000xi32>>) structured(false) name("data1") -> !fir.ref<!fir.array<5000xi32>>
! CHECK: acc.declare_enter dataOperands(%[[LINK]] : !fir.ref<!fir.array<5000xi32>>)
! CHECK: acc.terminator
diff --git a/flang/test/Lower/OpenACC/acc-declare-interface-body.f90 b/flang/test/Lower/OpenACC/acc-declare-interface-body.f90
index c119d527a0fb6..8cd2a3cb4ac3d 100644
--- a/flang/test/Lower/OpenACC/acc-declare-interface-body.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-interface-body.f90
@@ -37,7 +37,7 @@ end subroutine compute
! The actual definition lowers the present declare on its dummy argument.
! CHECK-LABEL: func.func @_QPcompute(
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}} {acc.declare = #acc.declare<dataClause = acc_present>, uniq_name = "_QFcomputeEa"}
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}} {acc.declare = #acc.declare<dataClause = <acc_present>>, uniq_name = "_QFcomputeEa"}
! CHECK: %[[PRES:.*]] = acc.present var(%[[DECL]]#0
! CHECK: acc.declare_enter dataOperands(%[[PRES]]
! CHECK: acc.declare_exit
diff --git a/flang/test/Lower/OpenACC/acc-declare-unified.f90 b/flang/test/Lower/OpenACC/acc-declare-unified.f90
index 871d10d94a2d6..2555eab9ac220 100644
--- a/flang/test/Lower/OpenACC/acc-declare-unified.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-unified.f90
@@ -59,7 +59,7 @@ subroutine use_mod()
! The using unit neither annotates the allocation nor declares external recipes.
! UNIFIED-USE: func.func @_QPuse_mod()
-! UNIFIED-USE: fir.global @_QMacc_declare_unified_modEgarr {acc.declare = #acc.declare<dataClause = acc_create>
+! UNIFIED-USE: fir.global @_QMacc_declare_unified_modEgarr {acc.declare = #acc.declare<dataClause = <acc_create>>
! DISCRETE-MOD: acc.declare_action = #acc.declare_action<postAlloc = @_QMacc_declare_unified_modEgarr_acc_declare_post_alloc>
! DISCRETE-MOD: acc.declare_action = #acc.declare_action<preDealloc = @_QMacc_declare_unified_modEgarr_acc_declare_pre_dealloc>
diff --git a/flang/test/Lower/OpenACC/acc-declare-use-associated-allocatable.f90 b/flang/test/Lower/OpenACC/acc-declare-use-associated-allocatable.f90
index 89338ebdc460b..8a53746827e5b 100644
--- a/flang/test/Lower/OpenACC/acc-declare-use-associated-allocatable.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-use-associated-allocatable.f90
@@ -29,7 +29,7 @@ subroutine use_mod()
! USE: func.func @_QPuse_mod() {
! USE: acc.declare_action = #acc.declare_action<postAlloc = @_QMacc_declare_alloc_modEdata_acc_declare_post_alloc>
-! USE: fir.global @_QMacc_declare_alloc_modEdata {acc.declare = #acc.declare<dataClause = acc_create>
+! USE: fir.global @_QMacc_declare_alloc_modEdata {acc.declare = #acc.declare<dataClause = <acc_create>>
! USE: func.func private @_QMacc_declare_alloc_modEdata_acc_declare_post_alloc()
! USE: func.func private @_QMacc_declare_alloc_modEdata_acc_declare_pre_dealloc()
! USE-NOT: acc.declare_enter
diff --git a/flang/test/Lower/OpenACC/acc-declare-use-associated.f90 b/flang/test/Lower/OpenACC/acc-declare-use-associated.f90
index 4d2275d67efb5..e39fd1166d8b2 100644
--- a/flang/test/Lower/OpenACC/acc-declare-use-associated.f90
+++ b/flang/test/Lower/OpenACC/acc-declare-use-associated.f90
@@ -25,5 +25,5 @@ subroutine use_mod()
end do
end subroutine
-! CHECK: fir.global @_QMacc_declare_modEaa {acc.declare = #acc.declare<dataClause = acc_create>, alignment = 64 : i64} : !fir.array<100xf32>
-! CHECK: fir.global @_QMacc_declare_modEcoef {acc.declare = #acc.declare<dataClause = acc_copyin>} : f32
+! CHECK: fir.global @_QMacc_declare_modEaa {acc.declare = #acc.declare<dataClause = <acc_create>>, alignment = 64 : i64} : !fir.array<100xf32>
+! CHECK: fir.global @_QMacc_declare_modEcoef {acc.declare = #acc.declare<dataClause = <acc_copyin>>} : f32
diff --git a/flang/test/Lower/OpenACC/acc-declare.f90 b/flang/test/Lower/OpenACC/acc-declare.f90
index 8232c82641eeb..35c1d8581f50e 100644
--- a/flang/test/Lower/OpenACC/acc-declare.f90
+++ b/flang/test/Lower/OpenACC/acc-declare.f90
@@ -17,7 +17,7 @@ subroutine acc_declare_copy()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_copy()
! CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QMacc_declareFacc_declare_copyEa"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copy>, uniq_name = "_QMacc_declareFacc_declare_copyEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copy>>, uniq_name = "_QMacc_declareFacc_declare_copyEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) dataClause(acc_copy) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[COPYIN]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -37,7 +37,7 @@ subroutine acc_declare_create()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_create() {
! CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QMacc_declareFacc_declare_createEa"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_create>, uniq_name = "_QMacc_declareFacc_declare_createEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_create>>, uniq_name = "_QMacc_declareFacc_declare_createEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -57,7 +57,7 @@ subroutine acc_declare_present(a)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_present(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "a"})
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_present>, uniq_name = "_QMacc_declareFacc_declare_presentEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_present>>, uniq_name = "_QMacc_declareFacc_declare_presentEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[PRESENT:.*]] = acc.present varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[PRESENT]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -92,9 +92,9 @@ subroutine acc_declare_copyin()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_copyin()
! CHECK: %[[A:.*]] = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QMacc_declareFacc_declare_copyinEa"}
-! CHECK: %[[ADECL:.*]]:2 = hlfir.declare %[[A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copyin>, uniq_name = "_QMacc_declareFacc_declare_copyinEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[ADECL:.*]]:2 = hlfir.declare %[[A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copyin>>, uniq_name = "_QMacc_declareFacc_declare_copyinEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[B:.*]] = fir.alloca !fir.array<10xi32> {bindc_name = "b", uniq_name = "_QMacc_declareFacc_declare_copyinEb"}
-! CHECK: %[[BDECL:.*]]:2 = hlfir.declare %[[B]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copyin_readonly>, uniq_name = "_QMacc_declareFacc_declare_copyinEb"} : (!fir.ref<!fir.array<10xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<10xi32>>, !fir.ref<!fir.array<10xi32>>)
+! CHECK: %[[BDECL:.*]]:2 = hlfir.declare %[[B]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copyin_readonly>>, uniq_name = "_QMacc_declareFacc_declare_copyinEb"} : (!fir.ref<!fir.array<10xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<10xi32>>, !fir.ref<!fir.array<10xi32>>)
! CHECK: %[[COPYIN_A:.*]] = acc.copyin varPtr(%[[ADECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[COPYIN_B:.*]] = acc.copyin varPtr(%[[BDECL]]#0 : !fir.ref<!fir.array<10xi32>>) dataClause(acc_copyin_readonly) name("b") -> !fir.ref<!fir.array<10xi32>>
! CHECK: acc.declare_enter dataOperands(%[[COPYIN_A]], %[[COPYIN_B]] : !fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<10xi32>>)
@@ -113,7 +113,7 @@ subroutine acc_declare_copyout()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_copyout()
! CHECK: %[[A:.*]] = fir.alloca !fir.array<100xi32> {bindc_name = "a", uniq_name = "_QMacc_declareFacc_declare_copyoutEa"}
-! CHECK: %[[ADECL:.*]]:2 = hlfir.declare %[[A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_copyout>, uniq_name = "_QMacc_declareFacc_declare_copyoutEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[ADECL:.*]]:2 = hlfir.declare %[[A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_copyout>>, uniq_name = "_QMacc_declareFacc_declare_copyoutEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[ADECL]]#0 : !fir.ref<!fir.array<100xi32>>) dataClause(acc_copyout) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -132,7 +132,7 @@ subroutine acc_declare_deviceptr(a)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_deviceptr(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "a"}) {
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_deviceptr>, uniq_name = "_QMacc_declareFacc_declare_deviceptrEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_deviceptr>>, uniq_name = "_QMacc_declareFacc_declare_deviceptrEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[DEVICEPTR:.*]] = acc.deviceptr varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: acc.declare_enter dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -148,7 +148,7 @@ subroutine acc_declare_link(a)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_link(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "a"})
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_declare_link>, uniq_name = "_QMacc_declareFacc_declare_linkEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_declare_link>>, uniq_name = "_QMacc_declareFacc_declare_linkEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[LINK:.*]] = acc.declare_link varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: acc.declare_enter dataOperands(%[[LINK]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -164,7 +164,7 @@ subroutine acc_declare_device_resident(a)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_device_resident(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "a"})
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>, uniq_name = "_QMacc_declareFacc_declare_device_residentEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>, uniq_name = "_QMacc_declareFacc_declare_device_residentEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[DEVICERES:.*]] = acc.declare_device_resident varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xi32>>) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[DEVICERES]] : !fir.ref<!fir.array<100xi32>>)
! CHECK: fir.do_loop %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
@@ -179,7 +179,7 @@ subroutine acc_declare_device_resident2()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_device_resident2()
! CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<100xf32> {bindc_name = "dataparam", uniq_name = "_QMacc_declareFacc_declare_device_resident2Edataparam"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_declare_device_resident>, uniq_name = "_QMacc_declareFacc_declare_device_resident2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_declare_device_resident>>, uniq_name = "_QMacc_declareFacc_declare_device_resident2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
! CHECK: %[[DEVICERES:.*]] = acc.declare_device_resident varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xf32>>) name("dataparam") -> !fir.ref<!fir.array<100xf32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[DEVICERES]] : !fir.ref<!fir.array<100xf32>>)
! CHECK: acc.declare_exit token(%[[TOKEN]]) dataOperands(%[[DEVICERES]] : !fir.ref<!fir.array<100xf32>>)
@@ -193,7 +193,7 @@ subroutine acc_declare_link2()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_link2()
! CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<100xf32> {bindc_name = "dataparam", uniq_name = "_QMacc_declareFacc_declare_link2Edataparam"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_declare_link>, uniq_name = "_QMacc_declareFacc_declare_link2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_declare_link>>, uniq_name = "_QMacc_declareFacc_declare_link2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
! CHECK: %[[LINK:.*]] = acc.declare_link varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xf32>>) name("dataparam") -> !fir.ref<!fir.array<100xf32>>
! CHECK: acc.declare_enter dataOperands(%[[LINK]] : !fir.ref<!fir.array<100xf32>>)
@@ -205,7 +205,7 @@ subroutine acc_declare_deviceptr2()
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_deviceptr2()
! CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<100xf32> {bindc_name = "dataparam", uniq_name = "_QMacc_declareFacc_declare_deviceptr2Edataparam"}
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_deviceptr>, uniq_name = "_QMacc_declareFacc_declare_deviceptr2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ALLOCA]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_deviceptr>>, uniq_name = "_QMacc_declareFacc_declare_deviceptr2Edataparam"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>)
! CHECK: %[[DEVICEPTR:.*]] = acc.deviceptr varPtr(%[[DECL]]#0 : !fir.ref<!fir.array<100xf32>>) name("dataparam") -> !fir.ref<!fir.array<100xf32>>
! CHECK: acc.declare_enter dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<100xf32>>)
@@ -232,7 +232,7 @@ function acc_declare_in_func2(i)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_in_func2(%arg0: !fir.ref<i32> {fir.bindc_name = "i"}) -> f32 {
! CHECK: %[[ALLOCA_A:.*]] = fir.alloca !fir.array<1024xf32> {bindc_name = "a", uniq_name = "_QMacc_declareFacc_declare_in_func2Ea"}
-! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ALLOCA_A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = acc_create>, uniq_name = "_QMacc_declareFacc_declare_in_func2Ea"} : (!fir.ref<!fir.array<1024xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf32>>, !fir.ref<!fir.array<1024xf32>>)
+! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ALLOCA_A]](%{{.*}}) {acc.declare = #acc.declare<dataClause = <acc_create>>, uniq_name = "_QMacc_declareFacc_declare_in_func2Ea"} : (!fir.ref<!fir.array<1024xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf32>>, !fir.ref<!fir.array<1024xf32>>)
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[DECL_A]]#0 : !fir.ref<!fir.array<1024xf32>>) name("a") -> !fir.ref<!fir.array<1024xf32>>
! CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<1024xf32>>)
! CHECK: cf.br ^bb1
@@ -291,8 +291,8 @@ subroutine acc_declare_multiple_directive(a, b)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_multiple_directive(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "a"}, %[[ARG1:.*]]: !fir.ref<!fir.array<100xi32>> {fir.bindc_name = "b"}) {
-! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_copy>, uniq_name = "_QMacc_declareFacc_declare_multiple_directiveEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
-! CHECK: %[[DECL_B:.*]]:2 = hlfir.declare %[[ARG1]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_copyout>, uniq_name = "_QMacc_declareFacc_declare_multiple_directiveEb"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ARG0]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_copy>>, uniq_name = "_QMacc_declareFacc_declare_multiple_directiveEa"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
+! CHECK: %[[DECL_B:.*]]:2 = hlfir.declare %[[ARG1]](%{{.*}}) dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_copyout>>, uniq_name = "_QMacc_declareFacc_declare_multiple_directiveEb"} : (!fir.ref<!fir.array<100xi32>>, !fir.shape<1>, !fir.dscope) -> (!fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
! CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[DECL_A]]#0 : !fir.ref<!fir.array<100xi32>>) dataClause(acc_copy) name("a") -> !fir.ref<!fir.array<100xi32>>
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[DECL_B]]#0 : !fir.ref<!fir.array<100xi32>>) dataClause(acc_copyout) name("b") -> !fir.ref<!fir.array<100xi32>>
! CHECK: acc.declare_enter dataOperands(%[[COPYIN]], %[[CREATE]] : !fir.ref<!fir.array<100xi32>>, !fir.ref<!fir.array<100xi32>>)
@@ -313,7 +313,7 @@ subroutine acc_declare_array_section(a)
! CHECK-LABEL: func.func @_QMacc_declarePacc_declare_array_section(
! CHECK-SAME: %[[ARG0:.*]]: !fir.box<!fir.array<?xi32>> {fir.bindc_name = "a"}) {
-! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ARG0]] dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = acc_copy>, uniq_name = "_QMacc_declareFacc_declare_array_sectionEa"} : (!fir.box<!fir.array<?xi32>>, !fir.dscope) -> (!fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>>)
+! CHECK: %[[DECL_A:.*]]:2 = hlfir.declare %[[ARG0]] dummy_scope %{{[0-9]+}} arg {{[0-9]+}} {acc.declare = #acc.declare<dataClause = <acc_copy>>, uniq_name = "_QMacc_declareFacc_declare_array_sectionEa"} : (!fir.box<!fir.array<?xi32>>, !fir.dscope) -> (!fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>>)
! CHECK: %[[COPYIN:.*]] = acc.copyin var(%[[DECL_A]]#0 : !fir.box<!fir.array<?xi32>>) bounds(%{{.*}}) dataClause(acc_copy) name("a(1:10)") -> !fir.box<!fir.array<?xi32>>
! CHECK: acc.declare_enter dataOperands(%[[COPYIN]] : !fir.box<!fir.array<?xi32>>)
@@ -411,7 +411,7 @@ subroutine init_in_data(e)
! CHECK-NOT: acc.terminator {acc.declare_action =
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_allocatable_testEdata1_acc_ctor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_allocatable_testEdata1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_allocatable_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
! CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) dataClause(acc_create) structured(false) implicit(true) name("data1") -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
! CHECK: acc.declare_enter dataOperands(%[[COPYIN]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
! CHECK: acc.terminator
@@ -435,7 +435,7 @@ subroutine init_in_data(e)
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_allocatable_testEdata1_acc_dtor {
-! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_allocatable_testEdata1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+! CHECK: %[[GLOBAL_ADDR:.*]] = fir.address_of(@_QMacc_declare_allocatable_testEdata1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
! CHECK: %[[DEVICEPTR:.*]] = acc.getdeviceptr varPtr(%[[GLOBAL_ADDR]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) dataClause(acc_create) structured(false) name("data1") -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) dataClause(acc_create) structured(false) name("data1")
@@ -444,26 +444,26 @@ subroutine init_in_data(e)
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_equivalentEv2_acc_ctor {
-! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalentEv1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<40xi8>>
+! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalentEv1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<40xi8>>
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[ADDR]] : !fir.ref<!fir.array<40xi8>>) structured(false) name("v2") -> !fir.ref<!fir.array<40xi8>>
! CHECK: acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<40xi8>>)
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_equivalentEv2_acc_dtor {
-! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalentEv1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<40xi8>>
+! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalentEv1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<40xi8>>
! CHECK: %[[DEVICEPTR:.*]] = acc.getdeviceptr varPtr(%[[ADDR]] : !fir.ref<!fir.array<40xi8>>) dataClause(acc_create) structured(false) name("v2") -> !fir.ref<!fir.array<40xi8>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<40xi8>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.array<40xi8>>) dataClause(acc_create) structured(false) name("v2")
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_ctor @_QMacc_declare_equivalent2Ev2_acc_ctor {
-! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalent2Ev1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<40xi8>>
+! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalent2Ev1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<40xi8>>
! CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[ADDR]] : !fir.ref<!fir.array<40xi8>>) structured(false) name("v2") -> !fir.ref<!fir.array<40xi8>>
! CHECK: acc.declare_enter dataOperands(%[[CREATE]] : !fir.ref<!fir.array<40xi8>>)
! CHECK: acc.terminator
! CHECK: }
! CHECK-LABEL: acc.global_dtor @_QMacc_declare_equivalent2Ev2_acc_dtor {
-! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalent2Ev1) {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.ref<!fir.array<40xi8>>
+! CHECK: %[[ADDR:.*]] = fir.address_of(@_QMacc_declare_equivalent2Ev1) {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.ref<!fir.array<40xi8>>
! CHECK: %[[DEVICEPTR:.*]] = acc.getdeviceptr varPtr(%[[ADDR]] : !fir.ref<!fir.array<40xi8>>) dataClause(acc_create) structured(false) name("v2") -> !fir.ref<!fir.array<40xi8>>
! CHECK: acc.declare_exit dataOperands(%[[DEVICEPTR]] : !fir.ref<!fir.array<40xi8>>)
! CHECK: acc.delete accPtr(%[[DEVICEPTR]] : !fir.ref<!fir.array<40xi8>>) dataClause(acc_create) structured(false) name("v2")
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/OpenACC/acc-implicit-data.fir b/flang/test/Transforms/OpenACC/acc-implicit-data.fir
index 736c6f1e9bf2a..dcdcc4564a3cc 100644
--- a/flang/test/Transforms/OpenACC/acc-implicit-data.fir
+++ b/flang/test/Transforms/OpenACC/acc-implicit-data.fir
@@ -401,7 +401,7 @@ func.func private @_FortranAAllocatableSetBounds(!fir.ref<!fir.box<none>>, i32,
func.func @test_fir_declare_deviceptr_arg_in_parallel(%arg0: !fir.ref<!fir.array<10xf64>>) {
%c10 = arith.constant 10 : index
%shape = fir.shape %c10 : (index) -> !fir.shape<1>
- %arr_decl = fir.declare %arg0(%shape) {acc.declare = #acc.declare<dataClause = acc_deviceptr>, uniq_name = "_QFtestEa"} : (!fir.ref<!fir.array<10xf64>>, !fir.shape<1>) -> !fir.ref<!fir.array<10xf64>>
+ %arr_decl = fir.declare %arg0(%shape) {acc.declare = #acc.declare<dataClause = <acc_deviceptr>>, uniq_name = "_QFtestEa"} : (!fir.ref<!fir.array<10xf64>>, !fir.shape<1>) -> !fir.ref<!fir.array<10xf64>>
%arr_box = fir.embox %arr_decl(%shape) : (!fir.ref<!fir.array<10xf64>>, !fir.shape<1>) -> !fir.box<!fir.array<10xf64>>
%devptr = acc.deviceptr var(%arr_box : !fir.box<!fir.array<10xf64>>) name("a") -> !fir.box<!fir.array<10xf64>>
%token = acc.declare_enter dataOperands(%devptr : !fir.box<!fir.array<10xf64>>)
@@ -415,7 +415,7 @@ func.func @test_fir_declare_deviceptr_arg_in_parallel(%arg0: !fir.ref<!fir.array
}
// CHECK-LABEL: func.func @test_fir_declare_deviceptr_arg_in_parallel
-// CHECK: %[[DECL:.*]] = fir.declare %{{.*}}{{.*}}{acc.declare = #acc.declare<dataClause = acc_deviceptr>{{.*}}
+// CHECK: %[[DECL:.*]] = fir.declare %{{.*}}{{.*}}{acc.declare = #acc.declare<dataClause = <acc_deviceptr>>{{.*}}
// CHECK: %[[BOX:.*]] = fir.embox %[[DECL]]
// CHECK: %[[DEVPTR:.*]] = acc.deviceptr var(%[[BOX]] : !fir.box<!fir.array<10xf64>>) name("a") -> !fir.box<!fir.array<10xf64>>
// CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[DEVPTR]] : !fir.box<!fir.array<10xf64>>)
diff --git a/flang/test/Transforms/OpenACC/acc-implicit-declare-type-descriptor-create-box.fir b/flang/test/Transforms/OpenACC/acc-implicit-declare-type-descriptor-create-box.fir
index 6dd1c19615b2f..40e4b2f4d1cc9 100644
--- a/flang/test/Transforms/OpenACC/acc-implicit-declare-type-descriptor-create-box.fir
+++ b/flang/test/Transforms/OpenACC/acc-implicit-declare-type-descriptor-create-box.fir
@@ -23,4 +23,4 @@ module {
}
}
-// CHECK: fir.global {{.*}}@_QMmmE.dt.struct {acc.declare = #acc.declare<dataClause = acc_copyin>}{{.*}}constant target
+// CHECK: fir.global {{.*}}@_QMmmE.dt.struct {acc.declare = #acc.declare<dataClause = <acc_copyin>>}{{.*}}constant target
diff --git a/flang/test/Transforms/OpenACC/offload-target-verifier.fir b/flang/test/Transforms/OpenACC/offload-target-verifier.fir
index 5b76d8454afc3..4083cedb738da 100644
--- a/flang/test/Transforms/OpenACC/offload-target-verifier.fir
+++ b/flang/test/Transforms/OpenACC/offload-target-verifier.fir
@@ -103,7 +103,7 @@ func.func @test_fir_global_no_declare() {
// -----
// Test fir.global with acc.declare attribute - should pass
-fir.global @_global_array_declared {acc.declare = #acc.declare<dataClause = acc_create>} : !fir.array<10xf32> {
+fir.global @_global_array_declared {acc.declare = #acc.declare<dataClause = <acc_create>>} : !fir.array<10xf32> {
%0 = fir.zero_bits !fir.array<10xf32>
fir.has_value %0 : !fir.array<10xf32>
}
@@ -349,7 +349,7 @@ func.func @test_use_stmt_no_declare() {
// -----
// Test fir.use_stmt inside offload region referencing declared global
-fir.global @_QMmod2Evar2 {acc.declare = #acc.declare<dataClause = acc_create>} : f32 {
+fir.global @_QMmod2Evar2 {acc.declare = #acc.declare<dataClause = <acc_create>>} : f32 {
%0 = arith.constant 0.0 : f32
fir.has_value %0 : f32
}
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/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/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/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9671dca32bd20..cf473d4aa745b 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -75,7 +75,7 @@ def CacheEvictionPriority : I32Enum<"CacheEvictionPriority",
def CacheEvictionPriorityAttr : EnumAttr<NVVM_Dialect, CacheEvictionPriority,
"cache_eviction_priority"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
// Memory Space enum definitions
@@ -124,7 +124,7 @@ def MemScopeKind : I32EnumAttr<"MemScopeKind", "NVVM Memory Scope kind",
let cppNamespace = "::mlir::NVVM";
}
def MemScopeKindAttr : EnumAttr<NVVM_Dialect, MemScopeKind, "mem_scope"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
// Attrs to disambiguate the cta or cluster space within shared memory
@@ -575,7 +575,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 +723,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 +748,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 +799,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 +829,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 +932,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 +978,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 +1107,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 +1140,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;
@@ -1375,7 +1375,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));
}];
@@ -1492,7 +1492,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 +1524,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 +1573,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 +1631,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 +1679,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;
}
@@ -1884,7 +1884,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
@@ -3160,7 +3162,7 @@ def MMAFrag: I32EnumAttr<"MMAFrag", "NVVM MMA frag type",
let cppNamespace = "::mlir::NVVM";
}
def MMAFragAttr : EnumAttr<NVVM_Dialect, MMAFrag, "mma_frag"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_WMMALoadOp: NVVM_Op<"wmma.load">,
@@ -3242,7 +3244,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;
}
@@ -4339,7 +4341,7 @@ 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";
+ let assemblyFormat = "`<` $value `>`";
}
// List of modes supported for TMA Store and Reduction Ops
@@ -4355,7 +4357,7 @@ def TMAStoreMode : I32EnumAttr<"TMAStoreMode", "NVVM TMA Store Mode",
let cppNamespace = "::mlir::NVVM";
}
def TMAStoreModeAttr : EnumAttr<NVVM_Dialect, TMAStoreMode, "tma_store_mode"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
// Num CTAs in a group participating in the TMA/MMA operations.
@@ -4462,7 +4464,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 +4551,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 +4674,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 +4696,7 @@ def PrefetchCacheLevel : I32Enum<"PrefetchCacheLevel",
}
def PrefetchCacheLevelAttr : EnumAttr<NVVM_Dialect, PrefetchCacheLevel, "prefetch_cache_level"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_PrefetchOp : NVVM_Op<"prefetch",
@@ -4736,7 +4738,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 +4822,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 +4852,7 @@ def TMAReduxKind : I32EnumAttr<"TMAReduxKind", "NVVM TMA redux kind",
let cppNamespace = "::mlir::NVVM";
}
def TMAReduxKindAttr : EnumAttr<NVVM_Dialect, TMAReduxKind, "tma_redux_kind"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_CpAsyncBulkTensorReduceOp :
@@ -4877,7 +4885,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 +4973,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)
}];
@@ -5329,7 +5337,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 +5418,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 +5481,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)
}];
@@ -5769,7 +5778,7 @@ def Tcgen05CpShape : I32EnumAttr<"Tcgen05CpShape", "tcgen05 cp shapes",
let genSpecializedAttr = 0;
}
def Tcgen05CpShapeAttr : EnumAttr<NVVM_Dialect, Tcgen05CpShape, "tcgen05_cp_shape"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def Tcgen05CpMulticastNone: I32EnumAttrCase<"NONE", 0, "none">;
@@ -5784,7 +5793,7 @@ def Tcgen05CpMulticast : I32EnumAttr<"Tcgen05CpMulticast", "tcgen05 cp multicast
let genSpecializedAttr = 0;
}
def Tcgen05CpMulticastAttr : EnumAttr<NVVM_Dialect, Tcgen05CpMulticast, "tcgen05_cp_multicast"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def FormatB6x16_P32: I32EnumAttrCase<"B6x16_P32", 0, "b6x16_p32">;
@@ -5826,9 +5835,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 +5940,7 @@ def Tcgen05LdStShape: I32EnumAttr<
}
def Tcgen05LdStShapeAttr: EnumAttr<NVVM_Dialect, Tcgen05LdStShape, "tcgen05_ldst_shape"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -5954,7 +5963,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 +6054,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 +6141,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 = [{
@@ -6365,7 +6374,7 @@ def ClusterLaunchControlQueryType
def ClusterLaunchControlQueryTypeAttr
: EnumAttr<NVVM_Dialect,
ClusterLaunchControlQueryType, "cluster_launch_control_query_type"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_ClusterLaunchControlQueryCancelOp
@@ -6392,7 +6401,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 +6463,7 @@ def Tcgen05MMAKindAttr : EnumAttr<NVVM_Dialect, Tcgen05MMAKind, "tcgen05_mma_kin
+-------------+--------------------------------------------+
```
}];
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def Tcgen05MMACollectorOpDiscard : I32EnumAttrCase<"DISCARD", 0, "discard">;
@@ -6481,7 +6491,7 @@ def Tcgen05MMACollectorOpAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorOp, "t
* fill : Fill buffer
* use : Use buffer without modification
}];
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
defvar Tcgen05MMANonBlockScaleKindList =
@@ -6569,10 +6579,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 +6653,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 +6692,7 @@ def Tcgen05MMABlockScale
def Tcgen05MMABlockScaleAttr : EnumAttr<NVVM_Dialect, Tcgen05MMABlockScale,
"tcgen05_mma_block_scale"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
@@ -6741,11 +6753,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 +6826,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 +6866,7 @@ def Tcgen05MMACollectorBBuffer : I32EnumAttr<
}
def Tcgen05MMACollectorBBufferAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorBBuffer, "tcgen05_mma_collectorb"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
@@ -6909,10 +6923,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 +6989,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 +7074,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 +7148,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 +7381,7 @@ def TensormapField : I32Enum<"TensormapField", "NVVM Tensormap Field Kind",
}
def TensormapFieldAttr : EnumAttr<NVVM_Dialect, TensormapField, "tensormap_field"> {
- let assemblyFormat = "$value";
+ let assemblyFormat = "`<` $value `>`";
}
// Tensormap Elemtype Enum
@@ -7525,7 +7540,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/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/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 1cdd05c6656b5..a548d3c7ef901 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -115,7 +115,7 @@ def OpenACC_ReductionOperator
def OpenACC_ReductionOperatorAttr : EnumAttr<OpenACC_Dialect,
OpenACC_ReductionOperator,
"reduction_operator"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
// OpenACC variable type categorization. This is needed because OpenACC
@@ -227,7 +227,7 @@ def OpenACC_ParLevel : I32EnumAttr<"ParLevel",
def OpenACC_ParLevelAttr : EnumAttr<OpenACC_Dialect,
OpenACC_ParLevel,
"par_level"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def OpenACC_PrivateRecipe : I32EnumAttrCase<"private_recipe", 0>;
@@ -247,7 +247,7 @@ def OpenACC_RecipeKind : I32EnumAttr<
def OpenACC_RecipeKindAttr : EnumAttr<OpenACC_Dialect,
OpenACC_RecipeKind,
"recipe_kind"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
// Type used in operation below.
@@ -426,7 +426,7 @@ def OpenACC_DeviceType : I32EnumAttr<"DeviceType",
def OpenACC_DeviceTypeAttr : EnumAttr<OpenACC_Dialect,
OpenACC_DeviceType,
"device_type"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def DeviceTypeArrayAttr :
@@ -448,7 +448,7 @@ def OpenACC_GangArgType : I32EnumAttr<"GangArgType",
def OpenACC_GangArgTypeAttr : EnumAttr<OpenACC_Dialect,
OpenACC_GangArgType,
"gang_arg_type"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def GangArgTypeArrayAttr :
TypedArrayAttrBase<OpenACC_GangArgTypeAttr, "gang arg type array attribute"> {
@@ -470,7 +470,7 @@ def OpenACC_CombinedConstructsType : I32EnumAttr<"CombinedConstructsType",
def OpenACC_CombinedConstructsAttr : EnumAttr<OpenACC_Dialect,
OpenACC_CombinedConstructsType,
"combined_constructs"> {
- let assemblyFormat = [{ ```<` $value `>` }];
+ let assemblyFormat = "`<` $value `>`";
}
def OpenACC_ParallelConstruct : I64EnumAttrCase<"acc_construct_parallel", 0>;
@@ -817,7 +817,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 +1211,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 +1289,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 +1953,7 @@ def OpenACC_ParallelOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2100,7 +2100,7 @@ def OpenACC_SerialOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2296,7 +2296,7 @@ def OpenACC_KernelsOp
)
$region
oilist(
- `defaultAttr` `(` qualified($defaultAttr) `)`
+ `defaultAttr` `(` enum($defaultAttr) `)`
| `selfAttr` $selfAttr
)
attr-dict-with-keyword
@@ -2416,7 +2416,7 @@ def OpenACC_DataOp
$waitOnly)
)
$region
- (`defaultAttr` `(` qualified($defaultAttr)^ `)`)?
+ (`defaultAttr` `(` enum($defaultAttr)^ `)`)?
attr-dict-with-keyword
}];
let hasVerifier = 1;
@@ -3214,9 +3214,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 +3290,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..c976ddf502cc9 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPEnums.td
@@ -56,7 +56,7 @@ def DeclareTargetCaptureClause : OpenMP_I32EnumAttr<
def DeclareTargetCaptureClauseAttr : OpenMP_EnumAttr<DeclareTargetCaptureClause,
"capture_clause"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -74,7 +74,7 @@ def ClauseDepend : OpenMP_I32EnumAttr<
]>;
def ClauseDependAttr : OpenMP_EnumAttr<ClauseDepend, "clause_depend"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -94,7 +94,7 @@ def ClauseSeverity : OpenMP_I32EnumAttr<
]>;
def ClauseSeverityAttr : OpenMP_EnumAttr<ClauseSeverity, "clause_severity"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -198,7 +198,7 @@ def ClauseTaskDepend
def ClauseTaskDependAttr : OpenMP_EnumAttr<ClauseTaskDepend,
"clause_task_depend"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -218,7 +218,7 @@ def DataSharingClauseType : OpenMP_I32EnumAttr<
def DataSharingClauseTypeAttr : OpenMP_EnumAttr<DataSharingClauseType,
"data_sharing_type"> {
- let assemblyFormat = "`{` `type` `=` $value `}`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -239,7 +239,7 @@ def DeclareTargetDeviceType : OpenMP_I32EnumAttr<
def DeclareTargetDeviceTypeAttr : OpenMP_EnumAttr<DeclareTargetDeviceType,
"device_type"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -270,7 +270,7 @@ def LinearModifier : OpenMP_I32EnumAttr<"LinearModifier", "linear modifier",
LinearModifierUval]>;
def LinearModifierAttr : OpenMP_EnumAttr<LinearModifier, "linear_modifier"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -291,7 +291,7 @@ def ReductionModifier : OpenMP_I32EnumAttr<
def ReductionModifierAttr : OpenMP_EnumAttr<ReductionModifier,
"reduction_modifier"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -335,7 +335,7 @@ def TargetExecMode : OpenMP_I32EnumAttr<
]>;
def TargetExecModeAttr : OpenMP_EnumAttr<TargetExecMode, "kernel_type"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -358,7 +358,7 @@ def VariableCaptureKind : OpenMP_I32EnumAttr<
def VariableCaptureKindAttr : OpenMP_EnumAttr<VariableCaptureKind,
"variable_capture_kind"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -375,7 +375,7 @@ def AccessGroupModifier : OpenMP_I32EnumAttr<
def AccessGroupModifierAttr : OpenMP_EnumAttr<AccessGroupModifier,
"access_group_modifier"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -396,7 +396,7 @@ def FallbackModifier : OpenMP_I32EnumAttr<
def FallbackModifierAttr : OpenMP_EnumAttr<FallbackModifier,
"fallback_modifier"> {
- let assemblyFormat = "`(` $value `)`";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -414,7 +414,7 @@ def InteropType : OpenMP_I32EnumAttr<
]>;
def InteropTypeAttr : OpenMP_EnumAttr<InteropType, "interop_type"> {
- let assemblyFormat = "`(` $value `)`";
+ 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/Shard/IR/ShardBase.td b/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
index 9af607f7bca5b..765a4fc88ad1e 100644
--- a/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
+++ b/mlir/include/mlir/Dialect/Shard/IR/ShardBase.td
@@ -63,7 +63,7 @@ def Shard_ReductionKind : I32EnumAttr<"ReductionKind",
}
def Shard_ReductionKindAttr : EnumAttr<Shard_Dialect, Shard_ReductionKind, "partial"> {
- let assemblyFormat = "$value";
+ 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/XeGPU/IR/XeGPUAttrs.td b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
index 2b977fc418553..2ac0763dfcec2 100644
--- a/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
+++ b/mlir/include/mlir/Dialect/XeGPU/IR/XeGPUAttrs.td
@@ -88,7 +88,7 @@ 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";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -129,7 +129,7 @@ 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";
+ let assemblyFormat = "`<` $value `>`";
}
//===----------------------------------------------------------------------===//
@@ -150,7 +150,7 @@ 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";
+ 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..075ba38243cd2 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
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/ACCDeclareGPUModuleInsertion.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
index e5b93e053c87d..96321e4ca1171 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
@@ -25,16 +25,16 @@
// Before:
// module {
// memref.global @arr : memref<7xf32> = dense<0.0>
-// {acc.declare = #acc.declare<dataClause = acc_create>}
+// {acc.declare = #acc.declare<dataClause = <acc_create>>}
// }
//
// After:
// module attributes {gpu.container_module} {
// memref.global @arr : memref<7xf32> = dense<0.0>
-// {acc.declare = #acc.declare<dataClause = acc_create>}
+// {acc.declare = #acc.declare<dataClause = <acc_create>>}
// gpu.module @acc_gpu_module {
// memref.global @arr : memref<7xf32> = dense<0.0>
-// {acc.declare = #acc.declare<dataClause = acc_create>}
+// {acc.declare = #acc.declare<dataClause = <acc_create>>}
// }
// }
//
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/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
index 88ff5da990377..210a502d15428 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
@@ -108,7 +108,7 @@
//
// After:
// memref.global constant @g_const : memref<f32> = dense<1.0>
-// {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// func.func @test() {
// acc.serial {
// %addr = memref.get_global @g_const : memref<f32>
@@ -129,7 +129,7 @@
//
// After:
// memref.global @g_data : memref<f32> = dense<0.0>
-// {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// acc.routine @routine_0 func(@device_func)
// func.func @device_func() attributes {acc.routine_info = ...} {
// %addr = memref.get_global @g_data : memref<f32>
@@ -157,7 +157,7 @@
//
// After:
// memref.global @g_init : memref<f32> = dense<0.0>
-// {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// acc.private.recipe @priv_recipe : memref<f32> init {
// ^bb0(%arg0: memref<f32>):
// %alloc = memref.alloc() : memref<f32>
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/Conversion/XeGPUToXeVM/loadstore_1d.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstore_1d.mlir
index d92f4f5f64df7..92f95910bb5da 100644
--- a/mlir/test/Conversion/XeGPUToXeVM/loadstore_1d.mlir
+++ b/mlir/test/Conversion/XeGPUToXeVM/loadstore_1d.mlir
@@ -24,13 +24,13 @@ gpu.module @load_store_check {
%loaded = xegpu.load_nd %src_tdesc[96] <{l1_hint = #xegpu.cache_hint<cached>, l2_hint = #xegpu.cache_hint<uncached>}>
: !xegpu.tensor_desc<32xf32> -> vector<2xf32>
- %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<256xf32> -> !xegpu.tensor_desc<32xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<256xf32> -> !xegpu.tensor_desc<32xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
// CHECK: %[[ADDR1:.*]] = arith.addi %[[INTPTR1_I64]], %[[C512]] : i64
// CHECK: %[[PTR1:.*]] = llvm.inttoptr %[[ADDR1]] : i64 to !llvm.ptr<1>
// CHECK: xevm.blockstore %[[PTR1]], %[[LOAD]] <{cache_control = #xevm.store_cache_control<L1wb_L2uc_L3wb>}>
// CHECK-SAME: : (!llvm.ptr<1>, vector<2xi32>)
xegpu.store_nd %loaded, %dst_tdesc[128] <{l1_hint = #xegpu.cache_hint<write_back>, l2_hint = #xegpu.cache_hint<uncached>}>
- : vector<2xf32>, !xegpu.tensor_desc<32xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ : vector<2xf32>, !xegpu.tensor_desc<32xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
gpu.return
}
}
diff --git a/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd.mlir
index 408950c0d8f0f..c453d227ee2b9 100644
--- a/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd.mlir
+++ b/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd.mlir
@@ -23,13 +23,13 @@ gpu.module @load_store_check {
%tid_x_f32 = arith.sitofp %tid_x_i32 : i32 to f32
%loaded_modified = vector.insert %tid_x_f32, %loaded[0] : f32 into vector<8xf32>
- %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<8x16xf32> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<8x16xf32> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
//CHECK: xevm.blockstore2d %{{.*}}, %[[W_P_BYTES]], %[[H]], %[[W_P_BYTES]], %[[ZERO]], %[[ZERO]], %{{.*}} <{
//CHECK-SAME: cache_control = #xevm.store_cache_control<L1wb_L2uc_L3wb>, elem_size_in_bits = 32 : i32,
//CHECK-SAME: tile_height = 8 : i32, tile_width = 16 : i32}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32, vector<8xi32>)
xegpu.store_nd %loaded_modified, %dst_tdesc[0, 0] <{l1_hint = #xegpu.cache_hint<write_back>, l2_hint = #xegpu.cache_hint<uncached>}>
- : vector<8xf32>, !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ : vector<8xf32>, !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
gpu.return
}
@@ -55,13 +55,13 @@ gpu.module @load_store_check {
%tid_x_f32 = arith.sitofp %tid_x_i32 : i32 to f32
%loaded_modified = vector.insert %tid_x_f32, %loaded[0] : f32 into vector<8xf32>
- %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<8x16xf32> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<8x16xf32> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
//CHECK: xevm.blockstore2d %{{.*}}, %[[W_P_BYTES]], %[[H]], %[[W_P_BYTES]], %[[ZERO]], %[[ZERO]], %{{.*}} <{
//CHECK-SAME: cache_control = #xevm.store_cache_control<Use_Default>, elem_size_in_bits = 32 : i32,
//CHECK-SAME: tile_height = 8 : i32, tile_width = 16 : i32}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32, vector<8xi32>)
xegpu.store_nd %loaded_modified, %dst_tdesc[0, 0] <{}>
- : vector<8xf32>, !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = global>>
+ : vector<8xf32>, !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <global>>>
gpu.return
}
diff --git a/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd_sub_byte.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd_sub_byte.mlir
index a8b5e695d4c38..02af4453f2aa7 100644
--- a/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd_sub_byte.mlir
+++ b/mlir/test/Conversion/XeGPUToXeVM/loadstore_nd_sub_byte.mlir
@@ -41,7 +41,7 @@ gpu.module @load_store_check {
// CHECK: %[[PAYLOAD1_DST:.*]] = vector.insert %[[C128_I32]], %[[BITCAST1_DST]] [2] : i32 into vector<8xi32>
// CHECK: %[[PAYLOAD2_DST:.*]] = vector.insert %[[C16_I32]], %[[PAYLOAD1_DST]] [3] : i32 into vector<8xi32>
// CHECK: %[[PAYLOAD3_DST:.*]] = vector.insert %[[C128_I32]], %[[PAYLOAD2_DST]] [4] : i32 into vector<8xi32>
- %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<16x128xi4> -> !xegpu.tensor_desc<8x64xi4, #xegpu.block_tdesc_attr<memory_space = global>>
+ %dst_tdesc = xegpu.create_nd_tdesc %dstte : memref<16x128xi4> -> !xegpu.tensor_desc<8x64xi4, #xegpu.block_tdesc_attr<memory_space = <global>>>
// CHECK: %[[BITCAST2_DST:.*]] = vector.bitcast %[[PAYLOAD3_DST]] : vector<8xi32> to vector<4xi64>
// CHECK: %[[DSTPTR64:.*]] = vector.extract %[[BITCAST2_DST]][0] : i64 from vector<4xi64>
@@ -51,7 +51,7 @@ gpu.module @load_store_check {
// CHECK-SAME: cache_control = #xevm.store_cache_control<L1wb_L2uc_L3wb>, elem_size_in_bits = 16 : i32,
// CHECK-SAME: tile_height = 8 : i32, tile_width = 16 : i32}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32, vector<8xi16>)
xegpu.store_nd %loaded, %dst_tdesc[8, 64] <{l1_hint = #xegpu.cache_hint<write_back>, l2_hint = #xegpu.cache_hint<uncached>}>
- : vector<32xi4>, !xegpu.tensor_desc<8x64xi4, #xegpu.block_tdesc_attr<memory_space = global>>
+ : vector<32xi4>, !xegpu.tensor_desc<8x64xi4, #xegpu.block_tdesc_attr<memory_space = <global>>>
gpu.return
}
diff --git a/mlir/test/Conversion/XeGPUToXeVM/prefetch_nd.mlir b/mlir/test/Conversion/XeGPUToXeVM/prefetch_nd.mlir
index 4163fccc9934f..87bc3875ecb7c 100644
--- a/mlir/test/Conversion/XeGPUToXeVM/prefetch_nd.mlir
+++ b/mlir/test/Conversion/XeGPUToXeVM/prefetch_nd.mlir
@@ -8,7 +8,7 @@ gpu.module @prefetch_nd_check {
// CHECK: %[[BASE_H:.*]] = arith.constant 8 : i32
%srcce = memref.memory_space_cast %src : memref<8x16xf32, 1> to memref<8x16xf32>
%src_tdesc = xegpu.create_nd_tdesc %srcce : memref<8x16xf32> -> !xegpu.tensor_desc<8x16xf32,
- #xegpu.block_tdesc_attr<memory_space = global>, #xegpu.layout<lane_layout = [1, 16], lane_data = [1, 1]>>
+ #xegpu.block_tdesc_attr<memory_space = <global>>, #xegpu.layout<lane_layout = [1, 16], lane_data = [1, 1]>>
//CHECK: %[[LLVMPTR:.*]] = llvm.inttoptr %{{.*}} : i64 to !llvm.ptr<1>
//CHECK: xevm.blockprefetch2d %[[LLVMPTR]], %[[BASE_WIDTH_PITCH_BYTES]], %[[BASE_H]],
@@ -17,7 +17,7 @@ gpu.module @prefetch_nd_check {
//CHECK-SAME: tile_height = 8 : i32, tile_width = 16 : i32, v_blocks = 1 : i32}>
//CHECK-SAME: : (!llvm.ptr<1>, i32, i32, i32, i32, i32)
xegpu.prefetch_nd %src_tdesc[0, 0] <{l1_hint = #xegpu.cache_hint<cached>, l2_hint = #xegpu.cache_hint<uncached>}>
- : !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = global>,
+ : !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <global>>,
#xegpu.layout<lane_layout = [1, 16], lane_data = [1, 1]>>
gpu.return
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/acc-declare-ctor-dtor-conversion.mlir b/mlir/test/Dialect/OpenACC/acc-declare-ctor-dtor-conversion.mlir
index ad5e8414a3497..18a4e8ae1782b 100644
--- a/mlir/test/Dialect/OpenACC/acc-declare-ctor-dtor-conversion.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-declare-ctor-dtor-conversion.mlir
@@ -17,35 +17,35 @@
// NODTOR: llvm.mlir.global_ctors ctors = [@arr_acc_ctor, @other_arr_acc_ctor], priorities = [102 : i32, 102 : i32], data = [#llvm.zero, #llvm.zero]
// NODTOR-NOT: llvm.mlir.global_dtors
-llvm.mlir.global external @arr() {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.array<7 x f32> {
+llvm.mlir.global external @arr() {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.array<7 x f32> {
%0 = llvm.mlir.zero : !llvm.array<7 x f32>
llvm.return %0 : !llvm.array<7 x f32>
}
-llvm.mlir.global external @other_arr() {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.array<3 x f32> {
+llvm.mlir.global external @other_arr() {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.array<3 x f32> {
%0 = llvm.mlir.zero : !llvm.array<3 x f32>
llvm.return %0 : !llvm.array<3 x f32>
}
acc.global_ctor @arr_acc_ctor {
- %0 = llvm.mlir.addressof @arr {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+ %0 = llvm.mlir.addressof @arr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
%1 = acc.create varPtr(%0 : !llvm.ptr) varType(!llvm.array<7 x f32>) -> !llvm.ptr
acc.declare_enter dataOperands(%1 : !llvm.ptr)
acc.terminator
}
acc.global_ctor @other_arr_acc_ctor {
- %0 = llvm.mlir.addressof @other_arr {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+ %0 = llvm.mlir.addressof @other_arr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
%1 = acc.create varPtr(%0 : !llvm.ptr) varType(!llvm.array<3 x f32>) -> !llvm.ptr
acc.declare_enter dataOperands(%1 : !llvm.ptr)
acc.terminator
}
acc.global_dtor @arr_acc_dtor {
- %0 = llvm.mlir.addressof @arr {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+ %0 = llvm.mlir.addressof @arr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
%1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) varType(!llvm.array<7 x f32>) dataClause(acc_create) -> !llvm.ptr
acc.declare_exit dataOperands(%1 : !llvm.ptr)
acc.delete accPtr(%1 : !llvm.ptr)
acc.terminator
}
acc.global_dtor @other_arr_acc_dtor {
- %0 = llvm.mlir.addressof @other_arr {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+ %0 = llvm.mlir.addressof @other_arr {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
%1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) varType(!llvm.array<3 x f32>) dataClause(acc_create) -> !llvm.ptr
acc.declare_exit dataOperands(%1 : !llvm.ptr)
acc.delete accPtr(%1 : !llvm.ptr)
@@ -60,7 +60,7 @@ acc.global_dtor @other_arr_acc_dtor {
// CHECK: llvm.func internal @merged_acc_ctor
// CHECK: llvm.mlir.global_ctors ctors = [@existing_ctor, @merged_acc_ctor], priorities = [0 : i32, 102 : i32], data = [#llvm.zero, #llvm.zero]
-llvm.mlir.global external @merged_var() {acc.declare = #acc.declare<dataClause = acc_create>} : i32 {
+llvm.mlir.global external @merged_var() {acc.declare = #acc.declare<dataClause = <acc_create>>} : i32 {
%0 = llvm.mlir.constant(0 : i32) : i32
llvm.return %0 : i32
}
@@ -69,7 +69,7 @@ llvm.func internal @existing_ctor() {
}
llvm.mlir.global_ctors ctors = [@existing_ctor], priorities = [0 : i32], data = [#llvm.zero]
acc.global_ctor @merged_acc_ctor {
- %0 = llvm.mlir.addressof @merged_var {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+ %0 = llvm.mlir.addressof @merged_var {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
%1 = acc.create varPtr(%0 : !llvm.ptr) varType(i32) -> !llvm.ptr
acc.declare_enter dataOperands(%1 : !llvm.ptr)
acc.terminator
diff --git a/mlir/test/Dialect/OpenACC/acc-declare-gpu-module-insertion.mlir b/mlir/test/Dialect/OpenACC/acc-declare-gpu-module-insertion.mlir
index 2e5910d155641..f4fed5692d47d 100644
--- a/mlir/test/Dialect/OpenACC/acc-declare-gpu-module-insertion.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-declare-gpu-module-insertion.mlir
@@ -4,13 +4,13 @@
// The host global stays in the module; a copy is inserted into the GPU module.
// CHECK-LABEL: module
-// CHECK: memref.global @arr {{.*}} {acc.declare = #acc.declare<dataClause = acc_create>}
+// CHECK: memref.global @arr {{.*}} {acc.declare = #acc.declare<dataClause = <acc_create>>}
// CHECK: gpu.module @acc_gpu_module {
-// CHECK: memref.global @arr {{.*}} {acc.declare = #acc.declare<dataClause = acc_create>}
+// CHECK: memref.global @arr {{.*}} {acc.declare = #acc.declare<dataClause = <acc_create>>}
// CHECK: }
module {
- memref.global @arr : memref<7xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_create>}
+ memref.global @arr : memref<7xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = <acc_create>>}
}
// -----
@@ -18,13 +18,13 @@ module {
// If the GPU module already has the global (e.g. from CUDA Fortran pass before
// ACCImplicitDeclare marked the host), reuse it and propagate acc.declare.
// CHECK-LABEL: module attributes {gpu.container_module}
-// CHECK: memref.global @precloned {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global @precloned {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// CHECK: gpu.module @acc_gpu_module {
-// CHECK-NEXT: memref.global @precloned {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK-NEXT: memref.global @precloned {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// CHECK-NEXT: }
module attributes {gpu.container_module} {
- memref.global @precloned : memref<4xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_copyin>}
+ memref.global @precloned : memref<4xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
gpu.module @acc_gpu_module {
memref.global @precloned : memref<4xf32> = dense<0.0>
}
diff --git a/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir b/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
index 1ad0abb78b9ec..5dbd4c4b50eee 100644
--- a/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
@@ -265,7 +265,7 @@ func.func @test_device_global_in_parallel() {
// Test memref.view tagged with acc.declare deviceptr and used directly in region.
func.func @test_declare_deviceptr_arg_in_parallel(%arg0: memref<?xi8>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] {acc.declare = #acc.declare<dataClause = acc_deviceptr>} : memref<?xi8> to memref<10xf32>
+ %view = memref.view %arg0[%c0][] {acc.declare = #acc.declare<dataClause = <acc_deviceptr>>} : memref<?xi8> to memref<10xf32>
%devptr = acc.deviceptr varPtr(%view : memref<10xf32>) name("arg0") -> memref<10xf32>
%token = acc.declare_enter dataOperands(%devptr : memref<10xf32>)
acc.parallel {
@@ -278,7 +278,7 @@ func.func @test_declare_deviceptr_arg_in_parallel(%arg0: memref<?xi8>) {
}
// CHECK-LABEL: func.func @test_declare_deviceptr_arg_in_parallel
-// CHECK: %[[VIEW:.*]] = memref.view %{{.*}}[{{.*}}][] {acc.declare = #acc.declare<dataClause = acc_deviceptr>} : memref<?xi8> to memref<10xf32>
+// CHECK: %[[VIEW:.*]] = memref.view %{{.*}}[{{.*}}][] {acc.declare = #acc.declare<dataClause = <acc_deviceptr>>} : memref<?xi8> to memref<10xf32>
// CHECK: %[[DEVPTR:.*]] = acc.deviceptr varPtr(%[[VIEW]] : memref<10xf32>) name("arg0") -> memref<10xf32>
// CHECK: %[[TOKEN:.*]] = acc.declare_enter dataOperands(%[[DEVPTR]] : memref<10xf32>)
// CHECK: %[[IMPLICIT_DEVPTR:.*]] = acc.deviceptr varPtr(%{{.*}} : memref<?xi8>) implicit(true) name("") -> memref<?xi8>
diff --git a/mlir/test/Dialect/OpenACC/acc-implicit-declare.mlir b/mlir/test/Dialect/OpenACC/acc-implicit-declare.mlir
index 0998f1e9683e0..9f80b8a9168f4 100644
--- a/mlir/test/Dialect/OpenACC/acc-implicit-declare.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-implicit-declare.mlir
@@ -39,7 +39,7 @@ func.func @test_constant_in_serial() {
}
// This is expected to be `acc declare`'d since it is a constant.
-// CHECK: memref.global constant @gscalarconst {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global constant @gscalarconst {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// -----
@@ -79,7 +79,7 @@ func.func @test_scalar_in_accroutine() attributes {acc.routine_info = #acc.routi
}
// Global should be acc declare'd because it's in an acc routine
-// CHECK: memref.global @gscalar_routine {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global @gscalar_routine {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// -----
@@ -94,7 +94,7 @@ func.func @test_constant_in_accroutine() attributes {acc.routine_info = #acc.rou
return
}
-// CHECK: memref.global constant @gscalarconst_routine {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global constant @gscalarconst_routine {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// -----
@@ -127,7 +127,7 @@ func.func @test_private_recipe_referenced() {
}
// Global should be acc declare'd because the recipe is referenced
-// CHECK: memref.global @global_for_private {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global @global_for_private {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// -----
@@ -190,7 +190,7 @@ func.func @test_multiple_constructs() {
// Non-constant globals ARE hoisted before their compute regions
// Constant global should be marked with acc.declare
-// CHECK: memref.global constant @global_serial_const {{.*}} {acc.declare = #acc.declare<dataClause = acc_copyin>}
+// CHECK: memref.global constant @global_serial_const {{.*}} {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
// CHECK-LABEL: func.func @test_multiple_constructs
// CHECK: memref.get_global @global_parallel
// CHECK-NEXT: acc.parallel
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/offload-livein-value-canonicalization.mlir b/mlir/test/Dialect/OpenACC/offload-livein-value-canonicalization.mlir
index 7be62a789e6fa..4df55e728208a 100644
--- a/mlir/test/Dialect/OpenACC/offload-livein-value-canonicalization.mlir
+++ b/mlir/test/Dialect/OpenACC/offload-livein-value-canonicalization.mlir
@@ -92,7 +92,7 @@ func.func @test_accbounds_rematerialize(%arg0: memref<10xf32>) {
// -----
// Test memref.get_global with acc.declare sinking
-memref.global @memref_global_with_declare : memref<10xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_copyin>}
+memref.global @memref_global_with_declare : memref<10xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
func.func private @use_memref(memref<10xf32>) -> ()
@@ -113,7 +113,7 @@ func.func @test_memref_get_global_sink() {
// -----
// Test memref.reinterpret_cast traces through to get_global
-memref.global @memref_global_reinterpret : memref<2x5xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_copyin>}
+memref.global @memref_global_reinterpret : memref<2x5xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = <acc_copyin>>}
func.func private @use_memref_1d(memref<10xf32>) -> ()
diff --git a/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir b/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir
index aca6e22bae99a..188d67d213438 100644
--- a/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir
+++ b/mlir/test/Dialect/OpenACC/offload-target-verifier.mlir
@@ -133,7 +133,7 @@ func.func @test_memref_gpu_address_space() {
// -----
// Test global symbol with acc.declare attribute - should pass
-memref.global @global_array_declared : memref<10xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = acc_create>}
+memref.global @global_array_declared : memref<10xf32> = dense<0.0> {acc.declare = #acc.declare<dataClause = <acc_create>>}
func.func @test_global_symbol_with_declare() {
// expected-remark @below {{passed validity check}}
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir b/mlir/test/Dialect/OpenACC/ops.mlir
index 24668cf18c6d5..dc412a47842bd 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>>
// -----
@@ -1760,20 +1760,20 @@ func.func @testdeclareop(%a: memref<f32>, %b: memref<f32>, %c: memref<f32>) -> (
// -----
-llvm.mlir.global external @globalvar() { acc.declare = #acc.declare<dataClause = acc_create> } : i32 {
+llvm.mlir.global external @globalvar() { acc.declare = #acc.declare<dataClause = <acc_create>> } : i32 {
%0 = llvm.mlir.constant(0 : i32) : i32
llvm.return %0 : i32
}
acc.global_ctor @acc_constructor {
- %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr
+ %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = <acc_create>> } : !llvm.ptr
%1 = acc.create varPtr(%0 : !llvm.ptr) varType(i32) -> !llvm.ptr
acc.declare_enter dataOperands(%1 : !llvm.ptr)
acc.terminator
}
acc.global_dtor @acc_destructor {
- %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr
+ %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = <acc_create>> } : !llvm.ptr
%1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) varType(i32) dataClause(acc_create) -> !llvm.ptr<0>
acc.declare_exit dataOperands(%1 : !llvm.ptr)
acc.delete accPtr(%1 : !llvm.ptr)
@@ -1781,11 +1781,11 @@ acc.global_dtor @acc_destructor {
}
// CHECK-LABEL: acc.global_ctor @acc_constructor
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[ADDR]] : !llvm.ptr) varType(i32) -> !llvm.ptr
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !llvm.ptr)
// CHECK: acc.global_dtor @acc_destructor
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = <acc_create>>} : !llvm.ptr
// CHECK-NEXT: %[[DELETE:.*]] = acc.getdeviceptr varPtr(%[[ADDR]] : !llvm.ptr) varType(i32) dataClause(acc_create) -> !llvm.ptr
// CHECK-NEXT: acc.declare_exit dataOperands(%[[DELETE]] : !llvm.ptr)
// CHECK-NEXT: acc.delete accPtr(%[[DELETE]] : !llvm.ptr)
@@ -1893,7 +1893,7 @@ func.func @compute3(%a: memref<10x10xf32>, %b: memref<10x10xf32>, %c: memref<10x
%numWorkers = arith.constant 10 : i64
%c20 = arith.constant 20 : i32
- %alloc = llvm.alloca %c20 x i32 { acc.declare = #acc.declare<dataClause = acc_create, implicit = true> } : (i32) -> !llvm.ptr
+ %alloc = llvm.alloca %c20 x i32 { acc.declare = #acc.declare<dataClause = <acc_create>, implicit = true> } : (i32) -> !llvm.ptr
%createlocal = acc.create varPtr(%alloc : !llvm.ptr) varType(!llvm.array<20 x i32>) implicit(true) -> !llvm.ptr<0>
%pa = acc.present varPtr(%a : memref<10x10xf32>) varType(tensor<10x10xf32>) -> memref<10x10xf32>
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/Dialect/XeGPU/invalid.mlir b/mlir/test/Dialect/XeGPU/invalid.mlir
index f06fbac607474..293e9e32a291b 100644
--- a/mlir/test/Dialect/XeGPU/invalid.mlir
+++ b/mlir/test/Dialect/XeGPU/invalid.mlir
@@ -30,7 +30,7 @@ func.func @create_nd_tdesc_2(%src: memref<24x32xf32>) {
// -----
func.func @create_nd_tdesc_3(%src: memref<2x24x32xf32, 3>) {
// expected-error at +1 {{SLM is only supported for 1D block tensor}}
- %1 = xegpu.create_nd_tdesc %src : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = slm>>
+ %1 = xegpu.create_nd_tdesc %src : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<8x16xf32, #xegpu.block_tdesc_attr<memory_space = <slm>>>
return
}
diff --git a/mlir/test/Dialect/XeGPU/ops.mlir b/mlir/test/Dialect/XeGPU/ops.mlir
index 5e795017e2122..04ad56dc9dfb6 100644
--- a/mlir/test/Dialect/XeGPU/ops.mlir
+++ b/mlir/test/Dialect/XeGPU/ops.mlir
@@ -38,8 +38,8 @@ gpu.func @create_nd_tdesc_4(%src: memref<2x24x32xf32>) {
// CHECK: gpu.func @create_nd_tdesc_5(%[[arg0:.*]]: memref<2x24x32xf32, 3>) {
gpu.func @create_nd_tdesc_5(%src: memref<2x24x32xf32, 3>) {
- // CHECK: %[[REG:.*]] = xegpu.create_nd_tdesc %arg0 : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<16xf32, #xegpu.block_tdesc_attr<memory_space = slm>>
- %1 = xegpu.create_nd_tdesc %src : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<16xf32, #xegpu.block_tdesc_attr<memory_space = slm>>
+ // CHECK: %[[REG:.*]] = xegpu.create_nd_tdesc %arg0 : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<16xf32, #xegpu.block_tdesc_attr<memory_space = <slm>>>
+ %1 = xegpu.create_nd_tdesc %src : memref<2x24x32xf32, 3> -> !xegpu.tensor_desc<16xf32, #xegpu.block_tdesc_attr<memory_space = <slm>>>
gpu.return
}
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..5752a781553ee 100644
--- a/mlir/test/lib/Dialect/Test/TestAttrDefs.td
+++ b/mlir/test/lib/Dialect/Test/TestAttrDefs.td
@@ -352,17 +352,14 @@ 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">;
diff --git a/mlir/test/lib/Dialect/Test/TestOps.td b/mlir/test/lib/Dialect/Test/TestOps.td
index 3060de7ac2b4b..60d3a3fbd6c2c 100644
--- a/mlir/test/lib/Dialect/Test/TestOps.td
+++ b/mlir/test/lib/Dialect/Test/TestOps.td
@@ -497,7 +497,7 @@ def TestPrettyEnumAttr
// 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
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..5e18dc664e95d 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>
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..2f958021129bd 100644
--- a/mlir/test/mlir-tblgen/op-format.td
+++ b/mlir/test/mlir-tblgen/op-format.td
@@ -152,7 +152,7 @@ def TestNonKeywordEnumAttr
// 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 +160,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