Integrate LLVM to llvm/llvm-project@e89dc2172f0d (#24863)
Still carries a temporary revert of llvm/llvm-project@202ece62657c from
`sm-iree-integrates/llvm-20260821`. Should be dropped in the next
integrate via llvm/llvm-project#218693.
This cherry-picks an additional hotfix: llvm/llvm-project#219301.
Adaptations required:
llvm/llvm-project@979b72237338 - `SymbolOpInterface` no longer provides
implicit symbol traits, so a bunch of ops need updating with name and
visibility properties
llvm/llvm-project@180a0fc57c88 - symbol visibility query API has changed
llvm/llvm-project@325a853c1264 - using lowered result type for MLIR LLVM
int constants
llvm/llvm-project@8fe50937417b - attribute ordering for named Linalg ops
llvm/llvm-project@ad9dc60252ed - discardable/inherent attribute ordering
llvm/llvm-project@bfd1a003fb33 - `alloc_tensor` properties separated out
llvm/llvm-project@3e3927f7166c - change in printer defaults for property
dictionaries
Adaptations to llvm/llvm-project@979b72237338 affect Torch-MLIR, so the
`iree-org` fork was advanced to `sm-iree-integrates/torch-20260827`
accordingly.
Assisted-by: OpenAI Codex
---------
Signed-off-by: Artem Gindinson <gindinson@roofline.ai>
diff --git a/compiler/plugins/target/ROCM/test/lower_rocm_ukernel_descriptor.mlir b/compiler/plugins/target/ROCM/test/lower_rocm_ukernel_descriptor.mlir
index 542ff02..7dcf15a 100644
--- a/compiler/plugins/target/ROCM/test/lower_rocm_ukernel_descriptor.mlir
+++ b/compiler/plugins/target/ROCM/test/lower_rocm_ukernel_descriptor.mlir
@@ -74,7 +74,7 @@
// CHECK-SAME: %[[ARG0:[a-zA-Z0-9]+]]: tensor<1x2x8x4x16x2x8xi8>
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]: tensor<1x2x4x2x4x16x2x8xi8>
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9]+]]: tensor<1x1x4x8x2x4x16x4xi32>
-// CHECK: %[[ALLOC:.*]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<8192xi8>
+// CHECK: %[[ALLOC:.*]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<8192xi8>
// CHECK: %[[C1_INDEX:.*]] = arith.constant 1 : index
// CHECK: %[[DIM:.*]] = tensor.dim %[[ARG0]], %[[C1_INDEX]] : tensor<1x2x8x4x16x2x8xi8>
// CHECK: %[[DIM_CAST:.*]] = arith.index_cast %[[DIM]] : index to i32
diff --git a/compiler/src/iree/compiler/API/api_exports.c b/compiler/src/iree/compiler/API/api_exports.c
index e0e31c3..eae3d43 100644
--- a/compiler/src/iree/compiler/API/api_exports.c
+++ b/compiler/src/iree/compiler/API/api_exports.c
@@ -1077,8 +1077,8 @@
extern void mlirSymbolTableCreate();
extern void mlirSymbolTableDestroy();
extern void mlirSymbolTableErase();
+extern void mlirSymbolTableGetDefaultVisibilityAttributeName();
extern void mlirSymbolTableGetSymbolAttributeName();
-extern void mlirSymbolTableGetVisibilityAttributeName();
extern void mlirSymbolTableInsert();
extern void mlirSymbolTableLookup();
extern void mlirSymbolTableReplaceAllSymbolUses();
@@ -2319,8 +2319,8 @@
x += (uintptr_t)&mlirSymbolTableCreate;
x += (uintptr_t)&mlirSymbolTableDestroy;
x += (uintptr_t)&mlirSymbolTableErase;
+ x += (uintptr_t)&mlirSymbolTableGetDefaultVisibilityAttributeName;
x += (uintptr_t)&mlirSymbolTableGetSymbolAttributeName;
- x += (uintptr_t)&mlirSymbolTableGetVisibilityAttributeName;
x += (uintptr_t)&mlirSymbolTableInsert;
x += (uintptr_t)&mlirSymbolTableLookup;
x += (uintptr_t)&mlirSymbolTableReplaceAllSymbolUses;
diff --git a/compiler/src/iree/compiler/API/api_exports.def b/compiler/src/iree/compiler/API/api_exports.def
index a4b5fe7..f774f6a 100644
--- a/compiler/src/iree/compiler/API/api_exports.def
+++ b/compiler/src/iree/compiler/API/api_exports.def
@@ -1067,8 +1067,8 @@
mlirSymbolTableCreate
mlirSymbolTableDestroy
mlirSymbolTableErase
+ mlirSymbolTableGetDefaultVisibilityAttributeName
mlirSymbolTableGetSymbolAttributeName
- mlirSymbolTableGetVisibilityAttributeName
mlirSymbolTableInsert
mlirSymbolTableLookup
mlirSymbolTableReplaceAllSymbolUses
diff --git a/compiler/src/iree/compiler/API/api_exports.ld b/compiler/src/iree/compiler/API/api_exports.ld
index bf447bd..b7e32c1 100644
--- a/compiler/src/iree/compiler/API/api_exports.ld
+++ b/compiler/src/iree/compiler/API/api_exports.ld
@@ -1068,8 +1068,8 @@
mlirSymbolTableCreate;
mlirSymbolTableDestroy;
mlirSymbolTableErase;
+ mlirSymbolTableGetDefaultVisibilityAttributeName;
mlirSymbolTableGetSymbolAttributeName;
- mlirSymbolTableGetVisibilityAttributeName;
mlirSymbolTableInsert;
mlirSymbolTableLookup;
mlirSymbolTableReplaceAllSymbolUses;
diff --git a/compiler/src/iree/compiler/API/api_exports.macos.lst b/compiler/src/iree/compiler/API/api_exports.macos.lst
index 1db3db7..b3229c5 100644
--- a/compiler/src/iree/compiler/API/api_exports.macos.lst
+++ b/compiler/src/iree/compiler/API/api_exports.macos.lst
@@ -1066,8 +1066,8 @@
_mlirSymbolTableCreate
_mlirSymbolTableDestroy
_mlirSymbolTableErase
+_mlirSymbolTableGetDefaultVisibilityAttributeName
_mlirSymbolTableGetSymbolAttributeName
-_mlirSymbolTableGetVisibilityAttributeName
_mlirSymbolTableInsert
_mlirSymbolTableLookup
_mlirSymbolTableReplaceAllSymbolUses
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_alloc_private_memory_for_dps_ops.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_alloc_private_memory_for_dps_ops.mlir
index c91370d..670ac9c 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_alloc_private_memory_for_dps_ops.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_alloc_private_memory_for_dps_ops.mlir
@@ -2,7 +2,7 @@
// RUN: --pass-pipeline="builtin.module(func.func(iree-codegen-gpu-alloc-private-memory-for-dps-ops))" | FileCheck %s
// CHECK-LABEL: func.func @unused_result_copied
-// CHECK-DAG: %[[SRC:.*]] = bufferization.alloc_tensor() copy(%{{.*}}) {memory_space = #gpu.address_space<private>} : tensor<1x10xf32>
+// CHECK-DAG: %[[SRC:.*]] = bufferization.alloc_tensor() copy(%{{.*}}) <{memory_space = #gpu.address_space<private>}> : tensor<1x10xf32>
// CHECK-DAG: iree_linalg_ext.sort{{.*}} dimension(1) outs(%[[SRC]], %{{.*}} : tensor<1x10xf32>, tensor<1x10xi64>)
func.func @unused_result_copied(%arg0: !iree_tensor_ext.dispatch.tensor<readonly:tensor<1x10xf32>>, %arg1: tensor<1x10xi64>) -> tensor<1x10xi64> {
%2 = iree_tensor_ext.dispatch.tensor.load %arg0, offsets = [0, 0], sizes = [1, 10], strides = [1, 1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<1x10xf32>> -> tensor<1x10xf32>
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir
index f5e1b1f..ff1f36c 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir
@@ -1005,9 +1005,9 @@
}
// CHECK-LABEL: func @swizzle_with_fusion(
-// CHECK: %[[ALLOC1:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf16>
+// CHECK: %[[ALLOC1:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf16>
// CHECK: %[[SWIZZLE1:.+]] = iree_codegen.swizzle_hint %[[ALLOC1]][#iree_codegen.xor_shuffle<64, 8>] : tensor<128x128xf16>
-// CHECK: %[[ALLOC2:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf16>
+// CHECK: %[[ALLOC2:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf16>
// CHECK: %[[SWIZZLE2:.+]] = iree_codegen.swizzle_hint %[[ALLOC2]][#iree_codegen.xor_shuffle<64, 16>] : tensor<128x128xf16>
// CHECK: %[[EMPTY:.+]] = tensor.empty() : tensor<128x128xf16>
// CHECK: scf.forall {{.*}} shared_outs(%[[OUT:.+]] = %[[EMPTY]]) -> (tensor<128x128xf16>) {
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_infer_memory_space.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_infer_memory_space.mlir
index 48d1c2d..29a68c0 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_infer_memory_space.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_infer_memory_space.mlir
@@ -15,7 +15,7 @@
}
// CHECK-LABEL: func @write_in_lane_forall
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<private>}
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<private>}>
// CHECK: vector.transfer_write %{{.*}}, %[[ALLOC]]
// -----
@@ -31,24 +31,24 @@
}
// CHECK-LABEL: func @forall_shared_dest
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>}
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}>
// CHECK: scf.forall {{.*}} shared_outs(%{{.*}} = %[[ALLOC]])
// -----
func.func @already_annotated_alloc() -> tensor<2x3xi32> {
- %alloc = bufferization.alloc_tensor() {memory_space = #gpu.address_space<private>} : tensor<2x3xi32>
+ %alloc = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<private>}> : tensor<2x3xi32>
return %alloc : tensor<2x3xi32>
}
// CHECK-LABEL: func @already_annotated_alloc
-// CHECK: bufferization.alloc_tensor() {memory_space = #gpu.address_space<private>}
+// CHECK: bufferization.alloc_tensor() <{memory_space = #gpu.address_space<private>}>
// -----
// expected-error@+1 {{failed to set the gpu memory space for all `bufferization.alloc_tensor` ops}}
func.func @unknown_memory_space() -> tensor<2x3xi32> {
// expected-error@+1 {{unexpected gpu memory space must be private or workgroup.}}
- %alloc = bufferization.alloc_tensor() {memory_space = "bad"} : tensor<2x3xi32>
+ %alloc = bufferization.alloc_tensor() <{memory_space = "bad"}> : tensor<2x3xi32>
return %alloc : tensor<2x3xi32>
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_alloc.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_alloc.mlir
index 1f76718..9a7eb5e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_alloc.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_vector_alloc.mlir
@@ -28,7 +28,7 @@
// CHECK: gpu.barrier memfence [#gpu.address_space<workgroup>]
// CHECK: %[[GLOBAL:.+]] = vector.transfer_read %{{.*}} : memref<16x16xf16>, vector<16x16xf16>
// CHECK: %[[READ_LAYOUT_VALUE:.+]] = iree_vector_ext.to_layout %[[GLOBAL]] to layout(#[[$READ_LAYOUT]]) : vector<16x16xf16>
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<16x16xf16, #gpu.address_space<workgroup>>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<16x16xf16, #gpu.address_space<workgroup>>
// CHECK: %[[WRITE:.+]] = vector.transfer_write %[[READ_LAYOUT_VALUE]], %[[ALLOC]]
// CHECK: %[[BARRIER:.+]] = iree_gpu.value_barrier %[[WRITE]]
// CHECK: %[[LDS_READ:.+]] = vector.transfer_read %[[BARRIER]]
@@ -100,7 +100,7 @@
// CHECK: gpu.barrier memfence [#gpu.address_space<workgroup>]
// CHECK: %[[GLOBAL:.+]] = vector.gather %{{.*}} : memref<16x16xf16>, vector<16x16xindex>, vector<16x16xi1>, vector<16x16xf16> into vector<16x16xf16>
// CHECK: %[[READ_LAYOUT_VALUE:.+]] = iree_vector_ext.to_layout %[[GLOBAL]] to layout(#[[$GATHER_READ_LAYOUT]]) : vector<16x16xf16>
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<16x16xf16, #gpu.address_space<workgroup>>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<16x16xf16, #gpu.address_space<workgroup>>
// CHECK: %[[WRITE:.+]] = vector.transfer_write %[[READ_LAYOUT_VALUE]], %[[ALLOC]]
// CHECK: %[[BARRIER:.+]] = iree_gpu.value_barrier %[[WRITE]]
// CHECK: %[[LDS_READ:.+]] = vector.transfer_read %[[BARRIER]]
@@ -179,7 +179,7 @@
// CHECK-LABEL: func.func @materialize_layout_conflict
// CHECK: gpu.barrier memfence [#gpu.address_space<workgroup>]
// CHECK: %[[A:.+]] = iree_vector_ext.to_layout %{{.*}} to layout(#[[$CONFLICT_LAYOUT_A]]) : vector<16x16xf16>
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<16x16xf16, #gpu.address_space<workgroup>>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<16x16xf16, #gpu.address_space<workgroup>>
// CHECK: %[[WRITE:.+]] = vector.transfer_write %[[A]], %[[ALLOC]]
// CHECK: %[[BARRIER:.+]] = iree_gpu.value_barrier %[[WRITE]]
// CHECK: %[[LDS_READ:.+]] = vector.transfer_read %[[BARRIER]]
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUOps.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUOps.td
index 9e95c7d..bc5479f 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUOps.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUOps.td
@@ -61,7 +61,7 @@
```mlir
%0 = scf.forall (%idy, %idx) in (8, 8) -> (tensor<4x128xf32>) {
- %alloc = bufferization.alloc_tensor {memory_space = #gpu.address_space<workgroup>}
+ %alloc = bufferization.alloc_tensor <{memory_space = #gpu.address_space<workgroup>}>
: tensor<4x128xf32>
%barrier = iree_gpu.barrier_region %alloc {
^bb0(%shared: tensor<4x128xf32>):
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/transform_fuse_forall.mlir b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/transform_fuse_forall.mlir
index 6834cb4..04994f6 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/transform_fuse_forall.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/transform_fuse_forall.mlir
@@ -45,7 +45,7 @@
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<128x128xf32>
// CHECK-DAG: %[[EMPTY:.+]] = tensor.empty() : tensor<128x128xf32>
-// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[IDX:.+]], %[[IDY:.+]]) in (8, 8) shared_outs(%[[INIT:.+]] = %[[EMPTY]]) -> (tensor<128x128xf32>) {
// CHECK: %[[BARRIER:.+]] = iree_gpu.barrier_region ins(%[[ALLOC]] : tensor<128x128xf32>)
@@ -121,7 +121,7 @@
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<128x128xf32>
// CHECK-DAG: %[[EMPTY:.+]] = tensor.empty() : tensor<128x128xf32>
-// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[IDX:.+]], %[[IDY:.+]]) in (8, 8) shared_outs(%[[INIT:.+]] = %[[EMPTY]]) -> (tensor<128x128xf32>) {
// CHECK: %[[SHUFFLE:.+]] = iree_gpu.barrier_region ins(%[[ALLOC]] : tensor<128x128xf32>)
// CHECK: %[[LOOP:.+]] = scf.for {{.*}} iter_args(%[[INIT:.+]] = %{{.*}})
@@ -178,7 +178,7 @@
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<128x128xf32>
// CHECK-DAG: %[[EMPTY:.+]] = tensor.empty() : tensor<128x128xf32>
-// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[IDX:.+]], %[[IDY:.+]]) in (8, 8) shared_outs(%[[INIT:.+]] = %[[EMPTY]]) -> (tensor<128x128xf32>) {
// CHECK: %[[BARRIER:.+]] = iree_gpu.barrier_region ins(%[[ALLOC]] : tensor<128x128xf32>)
// CHECK: ^bb0(%[[INTERMEDIATE:.+]]: tensor<128x128xf32>):
@@ -244,7 +244,7 @@
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<128x128xf32>
// CHECK-DAG: %[[EMPTY:.+]] = tensor.empty() : tensor<128x128xf32>
-// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK-DAG: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[W_IDX:.+]], %[[W_IDY:.+]]) in (2, 2) shared_outs(%[[INIT:.+]] = %[[EMPTY]]) -> (tensor<128x128xf32>) {
// CHECK: scf.forall (%[[L_IDX:.+]], %[[L_IDY:.+]]) in (4, 4) {{.*}} -> (tensor<64x64xf32>)
@@ -306,7 +306,7 @@
// CHECK-LABEL: func @fuse_forall_different_thread_count
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<128x128xf32>
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[IDX:.+]], %[[IDY:.+]]) in (8, 8) {{.*}} -> (tensor<128x128xf32>) {
// CHECK: iree_gpu.barrier_region ins(%[[ALLOC]]
// CHECK: %[[LINEARID:.+]] = affine.apply #[[$MAP1]](%[[IDX]], %[[IDY]])
@@ -361,7 +361,7 @@
// CHECK-SAME: %[[Y:[A-Za-z0-9]+]]: index
// CHECK-SAME: %[[Z:[A-Za-z0-9]+]]: index
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() {memory_space = #gpu.address_space<workgroup>} : tensor<128x128xf32>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor() <{memory_space = #gpu.address_space<workgroup>}> : tensor<128x128xf32>
// CHECK: scf.forall (%[[IDX:.+]], %[[IDY:.+]]) in (8, 8) {{.*}} -> (tensor<128x128xf32>) {
// CHECK: iree_gpu.barrier_region ins(%[[ALLOC]]
// CHECK-DAG: %[[LINEARID:.+]] = affine.apply #[[$MAP1]](%[[IDX]], %[[IDY]])
@@ -413,7 +413,7 @@
// CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]: tensor<?x128xf32>
// CHECK-SAME: %[[SIZE:[A-Za-z0-9]+]]: index
-// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor(%[[SIZE]])
-// CHECK-SAME: memory_space = #gpu.address_space<workgroup>} : tensor<?x128xf32>
+// CHECK: %[[ALLOC:.+]] = bufferization.alloc_tensor(%[[SIZE]]) <{
+// CHECK-SAME: memory_space = #gpu.address_space<workgroup>}> : tensor<?x128xf32>
// CHECK: scf.forall ({{.*}}) in (8, 8) {{.*}} -> (tensor<128x128xf32>) {
// CHECK: iree_gpu.barrier_region ins(%[[ALLOC]]
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/PCF/ExternalInterfaces/test/bufferize.mlir b/compiler/src/iree/compiler/Codegen/Dialect/PCF/ExternalInterfaces/test/bufferize.mlir
index 1ba7b65..cad2e83 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/PCF/ExternalInterfaces/test/bufferize.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/PCF/ExternalInterfaces/test/bufferize.mlir
@@ -2,7 +2,7 @@
util.func private @bufferize_generic(%d0: index, %d1: index, %d2: index, %d3: index) {
%0 = bufferization.alloc_tensor(%d0) : tensor<?xi32>
- %1 = bufferization.alloc_tensor(%d3) {memory_space = "foo"} : tensor<?xi32>
+ %1 = bufferization.alloc_tensor(%d3) <{memory_space = "foo"}> : tensor<?xi32>
%2:4 = pcf.generic scope(#pcf.test_scope)
execute(%ref = %0, %ref_1, %ref_2, %ref_3 = %1)[%id: index, %n: index]
: (!pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>)
@@ -93,7 +93,7 @@
util.func private @bufferize_loop(%d0: index, %d1: index, %d2: index, %d3: index, %n: index) {
%0 = bufferization.alloc_tensor(%d0) : tensor<?xi32>
- %1 = bufferization.alloc_tensor(%d3) {memory_space = "foo"} : tensor<?xi32>
+ %1 = bufferization.alloc_tensor(%d3) <{memory_space = "foo"}> : tensor<?xi32>
%2:4 = pcf.loop scope(#pcf.test_scope) count(%n)
execute(%ref = %0, %ref_1, %ref_2, %ref_3 = %1)[%id: index]
: (!pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>, !pcf.sref<?xi32, #pcf.test_scope>)
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index 04e6ab3..0c14a31 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -145,7 +145,7 @@
LLVM::LLVMFuncOp::getAttributeNames().end());
SmallVector<NamedAttribute> funcAttrs;
for (auto attr : stdFuncOp->getAttrs()) {
- if (attr.getName() == SymbolTable::getSymbolAttrName() ||
+ if (attr.getName() == stdFuncOp.getSymNameAttrName() ||
attr.getName() == stdFuncOp.getFunctionTypeAttrName()) {
continue;
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
index dd46633..b28e241 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
@@ -323,7 +323,7 @@
// Construct newFunc with all attributes except return type & symbol name.
SmallVector<NamedAttribute> funcAttrs;
for (auto attr : funcOp->getAttrs()) {
- if (attr.getName() == SymbolTable::getSymbolAttrName() ||
+ if (attr.getName() == funcOp.getSymNameAttrName() ||
attr.getName() == funcOp.getFunctionTypeAttrName()) {
continue;
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/annotate_kernel_for_translation.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/annotate_kernel_for_translation.mlir
index 5351d33..c275c62 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/annotate_kernel_for_translation.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/annotate_kernel_for_translation.mlir
@@ -176,9 +176,9 @@
}
// CHECK-LABEL: llvm.func @test_kern_arg
-// CHECK: denormal_fpenv = #llvm.denormal_fpenv<default_output_mode = ieee, default_input_mode = ieee, float_output_mode = preservesign, float_input_mode = preservesign>
+// CHECK-SAME: llvm_func_attrs = {check_attr
+// CHECK-SAME: denormal_fpenv = #llvm.denormal_fpenv<default_output_mode = ieee, default_input_mode = ieee, float_output_mode = preservesign, float_input_mode = preservesign>
// CHECK-NOT: iree_codegen.denormal_fp_math_f32
-// CHECK: llvm_func_attrs = {check_attr
// -----
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir
index c2d4300..ecb26f6 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir
@@ -219,7 +219,7 @@
// CHECK-LABEL: func.func @conv_nhwc(
// CHECK-SAME: #iree_codegen.translation_info<pipeline = #iree_gpu.pipeline<TileAndFuse> workgroup_size = [64, 1, 1] subgroup_size = 64>
-// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}} lowering_config = #iree_gpu.lowering_config
+// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}}lowering_config = #iree_gpu.lowering_config
// CHECK-SAME: reduction = [0, 0, 0, 0, 1, 3, 4]
// CHECK-SAME: thread = [1, 1, 1, 1, 0, 0, 0]
// CHECK-SAME: workgroup = [1, 1, 1, 64, 0, 0, 0]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx942.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx942.mlir
index 1c82d60..632a031 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx942.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx942.mlir
@@ -68,7 +68,7 @@
}
// CHECK-LABEL: func.func @conv_nhwc()
-// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}} lowering_config = #iree_gpu.lowering_config
+// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}}lowering_config = #iree_gpu.lowering_config
// CHECK-SAME: mma_kind = #iree_gpu.mma_layout<MFMA_F32_16x16x16_F16>
// CHECK-SAME: reduction = [0, 0, 0, 0, 1, 1, 32]
// CHECK-SAME{LITERAL}: subgroup_basis = [[1, 1, 2, 2, 1, 1, 1], [0, 1, 2, 3, 4, 5, 6]]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx950.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx950.mlir
index ec99fd7..ad44291 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx950.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute_gfx950.mlir
@@ -68,7 +68,7 @@
}
// CHECK-LABEL: func.func @conv_nhwc()
-// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}} lowering_config = #iree_gpu.lowering_config
+// CHECK: linalg.conv_2d_nhwc_hwcf {{.*}}lowering_config = #iree_gpu.lowering_config
// CHECK-SAME: mma_kind = #iree_gpu.mma_layout<MFMA_F32_16x16x32_F16>
// CHECK-SAME: reduction = [0, 0, 0, 0, 1, 1, 64]
// CHECK-SAME{LITERAL}: subgroup_basis = [[1, 1, 2, 2, 1, 1, 1], [0, 1, 2, 3, 4, 5, 6]]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
index 84e5d16..77fcac8 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
@@ -94,8 +94,8 @@
// CHECK-DAG: %[[C4096_i32:.+]] = llvm.mlir.constant(4096 : i32) : i32
// CHECK-DAG: %[[C0_i32:.+]] = llvm.mlir.constant(0 : i32) : i32
// CHECK-DAG: %[[C32_i64:.+]] = llvm.mlir.constant(32 : i64) : i64
-// CHECK-DAG: %[[C5185728_i64:.+]] = llvm.mlir.constant(5185728 : index) : i64
-// CHECK-DAG: %[[C4438911803328_i64:.+]] = llvm.mlir.constant(4438911803328 : index) : i64
+// CHECK-DAG: %[[C5185728_i64:.+]] = llvm.mlir.constant(5185728 : i64) : i64
+// CHECK-DAG: %[[C4438911803328_i64:.+]] = llvm.mlir.constant(4438911803328 : i64) : i64
// CHECK-DAG: %[[OFFSET_LO:.+]] = llvm.zext %[[ARG3]] : i32 to i64
// CHECK-DAG: %[[OFFSET_HI:.+]] = llvm.zext %[[ARG4]] : i32 to i64
// CHECK-DAG: %[[OFFSET_HI_SHL:.+]] = llvm.shl %[[OFFSET_HI]], %[[C32_i64]] : i64
@@ -194,7 +194,7 @@
// CHECK-SAME: %{{.*}}: !llvm.ptr {llvm.align = 16 : i32, llvm.noalias, llvm.nonnull, llvm.noundef})
// CHECK: %[[BYTES_PER_BIT:.+]] = llvm.mlir.constant(8 : i64) : i64
// CHECK: %[[BITS_PER_ELEM:.+]] = llvm.mlir.constant(32 : i64) : i64
-// CHECK: %[[BYTE_OFFSET:.+]] = llvm.mlir.constant(128 : index) : i64
+// CHECK: %[[BYTE_OFFSET:.+]] = llvm.mlir.constant(128 : i64) : i64
// CHECK: %[[OFFSET_BITS:.+]] = llvm.mul %[[BYTE_OFFSET]], %[[BYTES_PER_BIT]]
// CHECK: %[[OFFSET_ELEMS:.+]] = llvm.udiv %[[OFFSET_BITS]], %[[BITS_PER_ELEM]]
// CHECK: nvvm.read.ptx.sreg.tid.x
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
index c7ec552..3424195 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -1213,7 +1213,7 @@
void ExecutableOp::build(OpBuilder &builder, OperationState &state,
StringRef name) {
ensureTerminator(*state.addRegion(), builder, state.location);
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
}
@@ -1260,9 +1260,7 @@
ArrayAttr tiedOperands,
ArrayRef<NamedAttribute> attributes) {
StringRef executableOpSymName =
- exportOp->getParentOp()
- ->getAttrOfType<StringAttr>(SymbolTable::getSymbolAttrName())
- .getValue();
+ cast<ExecutableOp>(exportOp->getParentOp()).getName();
auto entryPoint =
SymbolRefAttr::get(builder.getContext(), executableOpSymName,
{SymbolRefAttr::get(exportOp)});
@@ -1389,9 +1387,9 @@
ArrayRef<NamedAttribute> attrs,
ArrayRef<DictionaryAttr> argAttrs,
ArrayRef<DictionaryAttr> resAttrs) {
- state.addAttribute(SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
- state.addAttribute(SymbolTable::getVisibilityAttrName(),
+ state.addAttribute(getSymVisibilityAttrName(state.name),
builder.getStringAttr("private"));
state.addAttribute("function_type", TypeAttr::get(type));
state.attributes.append(attrs.begin(), attrs.end());
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td
index 217aabd..ab33a43 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -426,6 +426,8 @@
def FLOW_ExecutableOp : FLOW_Op<"executable", [
IsolatedFromAbove,
SingleBlockImplicitTerminator<"IREE::Flow::ExecutableEndOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
Util_ObjectLike,
@@ -479,6 +481,8 @@
def FLOW_ExecutableExportOp : FLOW_Op<"executable.export", [
HasParent<"IREE::Flow::ExecutableOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
IsolatedFromAbove,
]> {
@@ -648,6 +652,8 @@
def FLOW_FuncOp : FLOW_Op<"func", [
CallableOpInterface,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
IsolatedFromAbove,
]> {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir
index ddbe031..1bd04e7 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir
@@ -2,7 +2,7 @@
// CHECK-DAG: vm.import private @hal.command_buffer.dispatch{{\(.+}}attributes {minimum_version = 7 : i32}
// CHECK-DAG: vm.import private @hal.command_buffer.dispatch.indirect{{\(.+}}attributes {minimum_version = 7 : i32}
-// CHECK-DAG: vm.import private @hal.executable.lookup.function{{.*}}attributes {minimum_version = 7 : i32, nosideeffects}
+// CHECK-DAG: vm.import private @hal.executable.lookup.function{{.*}}attributes {nosideeffects, minimum_version = 7 : i32}
// CHECK-LABEL: @command_buffer_create
// CHECK-SAME: (%[[DEVICE:.+]]: !vm.ref<!hal.device>, %[[AFFINITY:.+]]: i64)
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
index 99997ad..2407bee 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
@@ -1030,7 +1030,9 @@
rewriter.setInsertionPoint(blockOps.front());
auto fusedLoc = rewriter.getFusedLoc(blockLocs);
auto newBlockOp = ExecutableConstantBlockOp::create(
- rewriter, fusedLoc, rewriter.getFunctionType(inputTypes, resultTypes),
+ rewriter, fusedLoc, /*sym_name=*/StringAttr(),
+ /*sym_visibility=*/StringAttr(),
+ rewriter.getFunctionType(inputTypes, resultTypes),
rewriter.getArrayAttr(resultKeys), /*arg_attrs=*/ArrayAttr(),
/*res_attrs=*/ArrayAttr());
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
index 91aec1e..c5f3d97 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
@@ -1683,7 +1683,7 @@
void ExecutableOp::build(OpBuilder &builder, OperationState &state,
StringRef name) {
ensureTerminator(*state.addRegion(), builder, state.location);
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
}
@@ -1961,7 +1961,7 @@
StringRef symName,
IREE::HAL::ExecutableTargetAttr target) {
ensureTerminator(*state.addRegion(), builder, state.location);
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(symName));
state.addAttribute("target", target);
}
@@ -2228,7 +2228,7 @@
void ExecutableBinaryOp::build(OpBuilder &builder, OperationState &state,
StringRef symName, StringRef format,
std::vector<uint8_t> data) {
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(symName));
state.addAttribute("format", builder.getStringAttr(format));
state.addAttribute("data",
@@ -2241,7 +2241,7 @@
void ExecutableBinaryOp::build(OpBuilder &builder, OperationState &state,
StringRef symName, StringAttr format,
DenseIntElementsAttr data) {
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(symName));
state.addAttribute("format", format);
state.addAttribute("data", data);
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.td b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.td
index 06a04f1..3460f27 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.td
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -2615,6 +2615,8 @@
def HAL_ExecutableSourceOp : HAL_Op<"executable.source", [
IsolatedFromAbove,
SingleBlockImplicitTerminator<"IREE::HAL::ExecutableSourceEndOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
Util_ObjectLike,
@@ -2684,6 +2686,8 @@
def HAL_ExecutableOp : HAL_Op<"executable", [
IsolatedFromAbove,
SingleBlockImplicitTerminator<"IREE::HAL::ExecutableEndOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
Util_ObjectLike,
@@ -2740,6 +2744,8 @@
}
def HAL_ExecutableExportOp : HAL_Op<"executable.export", [
+ SymbolName,
+ SymbolVisibility,
Symbol,
DeclareOpInterfaceMethods<SymbolUserOpInterface>,
ParentOneOf<[
@@ -2868,6 +2874,8 @@
IsolatedFromAbove,
HasParent<"IREE::HAL::ExecutableOp">,
SingleBlockImplicitTerminator<"IREE::HAL::ExecutableVariantEndOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
]> {
@@ -2963,6 +2971,8 @@
def HAL_ExecutableConditionOp : HAL_Op<"executable.condition", [
IsolatedFromAbove,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
CallableOpInterface,
]> {
@@ -2976,6 +2986,8 @@
}];
let arguments = (ins
+ OptionalAttr<SymbolNameAttr>:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
TypeAttrOf<FunctionType>:$function_type,
OptionalAttr<DictArrayAttr>:$arg_attrs,
OptionalAttr<DictArrayAttr>:$res_attrs
@@ -3021,6 +3033,8 @@
]>,
IsolatedFromAbove,
CallableOpInterface,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
]> {
let summary = [{Executable constant block initializer.}];
@@ -3066,6 +3080,8 @@
}];
let arguments = (ins
+ OptionalAttr<SymbolNameAttr>:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
TypeAttrOf<FunctionType>:$function_type,
ArrayAttr:$keys,
OptionalAttr<DictArrayAttr>:$arg_attrs,
@@ -3134,6 +3150,8 @@
def HAL_ExecutableBinaryOp : HAL_Op<"executable.binary", [
HasParent<"IREE::HAL::ExecutableOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
]> {
let summary = [{Compiled executable binary data.}];
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.cpp
index dd6b745..e00509c 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.cpp
@@ -3121,9 +3121,7 @@
ArrayRef<int64_t> tiedOperands,
AffinityAttr affinityAttr) {
StringRef executableOpSymName =
- exportOp->getParentOp()
- ->getAttrOfType<StringAttr>(SymbolTable::getSymbolAttrName())
- .getValue();
+ cast<ExecutableOp>(exportOp->getParentOp()).getName();
auto entryPoint =
SymbolRefAttr::get(builder.getContext(), executableOpSymName,
{SymbolRefAttr::get(exportOp)});
@@ -3299,9 +3297,9 @@
ArrayAttr tiedOperands,
ArrayRef<DictionaryAttr> argAttrs,
ArrayRef<DictionaryAttr> resAttrs) {
- state.addAttribute(SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
- state.addAttribute(SymbolTable::getVisibilityAttrName(),
+ state.addAttribute(getSymVisibilityAttrName(state.name),
builder.getStringAttr("private"));
state.addAttribute("function_type", TypeAttr::get(type));
if (tiedOperands) {
@@ -4339,9 +4337,9 @@
void CmdFuncOp::build(OpBuilder &builder, OperationState &state, StringRef name,
FunctionType type, ArrayRef<DictionaryAttr> argAttrs,
ArrayRef<DictionaryAttr> resAttrs) {
- state.addAttribute(SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
- state.addAttribute(SymbolTable::getVisibilityAttrName(),
+ state.addAttribute(getSymVisibilityAttrName(state.name),
builder.getStringAttr("private"));
state.addAttribute("function_type", TypeAttr::get(type));
state.addRegion();
@@ -5126,7 +5124,7 @@
void ExecutableOp::build(OpBuilder &builder, OperationState &state,
StringRef sym_name) {
ensureTerminator(*state.addRegion(), builder, state.location);
- state.addAttribute(mlir::SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(sym_name));
}
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.td b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.td
index 4d1fb78..68dff9e 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.td
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOps.td
@@ -2837,6 +2837,8 @@
def Stream_AsyncFuncOp : Stream_Op<"async.func", [
CallableOpInterface,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
IsolatedFromAbove,
Stream_AsyncPhaseOp,
@@ -3931,6 +3933,8 @@
def Stream_CmdFuncOp : Stream_Op<"cmd.func", [
CallableOpInterface,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
IsolatedFromAbove,
Stream_CmdPhaseOp,
@@ -5018,6 +5022,8 @@
def Stream_ExecutableOp : Stream_Op<"executable", [
IsolatedFromAbove,
SingleBlockImplicitTerminator<"IREE::Stream::ExecutableEndOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
Util_ObjectLike,
@@ -5070,6 +5076,8 @@
def Stream_ExecutableExportOp : Stream_Op<"executable.export", [
HasParent<"IREE::Stream::ExecutableOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
IsolatedFromAbove,
]> {
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/Utils.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/Utils.cpp
index 8e5b4f8..086c37a 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/Utils.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/Utils.cpp
@@ -215,8 +215,7 @@
assert(dispatchSiteToExecutableOp.count(info));
auto executableOp = dispatchSiteToExecutableOp[info];
- auto newSym = SymbolRefAttr::get(executableOp->getAttrOfType<StringAttr>(
- SymbolTable::getSymbolAttrName()),
+ auto newSym = SymbolRefAttr::get(executableOp.getNameAttr(),
entryPoint.getNestedReferences());
newEntryPoints.push_back(newSym);
});
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.cpp b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.cpp
index 536d253..3e2c343 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.cpp
@@ -1826,9 +1826,9 @@
ArrayRef<NamedAttribute> attrs,
ArrayRef<DictionaryAttr> argAttrs,
ArrayRef<DictionaryAttr> resAttrs) {
- state.addAttribute(SymbolTable::getSymbolAttrName(),
+ state.addAttribute(getSymNameAttrName(state.name),
builder.getStringAttr(name));
- state.addAttribute(SymbolTable::getVisibilityAttrName(),
+ state.addAttribute(getSymVisibilityAttrName(state.name),
builder.getStringAttr("public"));
state.addAttribute("function_type", TypeAttr::get(type));
state.attributes.append(attrs.begin(), attrs.end());
@@ -1891,12 +1891,12 @@
return failure();
}
if (symVisibilityAttr) {
- result.addAttribute(SymbolTable::getVisibilityAttrName(),
+ result.addAttribute(getSymVisibilityAttrName(result.name),
symVisibilityAttr);
}
StringAttr nameAttr;
- if (parser.parseSymbolName(nameAttr, SymbolTable::getSymbolAttrName(),
+ if (parser.parseSymbolName(nameAttr, getSymNameAttrName(result.name),
result.attributes)) {
return failure();
}
@@ -1932,8 +1932,8 @@
return failure();
}
for (StringRef disallowed : {
- SymbolTable::getVisibilityAttrName(),
- SymbolTable::getSymbolAttrName(),
+ getSymVisibilityAttrName(result.name).getValue(),
+ getSymNameAttrName(result.name).getValue(),
StringRef("function_type"),
}) {
if (parsedAttributes.get(disallowed)) {
@@ -2244,7 +2244,7 @@
bool isMutable, Type type,
std::optional<TypedAttr> initialValue,
ArrayRef<NamedAttribute> attrs) {
- result.addAttribute(SymbolTable::getSymbolAttrName(),
+ result.addAttribute(getSymNameAttrName(result.name),
builder.getStringAttr(name));
if (isMutable) {
result.addAttribute("is_mutable", builder.getUnitAttr());
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
index f3a748f..0b735d1 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
@@ -658,6 +658,8 @@
def Util_InitializerOp : Util_Op<"initializer", [
IsolatedFromAbove,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
CallableOpInterface,
Util_InitializerOpInterface,
@@ -704,6 +706,8 @@
}];
let arguments = (ins
+ OptionalAttr<SymbolNameAttr>:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
TypeAttrOf<FunctionType>:$function_type,
OptionalAttr<DictArrayAttr>:$arg_attrs,
OptionalAttr<DictArrayAttr>:$res_attrs
@@ -757,6 +761,8 @@
def Util_FuncOp : Util_Op<"func", [
AffineScope,
AutomaticAllocationScope,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
IsolatedFromAbove,
OpAsmOpInterface,
@@ -1072,6 +1078,8 @@
let opDocGroup = OpGroupGlobalOps in {
def Util_GlobalOp : Util_Op<"global", [
+ SymbolName,
+ SymbolVisibility,
Symbol,
Util_GlobalOpInterface,
]> {
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/const_ops.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/const_ops.mlir
index bf07a67..09b4de4 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/const_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/const_ops.mlir
@@ -32,8 +32,8 @@
vm.func @const_ref_zero() {
// CHECK: %[[REF:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[REFPTR:.+]] = address_of %[[REF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
- // CHECK-NEXT: %[[SIZE:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: call_opaque "memset"(%[[REFPTR]], %[[SIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
+ // CHECK-NEXT: %[[SIZE:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: call_opaque "memset"(%[[REFPTR]], %[[SIZE]]) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
// CHECK-NEXT: call_opaque "iree_vm_ref_release"(%[[REFPTR]]) : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> ()
%null = vm.const.ref.zero : !vm.ref<?>
vm.return
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
index 7799af1..941aa21 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
@@ -356,7 +356,7 @@
// In case of fail, return status message.
// CHECK-NEXT: ^[[FAIL]]:
- // CHECK-NEXT: %[[MSG:.+]] = call_opaque "iree_make_cstring_view"() <{args = [#emitc.opaque<"\22message\22">]}>
+ // CHECK-NEXT: %[[MSG:.+]] = call_opaque "iree_make_cstring_view"() <args = [#emitc.opaque<"\22message\22">]>
// CHECK-SAME: : () -> !emitc.opaque<"iree_string_view_t">
// CHECK-NEXT: %[[MSG_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<!emitc.opaque<"iree_string_view_t">>
// CHECK-NEXT: assign %[[MSG]] : !emitc.opaque<"iree_string_view_t"> to %[[MSG_LVAL]] : <!emitc.opaque<"iree_string_view_t">>
@@ -366,7 +366,7 @@
// CHECK-NEXT: %[[MSGDATA_LVAL:.+]] = "emitc.member"(%[[MSG_LVAL]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_string_view_t">>) -> !emitc.lvalue<!emitc.ptr<!emitc.opaque<"const char">>>
// CHECK-NEXT: %[[MSGDATA:.+]] = load %[[MSGDATA_LVAL]] : <!emitc.ptr<!emitc.opaque<"const char">>>
// CHECK-NEXT: %[[FAILSTATUS:.+]] = call_opaque "iree_status_allocate_f"(%[[MSGSIZEINT]], %[[MSGDATA]])
- // CHECK-SAME: <{args = [#emitc.opaque<"IREE_STATUS_FAILED_PRECONDITION">, #emitc.opaque<"\22<vm>\22">, 0 : i32, #emitc.opaque<"\22%.*s\22">, 0 : index, 1 : index]}>
+ // CHECK-SAME: <args = [#emitc.opaque<"IREE_STATUS_FAILED_PRECONDITION">, #emitc.opaque<"\22<vm>\22">, 0 : i32, #emitc.opaque<"\22%.*s\22">, 0 : index, 1 : index]>
// CHECK-SAME: : (!emitc.opaque<"int">, !emitc.ptr<!emitc.opaque<"const char">>) -> !emitc.opaque<"iree_status_t">
// CHECK-NEXT: return %[[FAILSTATUS]] : !emitc.opaque<"iree_status_t">
vm.fail %arg0, "message"
@@ -408,7 +408,7 @@
// CHECK-NEXT: assign %[[ARGSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[ARGSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[ARGSDATA:.+]] = "emitc.member"(%[[ARGBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[ARGBYTESPANDATA]] : !emitc.ptr<ui8> to %[[ARGSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGALLOCASIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGALLOCASIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Allocate space for the result.
// CHECK-NEXT: %[[RESBYTESPAN_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
@@ -420,7 +420,7 @@
// CHECK-NEXT: assign %[[RESULTSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[RESSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[RESSDATA:.+]] = "emitc.member"(%[[RESBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[RESBYTESPANDATA]] : !emitc.ptr<ui8> to %[[RESSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESALLOCASIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESALLOCASIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Check that we don't pack anything into the argument struct.
// CHECK-NOT: "emitc.member"(%{{.+}}) <{member = "arguments"}>
@@ -460,22 +460,22 @@
// Calculate the size of the arguments.
// CHECK-NEXT: %[[ARGSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE01:.+]] = add %[[ARGSIZE0]], %[[ARGSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE2:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE012:.+]] = add %[[ARGSIZE01]], %[[ARGSIZE2]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE3:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE3:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE0123:.+]] = add %[[ARGSIZE012]], %[[ARGSIZE3]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE4:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE4:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE:.+]] = add %[[ARGSIZE0123]], %[[ARGSIZE4]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
// Calculate the size of the result.
// CHECK-NEXT: %[[RESULTSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[RESULTSIZE:.+]] = add %[[RESULTSIZE0]], %[[RESULTSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
@@ -498,7 +498,7 @@
// CHECK-NEXT: assign %[[ARGSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[ARGSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[ARGSDATA:.+]] = "emitc.member"(%[[ARGBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[ARGBYTESPANDATA]] : !emitc.ptr<ui8> to %[[ARGSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Allocate space for the result.
// CHECK-NEXT: %[[RESBYTESPAN_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
@@ -508,7 +508,7 @@
// CHECK-NEXT: assign %[[RESULTSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[RESSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[RESSDATA:.+]] = "emitc.member"(%[[RESBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[RESBYTESPANDATA]] : !emitc.ptr<ui8> to %[[RESSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Pack the arguments into the struct.
// Here we also create pointers for non-pointer types.
@@ -517,31 +517,31 @@
// CHECK-NEXT: %[[ARGSPTR:.+]] = load %[[ARGSPTR_LVAL]] : <!emitc.ptr<ui8>>
// CHECK-NEXT: %[[A1_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg2 : i32 to %[[A1_LVAL]] : <i32>
- // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A1PTR:.+]] = address_of %[[A1_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[ARGSPTR]], %[[A1PTR]], %[[A1SIZE]])
- // CHECK-NEXT: %[[ARGHOSTSIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[ARGHOSTSIZE2:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A1ADDR:.+]] = add %[[ARGSPTR]], %[[ARGHOSTSIZE2]]
// CHECK-SAME: : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
// CHECK-NEXT: %[[A2_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg3 : i32 to %[[A2_LVAL]] : <i32>
- // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A2PTR:.+]] = address_of %[[A2_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[A1ADDR]], %[[A2PTR]], %[[A1SIZE]])
- // CHECK-NEXT: %[[A1SIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A1SIZE2:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A2ADDR:.+]] = add %[[A1ADDR]], %[[A1SIZE2]]
// CHECK-SAME: : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
// CHECK-NEXT: %[[A3_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg4 : i32 to %[[A3_LVAL]] : <i32>
- // CHECK-NEXT: %[[A2SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A2SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A3PTR:.+]] = address_of %[[A3_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[A2ADDR]], %[[A3PTR]], %[[A2SIZE]])
- // CHECK-NEXT: %[[A2SIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A2SIZE2:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A3ADDR:.+]] = add %[[A2ADDR]], %[[A2SIZE2]]
// CHECK-SAME: : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
// CHECK-NEXT: %[[A4_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg5 : i32 to %[[A4_LVAL]] : <i32>
- // CHECK-NEXT: %[[A3SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A3SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A4PTR:.+]] = address_of %[[A4_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[A3ADDR]], %[[A4PTR]], %[[A3SIZE:.+]])
@@ -557,7 +557,7 @@
// CHECK: %[[RES_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
// CHECK-NEXT: %[[RESPTR_MEMBER:.+]] = "emitc.member"(%[[RES_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RESPTR:.+]] = load %[[RESPTR_MEMBER]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: %[[RESHOSTSIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[RESHOSTSIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: call_opaque "memcpy"(%arg6, %[[RESPTR]], %[[RESHOSTSIZE]])
// Return ok status.
@@ -582,16 +582,16 @@
// Calculate the size of the arguments.
// CHECK-NEXT: %[[ARGSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE01:.+]] = add %[[ARGSIZE0]], %[[ARGSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[ARGSIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[ARGSIZE2:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[ARGSIZE:.+]] = add %[[ARGSIZE01]], %[[ARGSIZE2]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
// Calculate the size of the result.
// CHECK-NEXT: %[[RESULTSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <{args = [i32]}>
+ // CHECK-NEXT: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <args = [i32]>
// CHECK-NEXT: %[[RESULTSIZE:.+]] = add %[[RESULTSIZE0]], %[[RESULTSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
@@ -614,7 +614,7 @@
// CHECK-NEXT: assign %[[ARGSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[ARGSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[ARGSDATA:.+]] = "emitc.member"(%[[ARGBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[ARGBYTESPANDATA]] : !emitc.ptr<ui8> to %[[ARGSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Allocate space for the result.
// CHECK-NEXT: %[[RESBYTESPAN_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
@@ -624,7 +624,7 @@
// CHECK-NEXT: assign %[[RESULTSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[RESSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[RESSDATA:.+]] = "emitc.member"(%[[RESBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[RESBYTESPANDATA]] : !emitc.ptr<ui8> to %[[RESSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Pack the arguments into the struct.
// Here we also create pointers for non-pointer types.
@@ -633,15 +633,15 @@
// CHECK-NEXT: %[[ARGSPTR:.+]] = load %[[ARGSPTR_LVAL]] : <!emitc.ptr<ui8>>
// CHECK-NEXT: %[[A1_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg2 : i32 to %[[A1_LVAL]] : <i32>
- // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A1PTR:.+]] = address_of %[[A1_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[ARGSPTR]], %[[A1PTR]], %[[A1SIZE]])
- // CHECK-NEXT: %[[ARGHOSTSIZE2:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[ARGHOSTSIZE2:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A1ADDR:.+]] = add %[[ARGSPTR]], %[[ARGHOSTSIZE2]]
// CHECK-SAME: : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
// CHECK-NEXT: %[[A2_LVAL:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<i32>
// CHECK-NEXT: assign %arg3 : i32 to %[[A2_LVAL]] : <i32>
- // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[A1SIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: %[[A2PTR:.+]] = address_of %[[A2_LVAL]] : !emitc.lvalue<i32>
// CHECK-NEXT: call_opaque "memcpy"(%[[A1ADDR]], %[[A2PTR]], %[[A1SIZE]])
@@ -657,7 +657,7 @@
// CHECK: %[[RES_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
// CHECK-NEXT: %[[RESPTR_MEMBER:.+]] = "emitc.member"(%[[RES_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RESPTR:.+]] = load %[[RESPTR_MEMBER]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: %[[RESHOSTSIZE:.+]] = call_opaque "sizeof"() <{args = [i32]}> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: %[[RESHOSTSIZE:.+]] = call_opaque "sizeof"() <args = [i32]> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK-NEXT: call_opaque "memcpy"(%arg4, %[[RESPTR]], %[[RESHOSTSIZE]])
// Return ok status.
@@ -683,7 +683,7 @@
// Calculate the size of the arguments (with alignment padding for refs).
// CHECK: %[[ARGSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK: call_opaque "iree_host_align"
- // CHECK: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}>
+ // CHECK: %[[ARGSIZE1:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]>
// CHECK: %[[ARGSIZE_PRE:.+]] = add %{{.+}}, %[[ARGSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
// CHECK: %[[ARGSIZE:.+]] = call_opaque "iree_host_align"(%[[ARGSIZE_PRE]]
@@ -691,7 +691,7 @@
// Calculate the size of the result (with alignment padding for refs).
// CHECK: %[[RESULTSIZE0:.+]] = "emitc.constant"() <{value = #emitc.opaque<"0">}> : () -> !emitc.opaque<"iree_host_size_t">
// CHECK: call_opaque "iree_host_align"
- // CHECK: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}>
+ // CHECK: %[[RESULTSIZE1:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]>
// CHECK: %[[RESULTSIZE_PRE:.+]] = add %{{.+}}, %[[RESULTSIZE1]]
// CHECK-SAME: : (!emitc.opaque<"iree_host_size_t">, !emitc.opaque<"iree_host_size_t">) -> !emitc.opaque<"iree_host_size_t">
// CHECK: %[[RESULTSIZE:.+]] = call_opaque "iree_host_align"(%[[RESULTSIZE_PRE]]
@@ -715,7 +715,7 @@
// CHECK-NEXT: assign %[[ARGSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[ARGSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[ARGSDATA:.+]] = "emitc.member"(%[[ARGBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[ARGBYTESPANDATA]] : !emitc.ptr<ui8> to %[[ARGSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Allocate space for the result.
// CHECK-NEXT: %[[RESBYTESPAN_MEMBER:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "results"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
@@ -725,7 +725,7 @@
// CHECK-NEXT: assign %[[RESULTSIZE]] : !emitc.opaque<"iree_host_size_t"> to %[[RESSDATALENGTH]] : <!emitc.opaque<"iree_host_size_t">>
// CHECK-NEXT: %[[RESSDATA:.+]] = "emitc.member"(%[[RESBYTESPAN_MEMBER]]) <{member = "data"}> : (!emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: assign %[[RESBYTESPANDATA]] : !emitc.ptr<ui8> to %[[RESSDATA]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}>
+ // CHECK-NEXT: call_opaque "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]>
// Pack the argument into the struct (with pointer alignment).
// CHECK: %[[ARGS:.+]] = "emitc.member"(%[[ARGSTRUCT]]) <{member = "arguments"}> : (!emitc.lvalue<!emitc.opaque<"iree_vm_function_call_t">>) -> !emitc.lvalue<!emitc.opaque<"iree_byte_span_t">>
@@ -835,8 +835,8 @@
// Create duplicate ref for
// CHECK-NEXT: %[[REF:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[REFPTR:.+]] = address_of %[[REF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
- // CHECK-NEXT: %[[REFSIZE:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: call_opaque "memset"(%[[REFPTR]], %[[REFSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
+ // CHECK-NEXT: %[[REFSIZE:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: call_opaque "memset"(%[[REFPTR]], %[[REFSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
// CHECK-NEXT: call_opaque "iree_vm_ref_move"(%arg4, %[[REFPTR]]) : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> ()
// CHECK-NEXT: call_opaque "iree_vm_ref_move"(%[[REFPTR]], %arg6) : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> ()
@@ -935,15 +935,15 @@
// Allocate ref variable for argument.
// CHECK-NEXT: %[[ARGREF:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[ARGREFPTR:.+]] = address_of %[[ARGREF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
- // CHECK-NEXT: %[[ARGREFSIZE:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: call_opaque "memset"(%[[ARGREFPTR]], %[[ARGREFSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
+ // CHECK-NEXT: %[[ARGREFSIZE:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: call_opaque "memset"(%[[ARGREFPTR]], %[[ARGREFSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
// CHECK-NEXT: call_opaque "iree_vm_ref_assign"(%arg3, %[[ARGREFPTR]]) : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> ()
// Allocate ref variable for result.
// CHECK-NEXT: %[[RESREF:.+]] = "emitc.variable"() <{value = #emitc.opaque<"">}> : () -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[RESREFPTR:.+]] = address_of %[[RESREF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
- // CHECK-NEXT: %[[RESREFSIZE:.+]] = call_opaque "sizeof"() <{args = [!emitc.opaque<"iree_vm_ref_t">]}> : () -> !emitc.opaque<"iree_host_size_t">
- // CHECK-NEXT: call_opaque "memset"(%[[RESREFPTR]], %[[RESREFSIZE]]) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
+ // CHECK-NEXT: %[[RESREFSIZE:.+]] = call_opaque "sizeof"() <args = [!emitc.opaque<"iree_vm_ref_t">]> : () -> !emitc.opaque<"iree_host_size_t">
+ // CHECK-NEXT: call_opaque "memset"(%[[RESREFPTR]], %[[RESREFSIZE]]) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> ()
// Call the import.
// CHECK-NEXT: %{{.+}} = call @my_module_call_[[YIELDABLEFN]](%arg0, %[[IMPORT]], %[[ARGREFPTR]], %[[RESREFPTR]])
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops.mlir
index 9b643ff..e5258e6 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops.mlir
@@ -9,7 +9,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_i32"(%[[RWDATA]]) <{args = [0 : index, 0 : ui32]}> : (!emitc.ptr<ui8>) -> i32
+ // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_i32"(%[[RWDATA]]) <args = [0 : index, 0 : ui32]> : (!emitc.ptr<ui8>) -> i32
%0 = vm.global.load.i32 @c42 : i32
vm.return %0 : i32
}
@@ -26,7 +26,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "vm_global_store_i32"(%[[RWDATA]], %arg3) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<ui8>, i32) -> ()
+ // CHECK-NEXT: call_opaque "vm_global_store_i32"(%[[RWDATA]], %arg3) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<ui8>, i32) -> ()
vm.global.store.i32 %arg0, @c107_mut : i32
vm.return
}
@@ -47,7 +47,7 @@
// CHECK-NEXT: %[[REF:.+]] = subscript %[[REFS]][%[[REF_INDEX]]] : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[REF_0:.+]] = address_of %[[REF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK: %[[C:.+]] = call_opaque "iree_vm_type_def_as_ref"(%{{.+}}) : (!emitc.opaque<"iree_vm_type_def_t">) -> !emitc.opaque<"iree_vm_ref_type_t">
- // CHECK: %{{.+}} = call_opaque "iree_vm_ref_retain_or_move_checked"(%[[REF_0]], %[[C]], %arg3) <{args = [false, 0 : index, 1 : index, 2 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_vm_ref_type_t">, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> !emitc.opaque<"iree_status_t">
+ // CHECK: %{{.+}} = call_opaque "iree_vm_ref_retain_or_move_checked"(%[[REF_0]], %[[C]], %arg3) <args = [false, 0 : index, 1 : index, 2 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_vm_ref_type_t">, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> !emitc.opaque<"iree_status_t">
%0 = vm.global.load.ref @g0 : !vm.buffer
vm.return %0 : !vm.buffer
}
@@ -68,7 +68,7 @@
// CHECK-NEXT: %[[REF:.+]] = subscript %[[REFS]][%[[REF_INDEX]]] : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %[[REF_0:.+]] = address_of %[[REF]] : !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK: %[[C:.+]] = call_opaque "iree_vm_type_def_as_ref"(%{{.+}}) : (!emitc.opaque<"iree_vm_type_def_t">) -> !emitc.opaque<"iree_vm_ref_type_t">
- // CHECK: %{{.+}} = call_opaque "iree_vm_ref_retain_or_move_checked"(%arg3, %[[C]], %[[REF_0]]) <{args = [false, 0 : index, 1 : index, 2 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_vm_ref_type_t">, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> !emitc.opaque<"iree_status_t">
+ // CHECK: %{{.+}} = call_opaque "iree_vm_ref_retain_or_move_checked"(%arg3, %[[C]], %[[REF_0]]) <args = [false, 0 : index, 1 : index, 2 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_vm_ref_type_t">, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>) -> !emitc.opaque<"iree_status_t">
vm.global.store.ref %arg0, @g0_mut : !vm.buffer
vm.return
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_f32.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_f32.mlir
index 9d6bf49..c264f6c 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_f32.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_f32.mlir
@@ -9,7 +9,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_f32"(%[[RWDATA]]) <{args = [0 : index, 0 : ui32]}> : (!emitc.ptr<ui8>) -> f32
+ // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_f32"(%[[RWDATA]]) <args = [0 : index, 0 : ui32]> : (!emitc.ptr<ui8>) -> f32
%0 = vm.global.load.f32 @c42 : f32
vm.return %0 : f32
}
@@ -26,7 +26,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "vm_global_store_f32"(%[[RWDATA]], %arg3) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<ui8>, f32) -> ()
+ // CHECK-NEXT: call_opaque "vm_global_store_f32"(%[[RWDATA]], %arg3) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<ui8>, f32) -> ()
vm.global.store.f32 %arg0, @c107_mut : f32
vm.return
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_i64.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_i64.mlir
index a9d4425..7fa9bac 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_i64.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/global_ops_i64.mlir
@@ -9,7 +9,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_i64"(%[[RWDATA]]) <{args = [0 : index, 0 : ui32]}> : (!emitc.ptr<ui8>) -> i64
+ // CHECK-NEXT: %[[RES:.+]] = call_opaque "vm_global_load_i64"(%[[RWDATA]]) <args = [0 : index, 0 : ui32]> : (!emitc.ptr<ui8>) -> i64
%0 = vm.global.load.i64 @c42 : i64
vm.return %0 : i64
}
@@ -26,7 +26,7 @@
// CHECK-NEXT: assign %arg2 : !emitc.ptr<!emitc.opaque<"struct my_module_state_t">> to %[[STATE_LVAL]] : <!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>
// CHECK-NEXT: %[[RWDATA_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "rwdata"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.ptr<ui8>>
// CHECK-NEXT: %[[RWDATA:.+]] = load %[[RWDATA_LVAL]] : <!emitc.ptr<ui8>>
- // CHECK-NEXT: call_opaque "vm_global_store_i64"(%[[RWDATA]], %arg3) <{args = [0 : index, 0 : ui32, 1 : index]}> : (!emitc.ptr<ui8>, i64) -> ()
+ // CHECK-NEXT: call_opaque "vm_global_store_i64"(%[[RWDATA]], %arg3) <args = [0 : index, 0 : ui32, 1 : index]> : (!emitc.ptr<ui8>, i64) -> ()
vm.global.store.i64 %arg0, @c107_mut : i64
vm.return
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops.mlir
index 711d4e0..09b137c 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops.mlir
@@ -10,7 +10,7 @@
// CHECK-DAG: %[[ALLOCATOR_LVAL:.+]] = "emitc.member_of_ptr"(%[[STATE_LVAL]]) <{member = "allocator"}> : (!emitc.lvalue<!emitc.ptr<!emitc.opaque<"struct my_module_state_t">>>) -> !emitc.lvalue<!emitc.opaque<"iree_allocator_t">>
// CHECK-DAG: %[[ALLOCATOR:.+]] = load %[[ALLOCATOR_LVAL]] : <!emitc.opaque<"iree_allocator_t">>
- // CHECK: %[[TYPE_DEF:.+]] = call_opaque "iree_vm_make_value_type_def"() <{args = [#emitc.opaque<"IREE_VM_VALUE_TYPE_I32">]}> : () -> !emitc.opaque<"iree_vm_type_def_t">
+ // CHECK: %[[TYPE_DEF:.+]] = call_opaque "iree_vm_make_value_type_def"() <args = [#emitc.opaque<"IREE_VM_VALUE_TYPE_I32">]> : () -> !emitc.opaque<"iree_vm_type_def_t">
// CHECK-NEXT: %[[STATUS:.+]] = call_opaque "iree_vm_list_create"(%[[TYPE_DEF]], %arg3, %[[ALLOCATOR]], %[[LIST_PTR]]) : (!emitc.opaque<"iree_vm_type_def_t">, i32, !emitc.opaque<"iree_allocator_t">, !emitc.ptr<!emitc.ptr<!emitc.opaque<"iree_vm_list_t">>>) -> !emitc.opaque<"iree_status_t">
// CHECK: %[[LIST_TYPE_ID:.+]] = call_opaque "iree_vm_list_type"() : () -> !emitc.opaque<"iree_vm_ref_type_t">
@@ -78,7 +78,7 @@
// CHECK-NEXT: %3 = subscript %arg3[%2] : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %4 = load %3 : <!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %5 = call_opaque "iree_vm_list_deref"(%4) : (!emitc.opaque<"iree_vm_ref_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_list_t">>
- // CHECK: %{{.+}} = call_opaque "iree_vm_list_get_value_as"(%5, %arg4, %1) <{args = [0 : index, 1 : index, #emitc.opaque<"IREE_VM_VALUE_TYPE_I32">, 2 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_list_t">>, i32, !emitc.ptr<!emitc.opaque<"iree_vm_value_t">>) -> !emitc.opaque<"iree_status_t">
+ // CHECK: %{{.+}} = call_opaque "iree_vm_list_get_value_as"(%5, %arg4, %1) <args = [0 : index, 1 : index, #emitc.opaque<"IREE_VM_VALUE_TYPE_I32">, 2 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_list_t">>, i32, !emitc.ptr<!emitc.opaque<"iree_vm_value_t">>) -> !emitc.opaque<"iree_status_t">
%0 = vm.list.get.i32 %arg0, %arg1 : (!vm.list<i32>, i32) -> i32
vm.return %0 : i32
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops_i64.mlir b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops_i64.mlir
index 3728e8e..83993ba 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops_i64.mlir
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/list_ops_i64.mlir
@@ -9,7 +9,7 @@
// CHECK-NEXT: %3 = subscript %arg3[%2] : (!emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, !emitc.opaque<"iree_host_size_t">) -> !emitc.lvalue<!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %4 = load %3 : <!emitc.opaque<"iree_vm_ref_t">>
// CHECK-NEXT: %5 = call_opaque "iree_vm_list_deref"(%4) : (!emitc.opaque<"iree_vm_ref_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_list_t">>
- // CHECK: %{{.+}} = call_opaque "iree_vm_list_get_value_as"(%5, %arg4, %1) <{args = [0 : index, 1 : index, #emitc.opaque<"IREE_VM_VALUE_TYPE_I64">, 2 : index]}> : (!emitc.ptr<!emitc.opaque<"iree_vm_list_t">>, i32, !emitc.ptr<!emitc.opaque<"iree_vm_value_t">>) -> !emitc.opaque<"iree_status_t">
+ // CHECK: %{{.+}} = call_opaque "iree_vm_list_get_value_as"(%5, %arg4, %1) <args = [0 : index, 1 : index, #emitc.opaque<"IREE_VM_VALUE_TYPE_I64">, 2 : index]> : (!emitc.ptr<!emitc.opaque<"iree_vm_list_t">>, i32, !emitc.ptr<!emitc.opaque<"iree_vm_value_t">>) -> !emitc.opaque<"iree_status_t">
%0 = vm.list.get.i64 %arg0, %arg1 : (!vm.list<i64>, i32) -> i64
vm.return %0 : i64
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
index 199d9bc..4538ed7 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
@@ -104,7 +104,7 @@
StringRef name) {
ensureTerminator(*result.addRegion(), builder, result.location);
result.attributes.push_back(builder.getNamedAttr(
- mlir::SymbolTable::getSymbolAttrName(), builder.getStringAttr(name)));
+ getSymNameAttrName(result.name), builder.getStringAttr(name)));
}
LogicalResult ModuleOp::verify() {
@@ -149,7 +149,7 @@
FunctionType type, ArrayRef<NamedAttribute> attrs,
ArrayRef<DictionaryAttr> argAttrs) {
result.addRegion();
- result.addAttribute(SymbolTable::getSymbolAttrName(),
+ result.addAttribute(getSymNameAttrName(result.name),
builder.getStringAttr(name));
result.addAttribute("function_type", TypeAttr::get(type));
result.attributes.append(attrs.begin(), attrs.end());
@@ -284,8 +284,7 @@
result.addAttribute("is_optional", builder.getUnitAttr());
}
StringAttr nameAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
+ if (failed(parser.parseSymbolName(nameAttr, getSymNameAttrName(result.name),
result.attributes)) ||
failed(parser.parseLParen())) {
return parser.emitError(parser.getNameLoc()) << "invalid import name";
@@ -391,7 +390,7 @@
void ImportOp::build(OpBuilder &builder, OperationState &result, StringRef name,
FunctionType type, ArrayRef<NamedAttribute> attrs,
ArrayRef<DictionaryAttr> argAttrs) {
- result.addAttribute(SymbolTable::getSymbolAttrName(),
+ result.addAttribute(getSymNameAttrName(result.name),
builder.getStringAttr(name));
result.addAttribute("function_type", TypeAttr::get(type));
result.attributes.append(attrs.begin(), attrs.end());
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.td b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.td
index 2e6a99e..1082aea 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.td
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.td
@@ -33,6 +33,8 @@
def VM_ModuleOp : VM_Op<"module", [
IsolatedFromAbove,
SingleBlockImplicitTerminator<"IREE::VM::ModuleTerminatorOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
SymbolTable,
]> {
@@ -95,6 +97,8 @@
def VM_FuncOp : VM_Op<"func", [
IsolatedFromAbove,
HasParent<"IREE::VM::ModuleOp">,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
CallableOpInterface,
Symbol,
@@ -106,6 +110,8 @@
}];
let arguments = (ins
+ SymbolNameAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
TypeAttrOf<FunctionType>:$function_type,
OptionalAttr<VM_Ordinal>:$ordinal,
OptionalAttr<Util_InliningPolicyAttrInterface>:$inlining_policy,
@@ -196,6 +202,8 @@
}
def VM_ImportOp : VM_Op<"import", [
+ SymbolName,
+ SymbolVisibility,
Symbol,
FunctionOpInterface,
CallableOpInterface,
@@ -286,6 +294,8 @@
def VM_InitializerOp : VM_Op<"initializer", [
IsolatedFromAbove,
HasParent<"IREE::VM::ModuleOp">,
+ SymbolName,
+ SymbolVisibility,
FunctionOpInterface,
CallableOpInterface,
Util_InitializerOpInterface,
@@ -298,6 +308,8 @@
}];
let arguments = (ins
+ OptionalAttr<SymbolNameAttr>:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
TypeAttrOf<FunctionType>:$function_type,
OptionalAttr<DictArrayAttr>:$arg_attrs,
OptionalAttr<DictArrayAttr>:$res_attrs
@@ -365,6 +377,8 @@
VM_Op<mnemonic, !listconcat(traits, [
IsolatedFromAbove,
HasParent<"IREE::VM::ModuleOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
DeclareOpInterfaceMethods<Util_GlobalOpInterface, [
"createLoadOp",
@@ -397,7 +411,7 @@
CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs
),
[{
- $_state.addAttribute(SymbolTable::getSymbolAttrName(),
+ $_state.addAttribute(getSymNameAttrName($_state.name),
$_builder.getStringAttr(name));
if (isMutable) {
$_state.addAttribute("is_mutable", $_builder.getUnitAttr());
@@ -470,6 +484,8 @@
VM_Op<"global.ref", [
IsolatedFromAbove,
HasParent<"IREE::VM::ModuleOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
DeclareOpInterfaceMethods<Util_GlobalOpInterface>,
]> {
@@ -504,7 +520,7 @@
CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs
),
[{
- $_state.addAttribute(SymbolTable::getSymbolAttrName(),
+ $_state.addAttribute(getSymNameAttrName($_state.name),
$_builder.getStringAttr(name));
if (isMutable) {
$_state.addAttribute("is_mutable", $_builder.getUnitAttr());
@@ -1193,6 +1209,8 @@
def VM_RodataOp : VM_Op<"rodata", [
IsolatedFromAbove,
HasParent<"IREE::VM::ModuleOp">,
+ SymbolName,
+ SymbolVisibility,
Symbol,
]> {
let summary = [{Read-only data definition operation.}];
diff --git a/compiler/src/iree/compiler/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir b/compiler/src/iree/compiler/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir
index 6a719f6..7d37379 100644
--- a/compiler/src/iree/compiler/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir
+++ b/compiler/src/iree/compiler/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir
@@ -2,7 +2,7 @@
// RUN: --iree-vm-conversion --canonicalize %s | FileCheck %s
// CHECK-DAG: vm.import private @hal_loader.executable.dispatch{{\(.+}}attributes {minimum_version = 1 : i32}
-// CHECK-DAG: vm.import private @hal_loader.executable.lookup.function{{.*}}attributes {minimum_version = 1 : i32, nosideeffects}
+// CHECK-DAG: vm.import private @hal_loader.executable.lookup.function{{.*}}attributes {nosideeffects, minimum_version = 1 : i32}
// CHECK-LABEL: @executableLoad
// CHECK-SAME: (%[[EXECUTABLE_DATA:.+]]: !vm.buffer)
diff --git a/compiler/src/iree/compiler/Utils/EquivalenceUtils.cpp b/compiler/src/iree/compiler/Utils/EquivalenceUtils.cpp
index ea0c9bb..7ecc334 100644
--- a/compiler/src/iree/compiler/Utils/EquivalenceUtils.cpp
+++ b/compiler/src/iree/compiler/Utils/EquivalenceUtils.cpp
@@ -17,8 +17,7 @@
OperationEquivalenceCache::OperationEquivalenceCache(MLIRContext *context)
: functionRefName(StringAttr::get(context, "function_ref")),
- symbolAttrName(
- StringAttr::get(context, SymbolTable::getSymbolAttrName())) {}
+ symbolAttrName(StringAttr::get(context, "sym_name")) {}
OperationEquivalenceCache::~OperationEquivalenceCache() {
for (auto *mapping : mappingFreeList) {
diff --git a/compiler/src/iree/compiler/Utils/EquivalenceUtils.h b/compiler/src/iree/compiler/Utils/EquivalenceUtils.h
index 9bf8fdb..74f2e1a 100644
--- a/compiler/src/iree/compiler/Utils/EquivalenceUtils.h
+++ b/compiler/src/iree/compiler/Utils/EquivalenceUtils.h
@@ -59,7 +59,7 @@
private:
StringAttr functionRefName; // "function_ref"
- StringAttr symbolAttrName; // SymbolTable::getSymbolAttrName()
+ StringAttr symbolAttrName; // Conventional "sym_name" attribute.
SmallVector<IRMapping *> mappingFreeList;
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 53e39b2..e65959f 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 53e39b225bd164efeae02c9a0721021469d5bc67
+Subproject commit e65959f1684bd2416e8aa2ce684b8c300c381902
diff --git a/third_party/torch-mlir b/third_party/torch-mlir
index ec8720b..77139e3 160000
--- a/third_party/torch-mlir
+++ b/third_party/torch-mlir
@@ -1 +1 @@
-Subproject commit ec8720bf26b0c7b35f98b93730b75fb2f4b0fccb
+Subproject commit 77139e314d560acbf482a5509237c5137d33f969
