Use cleaner symbol visibility syntax for hal types. (#6968)
Fixes https://github.com/google/iree/issues/4670
diff --git a/docs/developers/design_docs/codegen_passes.md b/docs/developers/design_docs/codegen_passes.md
index 7326091..d4d76d1 100644
--- a/docs/developers/design_docs/codegen_passes.md
+++ b/docs/developers/design_docs/codegen_passes.md
@@ -36,7 +36,7 @@
offset = %c0 : tensor<32x16xf32>
return
}
- hal.interface @legacy_io attributes {sym_visibility = "private"} {
+ hal.interface private @legacy_io {
hal.interface.binding @arg0, set=0, binding=0,
type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1,
@@ -71,7 +71,7 @@
offset = %c0 : tensor<10x15xf32>
return
}
- hal.interface @legacy_io attributes {sym_visibility = "private"} {
+ hal.interface private @legacy_io {
hal.interface.binding @arg0, set=0, binding=0,
type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1,
diff --git a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
index 6383fb5..4cdb10a 100644
--- a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
+++ b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
@@ -17,7 +17,7 @@
return
}
-hal.interface @interface_io attributes {sym_visibility = "private"} {
+hal.interface private @interface_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=0, type="StorageBuffer", access="Write|Discard"
}
@@ -42,7 +42,7 @@
return
}
-hal.interface @interface_io attributes {sym_visibility = "private"} {
+hal.interface private @interface_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=0, type="StorageBuffer", access="Write|Discard"
}
diff --git a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
index aba1c1b..59527c3 100644
--- a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
+++ b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
@@ -6,7 +6,7 @@
return %val: f32
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
}
@@ -28,7 +28,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_xw_external, set=0, binding=0, type="StorageBuffer", access="Write|Discard"
}
@@ -49,7 +49,7 @@
return %val: vector<4xf32>
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
}
@@ -65,7 +65,7 @@
return %val: f16
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
}
@@ -82,7 +82,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_xw_external, set=0, binding=0, type="StorageBuffer", access="Write|Discard"
}
@@ -176,7 +176,7 @@
return %val: f32
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
}
@@ -215,7 +215,7 @@
%val = vector.transfer_read %subspan[%arg1, %arg2, %arg3], %cst {in_bounds = [true]} : memref<6x7x8xf32>, vector<4xf32>
return %val: vector<4xf32>
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @ro, set=0, binding=0, type="StorageBuffer", access="Read"
}
// CHECK: #[[MAP:.+]] = affine_map<()[s0, s1, s2] -> (s0 * 56 + s1 * 8 + s2)>
@@ -237,7 +237,7 @@
vector.transfer_write %arg4, %subspan[%arg1, %arg2, %arg3] {in_bounds = [true]} : vector<4xf32>, memref<6x7x8xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @ro, set=0, binding=0, type="StorageBuffer", access="Read|Write"
}
// CHECK: #[[MAP:.+]] = affine_map<()[s0, s1, s2] -> (s0 * 56 + s1 * 8 + s2)>
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
index 899692f..7896654 100644
--- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
+++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -27,7 +27,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -77,7 +77,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -125,7 +125,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -182,7 +182,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -243,7 +243,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -300,7 +300,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -358,7 +358,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -409,7 +409,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -461,7 +461,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @TENSOR_LHS, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_RHS, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @TENSOR_INIT, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -523,7 +523,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @arg2, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -600,7 +600,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @arg2, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -647,7 +647,7 @@
flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -681,7 +681,7 @@
flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -720,7 +720,7 @@
flow.dispatch.tensor.store %4, %2, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
hal.interface.binding @ret1, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -760,7 +760,7 @@
flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -810,7 +810,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
@@ -845,7 +845,7 @@
flow.dispatch.tensor.store %7, %1, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -871,7 +871,7 @@
flow.dispatch.tensor.store %7, %1, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -902,7 +902,7 @@
flow.dispatch.tensor.store %11, %2, offsets = [%3, %5], sizes = [%6, %8], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
hal.interface.binding @ret1, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -931,7 +931,7 @@
flow.dispatch.tensor.store %6, %0, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<readwrite:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read|Write"
}
// CHECK-LABEL: func @slice_in_place()
@@ -953,7 +953,7 @@
flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<1x4xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -985,7 +985,7 @@
flow.dispatch.tensor.store %7, %2, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1015,7 +1015,7 @@
flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x9xi32> -> !flow.dispatch.tensor<writeonly:3x9xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -1036,7 +1036,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -1097,7 +1097,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -1145,7 +1145,7 @@
flow.dispatch.tensor.store %7, %2, offsets = [], sizes = [], strides = [] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1181,7 +1181,7 @@
flow.dispatch.tensor.store %9, %2, offsets = [], sizes = [], strides = [] : tensor<1x2x2x1xf32> -> !flow.dispatch.tensor<writeonly:1x2x2x1xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1256,7 +1256,7 @@
}
}
return
-}hal.interface @io attributes {sym_visibility = "private"} {
+}hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1345,7 +1345,7 @@
return
}
-hal.interface @interface_io attributes {sym_visibility = "private"} {
+hal.interface private @interface_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ro2, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -1442,7 +1442,7 @@
return
}
-hal.interface @interface_io attributes {sym_visibility = "private"} {
+hal.interface private @interface_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @wo1, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -1530,7 +1530,7 @@
flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = [] : tensor<?x?x?xf32> -> !flow.dispatch.tensor<readwrite:?x?x?xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
@@ -1580,7 +1580,7 @@
flow.dispatch.tensor.store %25, %3, offsets = [%c0, %c0], sizes = [%c1, %c1], strides = [%c1, %c1] : tensor<2x1xf32> -> !flow.dispatch.tensor<writeonly:2x4xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @arg2, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -1641,7 +1641,7 @@
flow.dispatch.tensor.store %25, %3, offsets = [%c0, %c0], sizes = [%c1, %c1], strides = [%c1, %c1] : tensor<2x1xf32> -> !flow.dispatch.tensor<readwrite:2x4xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=3, type="StorageBuffer", access="Write|Discard"
@@ -1716,7 +1716,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1845,7 +1845,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -1931,7 +1931,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -2003,7 +2003,7 @@
}
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -2091,7 +2091,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @ro1, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro2, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ro3, set=0, binding=2, type="StorageBuffer", access="Read"
@@ -2196,7 +2196,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @ro1, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro2, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo, set=0, binding=2, type="StorageBuffer", access="Read|Write"
@@ -2297,7 +2297,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -2362,7 +2362,7 @@
flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<4xi32> -> !flow.dispatch.tensor<writeonly:4xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @ro, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @wo, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -2409,7 +2409,7 @@
flow.dispatch.tensor.store %3, %0, offsets = [], sizes = [], strides = [] : tensor<4xi32> -> !flow.dispatch.tensor<readwrite:4xi32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @rw, set=0, binding=0, type="StorageBuffer", access="Read|Write"
}
// CHECK-LABEL: func @sort1D_inplace()
diff --git a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
index 4e62995..5f8e1c7 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
@@ -40,9 +40,7 @@
"test.sink"(%memref) : (memref<?x2xf32>) -> ()
return
}
-hal.interface @io attributes {push_constants = 2 : index, sym_visibility = "private"} {
+hal.interface private @io attributes {push_constants = 2 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write"
}
-
-
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
index 1fc4dbc..38471fc 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(iree-llvmcpu-lower-executable-target{test-lowering-configuration=true}))' -cse -canonicalize -split-input-file %s | IreeFileCheck %s
-hal.executable @matmul_tensors attributes {sym_visibility = "private"} {
+hal.executable private @matmul_tensors {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -61,7 +61,7 @@
// CHECK-DAG: #[[CONFIG:.+]] = {nativeVectorSize = [4, 4, 4], tileSizes = {{\[}}[64, 64], [32, 32, 32], [4, 4, 4]{{\]}}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
-// CHECK: hal.executable.entry_point @matmul_tensors
+// CHECK: hal.executable.entry_point public @matmul_tensors
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -76,11 +76,11 @@
// CHECK: #[[CONFIG:.+]] = {passPipeline = 0 : i32}
// CHECK-NOT: #config
-// CHECK: hal.executable.entry_point @add_no_config
+// CHECK: hal.executable.entry_point public @add_no_config
// CHECK-SAME: translation.info = #[[CONFIG]]
// CHECK-NOT: #config
-hal.executable @add_no_config attributes {sym_visibility = "private"} {
+hal.executable private @add_no_config {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -115,7 +115,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -126,7 +126,7 @@
// -----
-hal.executable @add attributes {sym_visibility = "private"} {
+hal.executable private @add {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -184,7 +184,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -194,7 +194,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[64, 64]{{\]}}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
-// CHECK: hal.executable.entry_point @add
+// CHECK: hal.executable.entry_point public @add
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -207,7 +207,7 @@
// -----
-hal.executable @add4D attributes {sym_visibility = "private"} {
+hal.executable private @add4D {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -288,7 +288,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -298,7 +298,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[0, 64, 64, 64]{{\]}}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
-// CHECK: hal.executable.entry_point @add4D
+// CHECK: hal.executable.entry_point public @add4D
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -311,7 +311,7 @@
// -----
-hal.executable @batch_matmul_tensors attributes {sym_visibility = "private"} {
+hal.executable private @batch_matmul_tensors {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -382,7 +382,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = {nativeVectorSize = [1, 4, 4, 4], tileSizes = {{\[}}[1, 32, 32], [1, 16, 16, 16], [1, 4, 4, 4]{{\]}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
-// CHECK: hal.executable.entry_point @batch_matmul_tensors
+// CHECK: hal.executable.entry_point public @batch_matmul_tensors
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9]+]]: index)
@@ -394,7 +394,7 @@
// -----
-hal.executable @preset_config_matmul_tensors attributes {sym_visibility = "private"} {
+hal.executable private @preset_config_matmul_tensors {
hal.executable.variant @system_elf_x86_64, target = #hal.executable.target<"llvm", "system-elf-x86_64"> {
hal.executable.entry_point @preset_config attributes {interface = @io, ordinal = 0 : index}
builtin.module {
@@ -434,7 +434,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
index 5bbb7ee..ffd662e 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
@@ -2,7 +2,7 @@
// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.variant(iree-llvmcpu-lower-executable-target{use-lowering-pipeline='builtin.func(iree-llvmcpu-vectorization{promote-workgroup-to-full-tiles}),cse'}))" -split-input-file %s | IreeFileCheck %s -check-prefix=CHECK-PROMOTED
#config = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
-hal.executable @dynamic_matmul attributes {sym_visibility = "private"} {
+hal.executable private @dynamic_matmul {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -115,7 +115,7 @@
// -----
#config = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
-hal.executable @matmul_i8_i8_i32 attributes {sym_visibility = "private"} {
+hal.executable private @matmul_i8_i8_i32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
diff --git a/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir b/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
index 9b2726e..e601390 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
@@ -38,7 +38,7 @@
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir b/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
index 4a91e6e..cdf78d1 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
@@ -40,7 +40,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
index 59a1314..00e5c31 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
@@ -19,7 +19,7 @@
memref.store %12, %2[%7] : memref<16xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=4, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=1, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
index fc6d70e..3dd2b57 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -17,7 +17,7 @@
memref.store %11, %0[%7] : memref<16xf32>
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
index 1d5faba..b24f3ff 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -7,7 +7,7 @@
#map2 = affine_map<(d0) -> (2, -d0 + 1024)>
#map3 = affine_map<(d0) -> (256, -d0 + 1024)>
#map4 = affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>
-hal.executable @dot_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @dot_dispatch_0 {
hal.executable.variant @cuda, target = #executable_target_cuda_nvptx_fb {
hal.executable.entry_point @dot_dispatch_0 attributes {
interface = @legacy_io,
@@ -50,7 +50,7 @@
}
return
}
- hal.interface @legacy_io attributes {sym_visibility = "private"} {
+ hal.interface private @legacy_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -59,8 +59,8 @@
}
}
-// CHECK-LABEL: hal.executable @dot_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable private @dot_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK: memref.global "private" @{{.*}} : memref<4x256xf32, 3>
// CHECK: memref.global "private" @{{.*}} : memref<2x4xf32, 3>
// CHECK-DAG: %[[C0:.+]] = constant 0 : index
@@ -112,7 +112,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -120,7 +120,7 @@
}
}
// CHECK: #[[CONFIG:.+]] = {tileSizes = {{\[}}[]{{\]}}}
-// CHECK: hal.executable @reduction_dispatch
+// CHECK: hal.executable public @reduction_dispatch
// CHECK: linalg.fill
// CHECK-SAME: lowering.config = #[[CONFIG]]
// CHECK: linalg.generic
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
index 4bbcb4b..47e7b94 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
@@ -7,7 +7,7 @@
// CHECK-DAG: #[[$MAP4:.*]] = affine_map<()[s0, s1, s2] -> (s0 + s1 * 32 + s2 * 128 + 128)>
// CHECK-DAG: #[[$MAP5:.*]] = affine_map<()[s0, s1, s2] -> (s0 * 4 + s1 * 128 + s2 * 512)>
-hal.executable @shared_mem_cpy attributes {sym_visibility = "private"} {
+hal.executable private @shared_mem_cpy {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @shared_mem_cpy attributes {
interface = @io,
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
index 0a02226..89ac158 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -24,7 +24,7 @@
flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16384xf32> -> !flow.dispatch.tensor<writeonly:16384xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -35,7 +35,7 @@
// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[128], [], [4]{{\]}}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)>
-// CHECK: hal.executable.entry_point @add_dispatch_0
+// CHECK: hal.executable.entry_point public @add_dispatch_0
// CHECK-SAME: passPipeline = 3 : i32
// CHECK-SAME: workloadPerWorkgroup = [128]
// CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index]
@@ -49,7 +49,7 @@
// -----
-hal.executable @dot_dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @dot_dispatch_1 {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @dot_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : index}
builtin.module {
@@ -84,7 +84,7 @@
}
return
}
- hal.interface @legacy_io attributes {sym_visibility = "private"} {
+ hal.interface private @legacy_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -95,7 +95,7 @@
// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[4, 2, 4], [], [1, 1]{{\]}}}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 2)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
-// CHECK: hal.executable.entry_point @dot_dispatch_1
+// CHECK: hal.executable.entry_point public @dot_dispatch_1
// CHECK-SAME: passPipeline = 4 : i32
// CHECK-SAME: workloadPerWorkgroup = [2, 4]
// CHECK-SAME: workgroup_size = [2 : index, 4 : index, 1 : index]
@@ -135,7 +135,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -145,7 +145,7 @@
// CHECK-DAG: #[[CONFIG0:.+]] = {passPipeline = 2 : i32}
// CHECK-DAG: #[[CONFIG1:.+]] = {tileSizes = {{\[}}[]{{\]}}}
-// CHECK: hal.executable.entry_point @predict_dispatch_153
+// CHECK: hal.executable.entry_point public @predict_dispatch_153
// CHECK-SAME: translation.info = #[[CONFIG0]]
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index,
diff --git a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
index d6025c4..9f438a7 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -27,7 +27,7 @@
flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -36,8 +36,8 @@
}
}
-// CHECK-LABEL: hal.executable @simpleMath_ex_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @simpleMath_ex_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK: llvm.fadd
// -----
@@ -45,7 +45,7 @@
#map0 = affine_map<()[s0, s1] -> (s0 * s1)>
#map1 = affine_map<(d0)[s0] -> (s0, -d0 + 1024)>
#map2 = affine_map<(d0)[s0] -> (-d0 + 1024, s0)>
-hal.executable @dot_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable @dot_dispatch_0 {
hal.interface @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -90,7 +90,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -99,8 +99,8 @@
}
}
-// CHECK-LABEL: hal.executable @dot_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @dot_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK-NOT: llvm.store
// CHECK-COUNT-3: llvm.load {{.*}} : !llvm.ptr<vector<4xf32>>
// CHECK: llvm.br
@@ -133,7 +133,7 @@
iterator_types = ["parallel", "parallel", "reduction"]
}
-hal.executable @dot_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable @dot_dispatch_0 {
hal.interface @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -183,7 +183,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -192,8 +192,8 @@
}
}
-// CHECK-LABEL: hal.executable @dot_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @dot_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK: llvm.br
// CHECK-COUNT-8: "llvm.intr.fmuladd"({{.*}}) : (vector<4xf32>, vector<4xf32>, vector<4xf32>) -> vector<4xf32>
// CHECK: llvm.br
@@ -201,7 +201,7 @@
// -----
-hal.executable @conv2d_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable @conv2d_dispatch_0 {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.interface @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -257,7 +257,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -266,8 +266,8 @@
}
}
-// CHECK-LABEL: hal.executable @conv2d_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @conv2d_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK-COUNT-3: llvm.load %{{.*}} : !llvm.ptr<f32>
// CHECK: lvm.fmul %{{.*}}, %{{.*}} : f32
// CHECK: llvm.fadd %{{.*}}, %{{.*}} : f32
@@ -298,7 +298,7 @@
flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -306,8 +306,8 @@
}
}
-// CHECK-LABEL: hal.executable @simpleMath_ex_dispatch_0
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @simpleMath_ex_dispatch_0
+// CHECK: hal.executable.variant public @cuda
// CHECK: llvm.mlir.global private constant @{{.*}}(dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00, 7.000000e+00, 8.000000e+00, 9.000000e+00, 1.000000e+01, 1.100000e+01, 1.200000e+01, 1.300000e+01, 1.400000e+01, 1.500000e+01, 1.600000e+01]> : tensor<16xf32>)
// CHECK: llvm.fadd
@@ -344,7 +344,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -352,8 +352,8 @@
}
}
-// CHECK-LABEL: hal.executable @reduction_dispatch
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @reduction_dispatch
+// CHECK: hal.executable.variant public @cuda
// CHECK: llvm.fadd
// -----
@@ -389,7 +389,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -398,8 +398,8 @@
}
}
-// CHECK-LABEL: hal.executable @vector_add_dispatch
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @vector_add_dispatch
+// CHECK: hal.executable.variant public @cuda
// CHECK: llvm.fadd %{{.*}}, %{{.*}} : vector<4xf32
// CHECK: llvm.store %{{.*}} : !llvm.ptr<vector<4xf32>>
@@ -442,7 +442,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
@@ -450,7 +450,7 @@
}
}
-// CHECK-LABEL: hal.executable @vector_reduction_dispatch
-// CHECK: hal.executable.variant @cuda
+// CHECK-LABEL: hal.executable public @vector_reduction_dispatch
+// CHECK: hal.executable.variant public @cuda
// CHECK-COUNT-4: llvm.fadd
// CHECK: llvm.store %{{.*}} : !llvm.ptr<vector<4xf32>>
diff --git a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
index 6612b5f..59bb295 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-llvmgpu-remove-single-iteration-loop))))' %s | IreeFileCheck %s
// CHECK-LABEL: func @dispatch_0()
-hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_0 {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @dispatch_0 attributes {
interface = @io,
diff --git a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
index 8d4ec2a..f2deab1 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
@@ -27,7 +27,7 @@
flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -36,8 +36,8 @@
}
}
-// CHECK-LABEL: hal.executable @simpleMath_ex_dispatch_0
-// CHECK: hal.executable.variant @rocm
+// CHECK-LABEL: hal.executable public @simpleMath_ex_dispatch_0
+// CHECK: hal.executable.variant public @rocm
// CHECK: llvm.fadd
// -----
@@ -45,7 +45,7 @@
#map0 = affine_map<()[s0, s1] -> (s0 * s1)>
#map1 = affine_map<(d0)[s0] -> (s0, -d0 + 1024)>
#map2 = affine_map<(d0)[s0] -> (-d0 + 1024, s0)>
-hal.executable @dot_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable @dot_dispatch_0 {
hal.interface @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -90,7 +90,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ro1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -99,8 +99,8 @@
}
}
-// CHECK-LABEL: hal.executable @dot_dispatch_0
-// CHECK: hal.executable.variant @rocm
+// CHECK-LABEL: hal.executable public @dot_dispatch_0
+// CHECK: hal.executable.variant public @rocm
// CHECK-NOT: llvm.store
// CHECK-COUNT-3: llvm.load {{.*}} : !llvm.ptr<vector<4xf32>>
// CHECK: llvm.br
diff --git a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
index eb46b1c..922fd2a 100644
--- a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(iree-convert-to-spirv)))' %s | IreeFileCheck %s
-hal.executable @push_constant attributes {sym_visibility = "private"} {
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.executable private @push_constant {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write"
}
@@ -25,7 +25,7 @@
return
}
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write"
}
@@ -35,8 +35,8 @@
// -----
-hal.executable @resource_bindings_in_same_func attributes {sym_visibility = "private"} {
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.executable private @resource_bindings_in_same_func {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=1, binding=3, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
@@ -83,7 +83,7 @@
return
}
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=1, binding=3, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
@@ -94,8 +94,8 @@
// -----
-hal.executable @resource_bindings_in_multi_entry_func attributes {sym_visibility = "private"} {
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.executable private @resource_bindings_in_multi_entry_func {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -144,7 +144,7 @@
return
}
- hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+ hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -154,8 +154,8 @@
// -----
-hal.executable @interface_binding attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @interface_binding {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -179,7 +179,7 @@
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -201,8 +201,8 @@
// -----
-hal.executable @interface_wg_id attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @interface_wg_id {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -219,7 +219,7 @@
%1 = hal.interface.workgroup.id[1] : index
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -240,8 +240,8 @@
// -----
-hal.executable @interface_wg_count attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @interface_wg_count {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -258,7 +258,7 @@
%1 = hal.interface.workgroup.count[1] : index
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir
index df01d23..9f4ea2a 100644
--- a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-distribute-to-global-id))))' -canonicalize -cse %s | IreeFileCheck %s
#map0 = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
-hal.executable @parallel_4D attributes {sym_visibility = "private"} {
+hal.executable private @parallel_4D {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -41,7 +41,7 @@
func private @parallel_4D__num_workgroups__
(!shapex.ranked_shape<[?,?,?,?]>, !shapex.ranked_shape<[?,?,?,?]>,
!shapex.ranked_shape<[?,?,?,?]>) -> (index, index, index)
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -81,7 +81,7 @@
// -----
#map0 = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
-hal.executable @parallel_4D_static attributes {sym_visibility = "private"} {
+hal.executable private @parallel_4D_static {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -114,7 +114,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -153,7 +153,7 @@
iterator_types = []
}
-hal.executable @scalar_add attributes {sym_visibility = "private"} {
+hal.executable private @scalar_add {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -184,7 +184,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -202,7 +202,7 @@
// -----
// TODO(GH-4901): Convert these tests back to use dynamic shapes when linalg on tensors becomes default.
-hal.executable @reduce_sum attributes {sym_visibility = "private"} {
+hal.executable private @reduce_sum {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -239,7 +239,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
index 9bba3e7..3c76311 100644
--- a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-fold-gpu-procid-uses))))' %s | IreeFileCheck %s
-hal.executable @fold_block_id attributes {sym_visibility = "private"} {
+hal.executable private @fold_block_id {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
@@ -35,7 +35,7 @@
// -----
-hal.executable @fold_interface_workgroup_id attributes {sym_visibility = "private"} {
+hal.executable private @fold_interface_workgroup_id {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
@@ -70,7 +70,7 @@
// -----
-hal.executable @fold_thread_id attributes {sym_visibility = "private"} {
+hal.executable private @fold_thread_id {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
@@ -100,7 +100,7 @@
// -----
-hal.executable @does_not_fold_mod attributes {sym_visibility = "private"} {
+hal.executable private @does_not_fold_mod {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
@@ -123,7 +123,7 @@
// -----
-hal.executable @does_not_fold_div attributes {sym_visibility = "private"} {
+hal.executable private @does_not_fold_div {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
@@ -146,7 +146,7 @@
// -----
-hal.executable @does_not_fold_symbol_mul_symbol attributes {sym_visibility = "private"} {
+hal.executable private @does_not_fold_symbol_mul_symbol {
hal.interface @io {
}
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
index 045ce99..4306690 100644
--- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-codegen-linalg-to-spirv-pipeline))' %s | IreeFileCheck %s
// TODO(#5608): iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-codegen-linalg-to-spirv-pipeline))' -iree-spirv-use-workgroup-memory %s | IreeFileCheck %s
-hal.executable @matmul_cooperative_matrix attributes {sym_visibility = "private"} {
+hal.executable private @matmul_cooperative_matrix {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -52,7 +52,7 @@
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
index dd4d216..a7716d9 100644
--- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
@@ -2,7 +2,7 @@
#config = {tileSizes = [[8, 64, 4], [], [8, 4, 4]]}
-hal.executable @fuse_and_vectorize_fill_matmul attributes {sym_visibility = "private"} {
+hal.executable private @fuse_and_vectorize_fill_matmul {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -51,7 +51,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -72,7 +72,7 @@
#config = {tileSizes = [[8, 64, 4], [], [8, 4, 4]]}
-hal.executable @fuse_and_vectorize_matmul_add attributes {sym_visibility = "private"} {
+hal.executable private @fuse_and_vectorize_matmul_add {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -132,7 +132,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_ro_external, set=0, binding=2, type="StorageBuffer", access="Read"
diff --git a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
index f15772d..3154dc3 100644
--- a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
@@ -1,7 +1,7 @@
// TODO(antiagainst): Fix promotion to workgroup and enable the test.
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute,iree-spirv-vectorize,canonicalize,cse))))' | IreeFileCheck %s
-hal.executable @matmul_promote_workgroup_memory attributes {sym_visibility = "private"} {
+hal.executable private @matmul_promote_workgroup_memory {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -46,7 +46,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -79,7 +79,7 @@
// -----
-hal.executable @conv_promote_workgroup_memory attributes {sym_visibility = "private"} {
+hal.executable private @conv_promote_workgroup_memory {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -118,7 +118,7 @@
outs(%15 : memref<1x?x?x14xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 2002 + s0 + d1 * 154 + d2 * 14 + d3)>>)
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
index 19b488e..82f8e3e 100644
--- a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-remove-one-trip-tiled-loop))))' %s | IreeFileCheck %s
-hal.executable @static_shaped_conv attributes {sym_visibility = "private"} {
+hal.executable private @static_shaped_conv {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -55,7 +55,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/set_lowering_config.mlir b/iree/compiler/Codegen/SPIRV/test/set_lowering_config.mlir
index 599d580..1f892d9 100644
--- a/iree/compiler/Codegen/SPIRV/test/set_lowering_config.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/set_lowering_config.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
-hal.executable @static_1d_sort attributes {sym_visibility = "private"} {
+hal.executable private @static_1d_sort {
hal.interface @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
}
@@ -25,7 +25,7 @@
flow.dispatch.tensor.store %2, %0, offsets = [], sizes = [], strides = [] : tensor<1000xi32> -> !flow.dispatch.tensor<readwrite:1000xi32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
}
}
@@ -34,7 +34,7 @@
// Check that the workgroup count and size are (1, 1, 1) for serializing the computation.
-// CHECK-LABEL: hal.executable.entry_point @static_1d_sort
+// CHECK-LABEL: hal.executable.entry_point public @static_1d_sort
// CHECK-SAME: translation.info = {passPipeline = 6 : i32}
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%{{.+}}: index, %{{.+}}: index, %{{.+}}: index):
@@ -47,7 +47,7 @@
// -----
-hal.executable @static_3d_sort attributes {sym_visibility = "private"} {
+hal.executable private @static_3d_sort {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
@@ -98,7 +98,7 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point @static_3d_sort
+// CHECK-LABEL: hal.executable.entry_point public @static_3d_sort
// CHECK-SAME: translation.info = {passPipeline = 5 : i32, workloadPerWorkgroup = [16, 1]}
// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
@@ -112,7 +112,7 @@
// -----
-hal.executable @static_1d_fft attributes {sym_visibility = "private"} {
+hal.executable private @static_1d_fft {
hal.interface @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
@@ -140,7 +140,7 @@
flow.dispatch.tensor.store %4#1, %1, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
@@ -150,7 +150,7 @@
// Check that the workgroup count and size are (1, 1, 1) for serializing the computation.
-// CHECK-LABEL: hal.executable.entry_point @static_1d_fft
+// CHECK-LABEL: hal.executable.entry_point public @static_1d_fft
// CHECK-SAME: translation.info = {passPipeline = 6 : i32}
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%{{.+}}: index, %{{.+}}: index, %{{.+}}: index):
@@ -163,7 +163,7 @@
// -----
-hal.executable @static_3d_fft attributes {sym_visibility = "private"} {
+hal.executable private @static_3d_fft {
hal.interface @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
@@ -191,7 +191,7 @@
flow.dispatch.tensor.store %4#1, %1, offsets = [], sizes = [], strides = [] : tensor<64x128x32xf32> -> !flow.dispatch.tensor<readwrite:64x128x32xf32>
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
@@ -201,7 +201,7 @@
// Right now n-D fft does not support tiling too.
-// CHECK-LABEL: hal.executable.entry_point @static_3d_fft
+// CHECK-LABEL: hal.executable.entry_point public @static_3d_fft
// CHECK-SAME: translation.info = {passPipeline = 6 : i32}
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%{{.+}}: index, %{{.+}}: index, %{{.+}}: index):
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
index 8decc55..e3a62d9 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute))))' %s | IreeFileCheck %s
-hal.executable @static_scatter_update_slice attributes {sym_visibility = "private"} {
+hal.executable private @static_scatter_update_slice {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -44,7 +44,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_rw_external, set=0, binding=2, type="StorageBuffer", access="Read|Write"
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
index ebfc35c..8bf3918 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute))))' %s | IreeFileCheck %s
-hal.executable @static_3d_sort attributes {sym_visibility = "private"} {
+hal.executable private @static_3d_sort {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
index cf522ac..f0a227c 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
@@ -10,7 +10,7 @@
#config = {tileSizes = [[8, 16, 0], [], [1, 1, 1]]}
-hal.executable @matmul attributes {sym_visibility = "private"} {
+hal.executable private @matmul {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -67,7 +67,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -92,7 +92,7 @@
#config = {tileSizes = [[1, 4, 32], [], [1, 1, 1]]}
-hal.executable @conv_1d attributes {sym_visibility = "private"} {
+hal.executable private @conv_1d {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -137,7 +137,7 @@
linalg.conv_1d_nwc_wcf { __internal_linalg_transform__ = "workgroup", lowering.config = #config, dilations = dense<1> : tensor<1xi64>, strides = dense<1> : tensor<1xi64>} ins(%8, %11 : memref<1x?x1xf32, affine_map<(d0, d1, d2)[s0] -> (d0 * 8 + s0 + d1 + d2)>>, memref<3x1x?xf32, affine_map<(d0, d1, d2)[s0] -> (d0 + s0 + d1 + d2)>>) outs(%16 : memref<1x?x?xf32, affine_map<(d0, d1, d2)[s0] -> (d0 * 6 + s0 + d1 + d2)>>)
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -183,7 +183,7 @@
#config = {tileSizes = [[0, 1, 4, 32], [], [0, 1, 1, 1]]}
-hal.executable @conv_no_padding attributes {sym_visibility = "private"} {
+hal.executable private @conv_no_padding {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -266,7 +266,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -318,7 +318,7 @@
#config = {tileSizes = [[0, 0, 1, 4, 32], [], [0, 0, 1, 1, 1]]}
-hal.executable @conv_3d attributes {sym_visibility = "private"} {
+hal.executable private @conv_3d {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -361,7 +361,7 @@
linalg.conv_3d_ndhwc_dhwcf {__internal_linalg_transform__ = "workgroup", lowering.config = #config, dilations = dense<1> : tensor<3xi64>, strides = dense<1> : tensor<3xi64>} ins(%10, %2 : memref<1x?x?x8x3xf32, affine_map<(d0, d1, d2, d3, d4)[s0] -> (d0 * 1536 + s0 + d1 * 192 + d2 * 24 + d3 * 3 + d4)>>, memref<2x2x2x3x2xf32>) outs(%15 : memref<1x?x?x7x2xf32, affine_map<(d0, d1, d2, d3, d4)[s0] -> (d0 * 686 + s0 + d1 * 98 + d2 * 14 + d3 * 2 + d4)>>)
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -399,7 +399,7 @@
#config = {tileSizes = [[1, 4, 32], [], [1, 1, 1]]}
module {
- hal.executable @pooling_nhwc_max attributes {sym_visibility = "private"} {
+ hal.executable private @pooling_nhwc_max {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -433,7 +433,7 @@
linalg.pooling_nhwc_max {__internal_linalg_transform__ = "workgroup", lowering.config = #config, dilations = dense<1> : vector<2xi64>, strides = dense<1> : vector<2xi64>} ins(%9, %1 : memref<2x?x?x6xf32, #map4>, memref<3x4xf32>) outs(%12 : memref<2x?x?x6xf32, #map7>)
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
index d17b02a..887dc7b 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -2,8 +2,8 @@
#config = {tileSizes = [[1, 8, 64, 4], [], [1, 8, 4, 4]]}
-hal.executable @batch_matmul_static_shape attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @batch_matmul_static_shape {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -55,7 +55,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -371,8 +371,8 @@
#config = {tileSizes = [[1, 8, 64, 4], [], [1, 8, 4, 4]]}
-hal.executable @fused_fill_batch_matmul attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @fused_fill_batch_matmul {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -424,7 +424,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
index 3fa302b..4b6dbae 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -2,7 +2,7 @@
#config = {tileSizes = [[0, 4, 4, 16], [], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]]}
-hal.executable @conv_static_shape_f32 attributes {sym_visibility = "private"} {
+hal.executable private @conv_static_shape_f32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -66,7 +66,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -102,7 +102,7 @@
#config = {tileSizes = [[0, 2, 2, 32], [], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]]}
-hal.executable @depthwise_conv_static_shape_f32 attributes {sym_visibility = "private"} {
+hal.executable private @depthwise_conv_static_shape_f32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -168,7 +168,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
index 545757a..3822bdf 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -2,8 +2,8 @@
#config = {tileSizes = [[8, 64, 4], [], [8, 4, 4]]}
-hal.executable @matmul_static_shape_f16 attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @matmul_static_shape_f16 {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -46,7 +46,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -69,8 +69,8 @@
#config = {tileSizes = [[8, 64, 4], [], [8, 4, 4]]}
-hal.executable @matmul_static_shape_f32 attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @matmul_static_shape_f32 {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -113,7 +113,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
index 922d627..07cec4f 100644
--- a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
@@ -5,7 +5,7 @@
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
-hal.executable @kernel_matmul attributes {sym_visibility = "private"} {
+hal.executable private @kernel_matmul {
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
gpu.container_module,
spv.target_env = #spv.target_env<#spv.vce<v1.0,
@@ -41,7 +41,7 @@
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
-hal.executable @kernel_matmul attributes {sym_visibility = "private"} {
+hal.executable private @kernel_matmul {
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
gpu.container_module,
spv.target_env = #spv.target_env<#spv.vce<v1.0,
@@ -88,7 +88,7 @@
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
-hal.executable @kernel_matmul attributes {sym_visibility = "private"} {
+hal.executable private @kernel_matmul {
hal.executable.variant @vulkan, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
gpu.container_module,
spv.target_env = #spv.target_env<#spv.vce<v1.0,
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
index 4f3f2fe..1a328e7 100644
--- a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
@@ -5,7 +5,7 @@
// CHECK: vector.transfer_read %{{.+}}[%c0], {{.+}} memref<4xf32, #{{.+}}>, vector<4xf32>
// CHECK: addf %{{.*}}, %{{.*}} : vector<4xf32>
// CHECK: vector.transfer_write {{.*}} : vector<4xf32>, memref<4xf32
-hal.executable @elementwise_static_shape attributes {sym_visibility = "private"} {
+hal.executable private @elementwise_static_shape {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -43,7 +43,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -60,7 +60,7 @@
// CHECK-NOT: vector.transfer_read
// CHECK: scf.for
// CHECK: scf.for
-hal.executable @elementwise_transpose attributes {sym_visibility = "private"} {
+hal.executable private @elementwise_transpose {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -98,7 +98,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -106,4 +106,3 @@
}
}
}
-
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
index a6f00e2..3bf7e59 100644
--- a/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
@@ -54,7 +54,7 @@
return
}
-hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -80,7 +80,7 @@
return
}
-hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -106,7 +106,7 @@
return
}
-hal.interface @io attributes {push_constants = 5 : index, sym_visibility = "private"} {
+hal.interface private @io attributes {push_constants = 5 : index} {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -182,7 +182,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=1, binding=2, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=3, binding=4, type="StorageBuffer", access="Write"
}
@@ -203,7 +203,7 @@
return
}
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
index f7ae29c..ff2578b 100644
--- a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
@@ -2,8 +2,8 @@
// TODO(antiagainst): Fix promotion to workgroup and enable the test.
// | IreeFileCheck %s -check-prefix=PROMOTE
-hal.executable @matmul_static_shape attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @matmul_static_shape {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -61,7 +61,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -264,8 +264,8 @@
// -----
-hal.executable @matmul_static_shape attributes {sym_visibility = "private"} {
- hal.interface @io attributes {sym_visibility = "private"} {
+hal.executable private @matmul_static_shape {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -323,7 +323,7 @@
}
return
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir
index eed33d2..7c70e95 100644
--- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir
+++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir
@@ -361,7 +361,7 @@
module attributes {hal.device.targets = [#hal.device.target<"vmvx">]} {
-hal.executable @ex attributes {sym_visibility = "private"} {
+hal.executable private @ex {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
@@ -409,7 +409,7 @@
module attributes {hal.device.targets = [#hal.device.target<"vmvx">]} {
-hal.executable @ex attributes {sym_visibility = "private"} {
+hal.executable private @ex {
hal.interface @io attributes {push_constants = 4 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
@@ -469,7 +469,7 @@
module attributes {hal.device.targets = [#hal.device.target<"vmvx">]} {
-hal.executable @pad_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @pad_dispatch_0 {
hal.interface @interface_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @wo1, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
@@ -483,7 +483,7 @@
}
}
-hal.executable @pad_dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @pad_dispatch_1 {
hal.interface @interface_io {
hal.interface.binding @ro0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @rw1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.cpp b/iree/compiler/Dialect/HAL/IR/HALOps.cpp
index 0cdb523..3652cf6 100644
--- a/iree/compiler/Dialect/HAL/IR/HALOps.cpp
+++ b/iree/compiler/Dialect/HAL/IR/HALOps.cpp
@@ -599,38 +599,6 @@
state.addAttribute("buffer_constraints", bufferConstraints);
}
-static ParseResult parseConstantPoolOp(OpAsmParser &parser,
- OperationState *result) {
- StringAttr nameAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
- result->attributes)) ||
- failed(parser.parseOptionalAttrDictWithKeyword(result->attributes))) {
- return failure();
- }
-
- // Parse the module body.
- auto *body = result->addRegion();
- if (failed(parser.parseRegion(*body, llvm::None, llvm::None))) {
- return failure();
- }
-
- // Ensure that this module has a valid terminator.
- ConstantPoolOp::ensureTerminator(*body, parser.getBuilder(),
- result->location);
- return success();
-}
-
-static void printConstantPoolOp(OpAsmPrinter &p, ConstantPoolOp op) {
- p << ' ';
- p.printSymbolName(op.sym_name());
- p.printOptionalAttrDictWithKeyword(
- op->getAttrs(),
- /*elidedAttrs=*/{mlir::SymbolTable::getSymbolAttrName()});
- p.printRegion(op.body(), /*printEntryBlockArgs=*/false,
- /*printBlockTerminators=*/false);
-}
-
//===----------------------------------------------------------------------===//
// hal.constant_pool.load
//===----------------------------------------------------------------------===//
@@ -842,37 +810,6 @@
builder.getStringAttr(name));
}
-static ParseResult parseExecutableOp(OpAsmParser &parser,
- OperationState *result) {
- StringAttr nameAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
- result->attributes)) ||
- failed(parser.parseOptionalAttrDictWithKeyword(result->attributes))) {
- return failure();
- }
-
- // Parse the module body.
- auto *body = result->addRegion();
- if (failed(parser.parseRegion(*body, llvm::None, llvm::None))) {
- return failure();
- }
-
- // Ensure that this module has a valid terminator.
- ExecutableOp::ensureTerminator(*body, parser.getBuilder(), result->location);
- return success();
-}
-
-static void printExecutableOp(OpAsmPrinter &p, ExecutableOp op) {
- p << ' ';
- p.printSymbolName(op.sym_name());
- p.printOptionalAttrDictWithKeyword(
- op->getAttrs(),
- /*elidedAttrs=*/{mlir::SymbolTable::getSymbolAttrName()});
- p.printRegion(op.body(), /*printEntryBlockArgs=*/false,
- /*printBlockTerminators=*/false);
-}
-
static LogicalResult verifyExecutableOp(ExecutableOp op) {
// TODO(benvanik): check export name conflicts.
return success();
@@ -884,6 +821,11 @@
static ParseResult parseExecutableEntryPointOp(OpAsmParser &parser,
OperationState *result) {
+ StringAttr visibilityAttr;
+ if (failed(parseSymbolVisibility(parser, visibilityAttr))) {
+ return failure();
+ }
+
StringAttr nameAttr;
if (failed(parser.parseSymbolName(nameAttr,
mlir::SymbolTable::getSymbolAttrName(),
@@ -907,6 +849,8 @@
static void printExecutableEntryPointOp(OpAsmPrinter &p,
ExecutableEntryPointOp op) {
p << ' ';
+ printSymbolVisibility(p, op, op->getAttrOfType<StringAttr>("sym_visibility"));
+ p << ' ';
p.printSymbolName(op.sym_name());
p.printOptionalAttrDictWithKeyword(op->getAttrs(),
/*elidedAttrs=*/{"sym_name"});
@@ -957,45 +901,6 @@
state.addAttribute("target", target);
}
-static ParseResult parseExecutableVariantOp(OpAsmParser &parser,
- OperationState *result) {
- auto *body = result->addRegion();
- StringAttr nameAttr;
- IREE::HAL::ExecutableTargetAttr targetAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
- result->attributes)) ||
- failed(parser.parseComma()) || failed(parser.parseKeyword("target")) ||
- failed(parser.parseEqual()) ||
- failed(parser.parseAttribute(targetAttr, "target", result->attributes)) ||
- failed(parser.parseOptionalAttrDictWithKeyword(result->attributes))) {
- return failure();
- }
-
- OptionalParseResult parseResult = parser.parseOptionalRegion(*body);
- if (parseResult.hasValue() && failed(*parseResult)) {
- return failure();
- }
-
- // Ensure that this module has a valid terminator.
- ExecutableVariantOp::ensureTerminator(*body, parser.getBuilder(),
- result->location);
- return success();
-}
-
-static void printExecutableVariantOp(OpAsmPrinter &p, ExecutableVariantOp op) {
- p << ' ';
- p.printSymbolName(op.sym_name());
- p << ", target = " << op.target();
- p.printOptionalAttrDictWithKeyword(
- op->getAttrs(),
- /*elidedAttrs=*/{mlir::SymbolTable::getSymbolAttrName(), "target"});
- if (!op.body().empty()) {
- p.printRegion(op.body(), /*printEntryBlockArgs=*/false,
- /*printBlockTerminators=*/false);
- }
-}
-
//===----------------------------------------------------------------------===//
// hal.executable.binary
//===----------------------------------------------------------------------===//
@@ -1022,26 +927,6 @@
state.addAttribute("data", data);
}
-static ParseResult parseExecutableBinaryOp(OpAsmParser &parser,
- OperationState *result) {
- StringAttr nameAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
- result->attributes)) ||
- failed(parser.parseOptionalAttrDictWithKeyword(result->attributes))) {
- return failure();
- }
- return success();
-}
-
-static void printExecutableBinaryOp(OpAsmPrinter &p, ExecutableBinaryOp op) {
- p << ' ';
- p.printSymbolName(op.sym_name());
- p.printOptionalAttrDictWithKeyword(
- op->getAttrs(),
- /*elidedAttrs=*/{mlir::SymbolTable::getSymbolAttrName()});
-}
-
//===----------------------------------------------------------------------===//
// hal.executable.create
//===----------------------------------------------------------------------===//
@@ -1074,37 +959,6 @@
}
}
-static ParseResult parseInterfaceOp(OpAsmParser &parser,
- OperationState *result) {
- StringAttr nameAttr;
- if (failed(parser.parseSymbolName(nameAttr,
- mlir::SymbolTable::getSymbolAttrName(),
- result->attributes)) ||
- failed(parser.parseOptionalAttrDictWithKeyword(result->attributes))) {
- return failure();
- }
-
- // Parse the module body.
- auto *body = result->addRegion();
- if (failed(parser.parseRegion(*body, llvm::None, llvm::None))) {
- return failure();
- }
-
- // Ensure that this module has a valid terminator.
- InterfaceOp::ensureTerminator(*body, parser.getBuilder(), result->location);
- return success();
-}
-
-static void printInterfaceOp(OpAsmPrinter &p, InterfaceOp op) {
- p << ' ';
- p.printSymbolName(op.sym_name());
- p.printOptionalAttrDictWithKeyword(
- op->getAttrs(),
- /*elidedAttrs=*/{mlir::SymbolTable::getSymbolAttrName()});
- p.printRegion(op.body(), /*printEntryBlockArgs=*/false,
- /*printBlockTerminators=*/false);
-}
-
ArrayAttr InterfaceOp::getExecutableSetLayoutsAttr() {
Builder builder(getContext());
SmallVector<SmallVector<Attribute, 4>, 4> setAttrs;
@@ -1153,6 +1007,11 @@
static ParseResult parseInterfaceBindingOp(OpAsmParser &parser,
OperationState *result) {
+ StringAttr visibilityAttr;
+ if (failed(parseSymbolVisibility(parser, visibilityAttr))) {
+ return failure();
+ }
+
StringAttr nameAttr;
IntegerAttr setAttr;
IntegerAttr bindingAttr;
@@ -1184,6 +1043,8 @@
static void printInterfaceBindingOp(OpAsmPrinter &p, InterfaceBindingOp op) {
p << ' ';
+ printSymbolVisibility(p, op, op->getAttrOfType<StringAttr>("sym_visibility"));
+ p << ' ';
p.printSymbolName(op.sym_name());
p << ", set=" << op.set();
p << ", binding=" << op.binding();
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.td b/iree/compiler/Dialect/HAL/IR/HALOps.td
index d779f78..3d14980 100644
--- a/iree/compiler/Dialect/HAL/IR/HALOps.td
+++ b/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -1237,7 +1237,7 @@
]> {
let summary = [{pool of constants with similar lifetimes}];
let description = [{
- A pool of constants that share a similiar lifetime and that should be stored
+ A pool of constants that share a similar lifetime and that should be stored
together both in the source files and at runtime. By logically grouping
constants by their frequency and locality of access we can reduce the number
of bindings required on hal.interface by sourcing constants from the same
@@ -1246,12 +1246,21 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
HAL_BufferConstraintsAttr:$buffer_constraints
);
let regions = (region SizedRegion<1>:$body);
+ let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
+ $sym_name
+ attr-dict-with-keyword
+ ``
+ regions
+ }];
+
let skipDefaultBuilders = 1;
let builders = [
OpBuilder<(ins
@@ -1284,13 +1293,24 @@
}];
let arguments = (ins
+ OptionalAttr<StrAttr>:$sym_visibility,
SymbolNameAttr:$sym_name,
Util_AnySerializableAttr:$value
);
let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
$sym_name `=` $value attr-dict-with-keyword
}];
+
+ let builders = [
+ OpBuilder<(ins
+ "StringRef":$sym_name,
+ "::mlir::Attribute":$value
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, value);
+ }]>
+ ];
}
def HAL_ConstantPoolSpanOp : HAL_Op<"constant_pool.span", [
@@ -1305,6 +1325,7 @@
}];
let arguments = (ins
+ OptionalAttr<StrAttr>:$sym_visibility,
SymbolNameAttr:$sym_name,
TypeAttr:$tensor_type,
SymbolRefAttr:$storage_buffer,
@@ -1314,11 +1335,26 @@
);
let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
$sym_name `:` $tensor_type
`=` $storage_buffer `[` $storage_range `]`
(`->` $runtime_buffer^ `[` $runtime_range `]`)?
attr-dict-with-keyword
}];
+
+ let builders = [
+ OpBuilder<(ins
+ "StringRef":$sym_name,
+ "::mlir::Type":$tensor_type,
+ "::mlir::FlatSymbolRefAttr":$storage_buffer,
+ "::mlir::Attribute":$storage_range,
+ "::mlir::SymbolRefAttr":$runtime_buffer,
+ "::mlir::Attribute":$runtime_range
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, tensor_type, storage_buffer,
+ storage_range, runtime_buffer, runtime_range);
+ }]>
+ ];
}
def HAL_ConstantPoolSplatOp : HAL_Op<"constant_pool.splat", [
@@ -1332,6 +1368,7 @@
}];
let arguments = (ins
+ OptionalAttr<StrAttr>:$sym_visibility,
SymbolNameAttr:$sym_name,
ElementsAttr:$value,
OptionalAttr<SymbolRefAttr>:$runtime_buffer,
@@ -1339,10 +1376,21 @@
);
let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
$sym_name `=` $value
(`->` $runtime_buffer^ `[` $runtime_range `]`)?
attr-dict-with-keyword
}];
+
+ let builders = [
+ OpBuilder<(ins
+ "StringRef":$sym_name,
+ "::mlir::ElementsAttr":$value
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, value, SymbolRefAttr{},
+ IREE::Util::ByteRangeAttr{});
+ }]>
+ ];
}
def HAL_ConstantPoolLoadOp : HAL_PureOp<"constant_pool.load", [
@@ -1388,13 +1436,24 @@
}];
let arguments = (ins
+ OptionalAttr<StrAttr>:$sym_visibility,
SymbolNameAttr:$sym_name,
Util_AnySerializableAttr:$value
);
let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
$sym_name `=` $value attr-dict-with-keyword
}];
+
+ let builders = [
+ OpBuilder<(ins
+ "StringRef":$sym_name,
+ "::mlir::Attribute":$value
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, value);
+ }]>
+ ];
}
def HAL_ConstantStorageLookupOp :
@@ -1750,12 +1809,21 @@
}];
let arguments = (ins
- StrAttr:$sym_name
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name
// TODO(benvanik): entry point types for verification.
);
let regions = (region SizedRegion<1>:$body);
+ let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
+ $sym_name
+ attr-dict-with-keyword
+ ``
+ regions
+ }];
+
let skipDefaultBuilders = 1;
let builders = [
OpBuilder<(ins "StringRef":$name)>,
@@ -1799,7 +1867,8 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
HAL_OrdinalAttr:$ordinal,
FlatSymbolRefAttr:$interface,
OptionalAttr<HAL_WorkgroupSizeAttr>:$workgroup_size,
@@ -1816,7 +1885,7 @@
"::mlir::ArrayAttr":$workgroup_size,
"::mlir::IntegerAttr":$workgroup_local_memory
), [{
- build($_builder, $_state, sym_name, ordinal, interface,
+ build($_builder, $_state, nullptr, sym_name, ordinal, interface,
workgroup_size, workgroup_local_memory, 0);
}]>,
OpBuilder<(ins
@@ -1826,8 +1895,19 @@
"::mlir::ArrayAttr":$workgroup_size,
"::mlir::IntegerAttr":$workgroup_local_memory
), [{
- build($_builder, $_state, sym_name, ordinal, interface,
+ build($_builder, $_state, nullptr, sym_name, ordinal, interface,
workgroup_size, workgroup_local_memory, 0);
+ }]>,
+ OpBuilder<(ins
+ "::mlir::StringAttr":$sym_name,
+ "::mlir::IntegerAttr":$ordinal,
+ "::mlir::FlatSymbolRefAttr":$interface,
+ "::mlir::ArrayAttr":$workgroup_size,
+ "::mlir::IntegerAttr":$workgroup_local_memory,
+ "int":$workgroup_count_regionCount
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, ordinal, interface,
+ workgroup_size, workgroup_local_memory, workgroup_count_regionCount);
}]>
];
@@ -1860,12 +1940,22 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
HAL_ExecutableTargetAttr:$target
);
let regions = (region SizedRegion<1>:$body);
+ let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
+ $sym_name
+ `,` `target` `=` $target
+ attr-dict-with-keyword
+ ``
+ regions
+ }];
+
let skipDefaultBuilders = 1;
let builders = [
OpBuilder<(ins "StringRef":$name, "IREE::HAL::ExecutableTargetAttr":$target)>,
@@ -1901,13 +1991,20 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
StrAttr:$format,
HAL_ExecutableDataAttr:$data,
OptionalAttr<StrAttr>:$mime_type
// TODO(benvanik): add compatibility and versioning attributes.
);
+ let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
+ $sym_name
+ attr-dict-with-keyword
+ }];
+
let skipDefaultBuilders = 1;
let builders = [
OpBuilder<(ins
@@ -2025,12 +2122,21 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
OptionalAttr<IndexAttr>:$push_constants
);
let regions = (region SizedRegion<1>:$body);
+ let assemblyFormat = [{
+ custom<SymbolVisibility>($sym_visibility)
+ $sym_name
+ attr-dict-with-keyword
+ ``
+ regions
+ }];
+
let skipDefaultBuilders = 1;
let builders = [
OpBuilder<(ins
@@ -2086,13 +2192,38 @@
}];
let arguments = (ins
- StrAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ SymbolNameAttr:$sym_name,
IndexAttr:$set,
IndexAttr:$binding,
HAL_DescriptorTypeAttr:$type,
HAL_MemoryAccessBitfieldAttr:$access
);
+ // TODO(scotttodd): enable assemblyFormat after figuring out how to keep
+ // quotes around "type" and "access". The custom printer uses stringify*
+ // let assemblyFormat = [{
+ // custom<SymbolVisibility>($sym_visibility)
+ // $sym_name
+ // `,` `set` `` `=` `` $set
+ // `,` `binding` `` `=` `` $binding
+ // `,` `type` `` `=` `` $type
+ // `,` `access` `` `=` `` $access
+ // attr-dict-with-keyword
+ // }];
+
+ let builders = [
+ OpBuilder<(ins
+ "StringRef":$sym_name,
+ "::llvm::APInt":$set,
+ "::llvm::APInt":$binding,
+ "IREE::HAL::DescriptorType":$type,
+ "IREE::HAL::MemoryAccessBitfield":$access
+ ), [{
+ build($_builder, $_state, nullptr, sym_name, set, binding, type, access);
+ }]>
+ ];
+
let extraClassDeclaration = [{
/// Returns a hash for the descriptor, considering the set, binding,
/// type, and access.
diff --git a/iree/compiler/Dialect/HAL/IR/test/constant_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/constant_ops.mlir
index d931ddf..59c2a6e 100644
--- a/iree/compiler/Dialect/HAL/IR/test/constant_ops.mlir
+++ b/iree/compiler/Dialect/HAL/IR/test/constant_ops.mlir
@@ -1,15 +1,15 @@
// RUN: iree-opt -split-input-file %s | iree-opt -split-input-file | IreeFileCheck %s
-// CHECK-LABEL: hal.constant_pool @pool0
+// CHECK-LABEL: hal.constant_pool public @pool0
hal.constant_pool @pool0 attributes {
buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824,
min_buffer_offset_alignment = 32,
max_buffer_range = 134217728,
min_buffer_range_alignment = 4>
} {
- // CHECK-NEXT: hal.constant_pool.value @cst0 = dense<0.{{.+}}> : tensor<2x3xf32>
+ // CHECK-NEXT: hal.constant_pool.value public @cst0 = dense<0.{{.+}}> : tensor<2x3xf32>
hal.constant_pool.value @cst0 = dense<0.0> : tensor<2x3xf32>
- // CHECK-NEXT: hal.constant_pool.value @cst1 = dense<1.{{.+}}> : tensor<3x2xf32>
+ // CHECK-NEXT: hal.constant_pool.value public @cst1 = dense<1.{{.+}}> : tensor<3x2xf32>
hal.constant_pool.value @cst1 = dense<1.0> : tensor<3x2xf32>
}
@@ -24,26 +24,26 @@
// -----
-// CHECK-LABEL: hal.constant_pool @storage_allocated
+// CHECK-LABEL: hal.constant_pool public @storage_allocated
hal.constant_pool @storage_allocated attributes {
buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824,
min_buffer_offset_alignment = 32,
max_buffer_range = 134217728,
min_buffer_range_alignment = 4>
} {
- // CHECK-NEXT: hal.constant_pool.span @cst0 : tensor<2x3xf32> = @_storage[#util.byte_range<0, 1024>]
+ // CHECK-NEXT: hal.constant_pool.span public @cst0 : tensor<2x3xf32> = @_storage[#util.byte_range<0, 1024>]
hal.constant_pool.span @cst0 : tensor<2x3xf32> = @_storage[#util.byte_range<0, 1024>]
- // CHECK-NEXT: hal.constant_pool.span @cst1 : tensor<3x2xf32> = @_storage[#util.byte_range<1024, 1024>]
+ // CHECK-NEXT: hal.constant_pool.span public @cst1 : tensor<3x2xf32> = @_storage[#util.byte_range<1024, 1024>]
hal.constant_pool.span @cst1 : tensor<3x2xf32> = @_storage[#util.byte_range<1024, 1024>]
- // CHECK-NEXT: hal.constant_pool.splat @cst2 = dense<1.000000e+00> : tensor<1xf32>
+ // CHECK-NEXT: hal.constant_pool.splat public @cst2 = dense<1.000000e+00> : tensor<1xf32>
hal.constant_pool.splat @cst2 = dense<1.000000e+00> : tensor<1xf32>
- // CHECK-NEXT: hal.constant_storage @_storage = dense<1> : vector<768xi8>
+ // CHECK-NEXT: hal.constant_storage public @_storage = dense<1> : vector<768xi8>
hal.constant_storage @_storage = dense<1> : vector<768xi8>
}
// -----
-// CHECK-LABEL: hal.constant_pool @pool
+// CHECK-LABEL: hal.constant_pool public @pool
// CHECK-SAME: buffer_constraints = #hal.buffer_constraints
hal.constant_pool @pool attributes {
buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824,
@@ -51,17 +51,17 @@
max_buffer_range = 134217728,
min_buffer_range_alignment = 4>
} {
- // CHECK-NEXT: hal.constant_pool.span @cst0 : tensor<4xf32> = @_storage0[#util.byte_range<0, 16>] -> @pool_storage0_buffer[#util.byte_range<0, 16>]
+ // CHECK-NEXT: hal.constant_pool.span public @cst0 : tensor<4xf32> = @_storage0[#util.byte_range<0, 16>] -> @pool_storage0_buffer[#util.byte_range<0, 16>]
hal.constant_pool.span @cst0 : tensor<4xf32> = @_storage0[#util.byte_range<0, 16>] -> @pool_storage0_buffer[#util.byte_range<0, 16>]
- // CHECK-NEXT: hal.constant_pool.span @cst1 : tensor<3xi8> = @_storage1[#util.byte_range<0, 3>] -> @pool_storage1_buffer[#util.byte_range<0, 3>]
+ // CHECK-NEXT: hal.constant_pool.span public @cst1 : tensor<3xi8> = @_storage1[#util.byte_range<0, 3>] -> @pool_storage1_buffer[#util.byte_range<0, 3>]
hal.constant_pool.span @cst1 : tensor<3xi8> = @_storage1[#util.byte_range<0, 3>] -> @pool_storage1_buffer[#util.byte_range<0, 3>]
- // CHECK-NEXT: hal.constant_pool.splat @cst2 = dense<1.000000e+00> : tensor<1xf32> -> @pool_splats[#util.byte_range<0, 4>]
+ // CHECK-NEXT: hal.constant_pool.splat public @cst2 = dense<1.000000e+00> : tensor<1xf32> -> @pool_splats[#util.byte_range<0, 4>]
hal.constant_pool.splat @cst2 = dense<1.000000e+00> : tensor<1xf32> -> @pool_splats[#util.byte_range<0, 4>]
- // CHECK-NEXT: hal.constant_pool.splat @cst3 = dense<1234567890> : tensor<8xi32> -> @pool_splats[#util.byte_range<32, 32>]
+ // CHECK-NEXT: hal.constant_pool.splat public @cst3 = dense<1234567890> : tensor<8xi32> -> @pool_splats[#util.byte_range<32, 32>]
hal.constant_pool.splat @cst3 = dense<1234567890> : tensor<8xi32> -> @pool_splats[#util.byte_range<32, 32>]
- // CHECK-NEXT: hal.constant_storage @_storage0 = dense<[102, 102, 6, 64, -51, -52, 76, 64, -102, -103, -119, 64, -51, -52, -84, 64]> : vector<16xi8>
+ // CHECK-NEXT: hal.constant_storage public @_storage0 = dense<[102, 102, 6, 64, -51, -52, 76, 64, -102, -103, -119, 64, -51, -52, -84, 64]> : vector<16xi8>
hal.constant_storage @_storage0 = dense<[102, 102, 6, 64, -51, -52, 76, 64, -102, -103, -119, 64, -51, -52, -84, 64]> : vector<16xi8>
- // CHECK-NEXT: hal.constant_storage @_storage1 = dense<[6, 7, 8, 0]> : vector<4xi8>
+ // CHECK-NEXT: hal.constant_storage public @_storage1 = dense<[6, 7, 8, 0]> : vector<4xi8>
hal.constant_storage @_storage1 = dense<[6, 7, 8, 0]> : vector<4xi8>
}
diff --git a/iree/compiler/Dialect/HAL/IR/test/executable_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/executable_ops.mlir
index 8ebdb17..d7e70a6 100644
--- a/iree/compiler/Dialect/HAL/IR/test/executable_ops.mlir
+++ b/iree/compiler/Dialect/HAL/IR/test/executable_ops.mlir
@@ -4,9 +4,9 @@
// CHECK-LABEL: @ex
hal.executable @ex {
- // CHECK: hal.executable.variant @backend, target = #executable_target_format
+ // CHECK: hal.executable.variant public @backend, target = #executable_target_format
hal.executable.variant @backend, target = #executable_target_format {
- // CHECK-DAG: hal.executable.entry_point @entry0 attributes {
+ // CHECK-DAG: hal.executable.entry_point public @entry0 attributes {
// CHECK-SAME: interface = @interface
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
@@ -16,11 +16,11 @@
workgroup_size = [4 : index, 1 : index, 1 : index]
}
}
- // CHECK-DAG: hal.interface @interface
+ // CHECK-DAG: hal.interface public @interface
hal.interface @interface {
- // CHECK-NEXT: hal.interface.binding @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
+ // CHECK-NEXT: hal.interface.binding public @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
- // CHECK-NEXT: hal.interface.binding @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
+ // CHECK-NEXT: hal.interface.binding public @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
// CHECK: hal.executable.binary
@@ -38,9 +38,9 @@
// CHECK-LABEL: @ex_with_workgroup_count_region
hal.executable @ex_with_workgroup_count_region {
- // CHECK: hal.executable.variant @backend, target = #executable_target_format
+ // CHECK: hal.executable.variant public @backend, target = #executable_target_format
hal.executable.variant @backend, target = #executable_target_format {
- // CHECK-DAG: hal.executable.entry_point @entry0 attributes {
+ // CHECK-DAG: hal.executable.entry_point public @entry0 attributes {
// CHECK-SAME: interface = @interface
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
@@ -53,11 +53,11 @@
hal.return %arg0, %arg1, %arg2 : index, index, index
}
}
- // CHECK-DAG: hal.interface @interface
+ // CHECK-DAG: hal.interface public @interface
hal.interface @interface {
- // CHECK-NEXT: hal.interface.binding @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
+ // CHECK-NEXT: hal.interface.binding public @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b0, set=0, binding=0, type="StorageBuffer", access="Read"
- // CHECK-NEXT: hal.interface.binding @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
+ // CHECK-NEXT: hal.interface.binding public @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
hal.interface.binding @s0b1, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
// CHECK: hal.executable.binary
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
index fff3c6f..3d2a07b 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
@@ -40,6 +40,6 @@
// PTX: add.rn.f32
// PTX: sqrt
-// CHECK: hal.executable.binary @cuda_nvptx_fb attributes {
+// CHECK: hal.executable.binary public @cuda_nvptx_fb attributes {
// CHECK-SAME: data = dense
// CHECK-SAME: format = "cuda-nvptx-fb"
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
index e6001f6..38d9e3c 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
@@ -34,6 +34,6 @@
}
-// CHECK: hal.executable.binary @embedded_elf_x86_64
+// CHECK: hal.executable.binary public @embedded_elf_x86_64
// CHECK-SAME: data = dense
// CHECK-SAME: format = "embedded-elf-x86_64"
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
index 6628de1..fd10ea9 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
@@ -36,6 +36,6 @@
}
-// CHECK: hal.executable.binary @metal_msl_fb attributes {
+// CHECK: hal.executable.binary public @metal_msl_fb attributes {
// CHECK-SAME: data = dense
// CHECK-SAME: format = "metal-msl-fb"
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
index d0c2ae4..a8fc4be 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
@@ -35,6 +35,6 @@
}
-// CHECK: hal.executable.binary @rocm_hsaco_fb attributes {
+// CHECK: hal.executable.binary public @rocm_hsaco_fb attributes {
// CHECK-SAME: data = dense
// CHECK-SAME: format = "rocm-hsaco-fb"
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir
index 7921db2..846ef60 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir
@@ -2,7 +2,7 @@
#vmvx_target = #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
-hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_0 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -20,7 +20,7 @@
}
}
}
-hal.executable @dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_1 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -38,7 +38,7 @@
}
}
}
-hal.executable @dispatch_2 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_2 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -69,25 +69,25 @@
// All executables (including their interfaces and entry points) should be
// linked together into a single executable.
-// CHECK-NOT: hal.executable @dispatch_0
-// CHECK-NOT: hal.executable @dispatch_1
-// CHECK-NOT: hal.executable @dispatch_2
-// CHECK: hal.executable @vmvx_linked attributes {sym_visibility = "private"} {
-// CHECK-NEXT: hal.interface @io_0 {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-NOT: hal.executable private @dispatch_0
+// CHECK-NOT: hal.executable private @dispatch_1
+// CHECK-NOT: hal.executable private @dispatch_2
+// CHECK: hal.executable private @vmvx_linked {
+// CHECK-NEXT: hal.interface public @io_0 {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.interface @io_1 {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg2, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-NEXT: hal.interface public @io_1 {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg2, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
-// CHECK-NEXT: hal.executable.entry_point @dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
-// CHECK-NEXT: hal.executable.entry_point @dispatch_1 attributes {interface = @io_0, ordinal = 1 : index}
-// CHECK-NEXT: hal.executable.entry_point @dispatch_2 attributes {interface = @io_1, ordinal = 2 : index}
+// CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+// CHECK-NEXT: hal.executable.entry_point public @dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
+// CHECK-NEXT: hal.executable.entry_point public @dispatch_1 attributes {interface = @io_0, ordinal = 1 : index}
+// CHECK-NEXT: hal.executable.entry_point public @dispatch_2 attributes {interface = @io_1, ordinal = 2 : index}
// CHECK-NEXT: module {
// CHECK-NEXT: vm.module public @linked_module {
// CHECK-NEXT: vm.func @dispatch_0() {
@@ -119,7 +119,7 @@
#cuda_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#vmvx_target = #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
-hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_0 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -141,7 +141,7 @@
}
}
}
-hal.executable @dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_1 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -183,19 +183,19 @@
// VMVX target should be pulled out from both executables leaving the originals
// untouched.
-// CHECK: hal.executable @vmvx_linked attributes {sym_visibility = "private"} {
-// CHECK-NEXT: hal.interface @io_0 {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK: hal.executable private @vmvx_linked {
+// CHECK-NEXT: hal.interface public @io_0 {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.interface @io_1 {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-NEXT: hal.interface public @io_1 {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
-// CHECK-NEXT: hal.executable.entry_point @dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
-// CHECK-NEXT: hal.executable.entry_point @dispatch_1 attributes {interface = @io_1, ordinal = 1 : index}
+// CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+// CHECK-NEXT: hal.executable.entry_point public @dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
+// CHECK-NEXT: hal.executable.entry_point public @dispatch_1 attributes {interface = @io_1, ordinal = 1 : index}
// CHECK-NEXT: module {
// CHECK-NEXT: vm.module public @linked_module {
// CHECK-NEXT: vm.func @dispatch_0() {
@@ -212,12 +212,12 @@
// CHECK-NEXT: }
//
// @dispatch_0/1 should remain, with just @cuda
-// CHECK: hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
-// CHECK: hal.interface @io
-// CHECK: hal.executable.variant @cuda, target = #executable_target_cuda
-// CHECK: hal.executable @dispatch_1 attributes {sym_visibility = "private"} {
-// CHECK: hal.interface @io
-// CHECK: hal.executable.variant @cuda, target = #executable_target_cuda
+// CHECK: hal.executable private @dispatch_0 {
+// CHECK: hal.interface public @io
+// CHECK: hal.executable.variant public @cuda, target = #executable_target_cuda
+// CHECK: hal.executable private @dispatch_1 {
+// CHECK: hal.interface public @io
+// CHECK: hal.executable.variant public @cuda, target = #executable_target_cuda
//
// CHECK: func @other_targets() {
// CHECK: hal.device.switch<%device : !hal.device>
@@ -241,7 +241,7 @@
#vmvx_target = #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
module {
- hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
+ hal.executable private @dispatch_0 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -254,7 +254,7 @@
}
}
}
- hal.executable @dispatch_1 attributes {sym_visibility = "private"} {
+ hal.executable private @dispatch_1 {
hal.interface @io attributes {push_constants = 2 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -267,7 +267,7 @@
}
}
}
- hal.executable @dispatch_2 attributes {sym_visibility = "private"} {
+ hal.executable private @dispatch_2 {
hal.interface @io attributes {push_constants = 2 : index} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -283,26 +283,26 @@
}
// Interfaces with different numbers of push constants should remain separate.
-// CHECK-NOT: hal.executable @dispatch_0
-// CHECK-NOT: hal.executable @dispatch_1
-// CHECK-NOT: hal.executable @dispatch_2
-// CHECK: hal.executable @vmvx_linked attributes {sym_visibility = "private"} {
-// CHECK-NEXT: hal.interface @io_0 {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-NOT: hal.executable private @dispatch_0
+// CHECK-NOT: hal.executable private @dispatch_1
+// CHECK-NOT: hal.executable private @dispatch_2
+// CHECK: hal.executable private @vmvx_linked {
+// CHECK-NEXT: hal.interface public @io_0 {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.interface @io_1 attributes {push_constants = 2 : index} {
-// CHECK-NEXT: hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-NEXT: hal.interface public @io_1 attributes {push_constants = 2 : index} {
+// CHECK-NEXT: hal.interface.binding public @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
// -----
#vmvx_target = #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
-hal.executable @dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_0 {
hal.interface @io {}
hal.executable.variant @vmvx, target = #vmvx_target {
hal.executable.entry_point @dispatch_0 attributes {interface = @io, ordinal = 0 : index}
@@ -325,7 +325,7 @@
}
}
}
-hal.executable @dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @dispatch_1 {
hal.interface @io {}
hal.executable.variant @vmvx, target = #vmvx_target {
hal.executable.entry_point @dispatch_1 attributes {interface = @io, ordinal = 0 : index}
@@ -356,10 +356,10 @@
// resolve conflicts.
// References to renamed symbols should be updated.
//
-// CHECK-NOT: hal.executable @dispatch_0
-// CHECK-NOT: hal.executable @dispatch_1
-// CHECK: hal.executable @vmvx_linked attributes {sym_visibility = "private"} {
-// CHECK: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+// CHECK-NOT: hal.executable private @dispatch_0
+// CHECK-NOT: hal.executable private @dispatch_1
+// CHECK: hal.executable private @vmvx_linked {
+// CHECK: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
// CHECK: module {
// CHECK-NEXT: vm.module public @linked_module {
// CHECK-NEXT: vm.rodata public @rodata_a dense<0> : tensor<1xi32>
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
index 51fb789..65da95e 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
@@ -34,14 +34,14 @@
}
-// CHECK-LABEL: hal.executable @add_dispatch_0
-// CHECK-NEXT: hal.interface @io {
-// CHECK-NEXT: hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @add_dispatch_0
+// CHECK-NEXT: hal.interface public @io {
+// CHECK-NEXT: hal.interface.binding public @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
-// CHECK-NEXT: hal.executable.entry_point @entry attributes {
+// CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+// CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @io,
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir
index 7ff3f3b..b94284f 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir
@@ -5,7 +5,7 @@
#executable_target_vulkan_spirv_fb = #hal.executable.target<"vulkan", "vulkan-spirv-fb">
-hal.executable @call_dispatch_0 attributes {sym_visibility = "private"} {
+hal.executable private @call_dispatch_0 {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
@@ -20,14 +20,14 @@
spv.EntryPoint "GLCompute" @call_dispatch_0
spv.ExecutionMode @call_dispatch_0 "LocalSize", 32, 1, 1
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
}
}
}
-hal.executable @call_dispatch_1 attributes {sym_visibility = "private"} {
+hal.executable private @call_dispatch_1 {
hal.interface @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -43,7 +43,7 @@
spv.EntryPoint "GLCompute" @call_dispatch_1
spv.ExecutionMode @call_dispatch_1 "LocalSize", 4, 4, 1
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -51,7 +51,7 @@
}
}
}
-hal.executable @call_dispatch_2 attributes {sym_visibility = "private"} {
+hal.executable private @call_dispatch_2 {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
@@ -66,14 +66,14 @@
spv.EntryPoint "GLCompute" @call_dispatch_2
spv.ExecutionMode @call_dispatch_2 "LocalSize", 32, 1, 1
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
}
}
}
}
-hal.executable @call_dispatch_3 attributes {sym_visibility = "private"} {
+hal.executable private @call_dispatch_3 {
hal.interface @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -95,7 +95,7 @@
spv.EntryPoint "GLCompute" @call_dispatch_3
spv.ExecutionMode @call_dispatch_3 "LocalSize", 8, 2, 2
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -103,7 +103,7 @@
}
}
}
-hal.executable @call_dispatch_4 attributes {sym_visibility = "private"} {
+hal.executable private @call_dispatch_4 {
hal.interface @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
@@ -119,7 +119,7 @@
spv.EntryPoint "GLCompute" @call_dispatch_4
spv.ExecutionMode @call_dispatch_4 "LocalSize", 2, 8, 1
}
- hal.interface @io attributes {sym_visibility = "private"} {
+ hal.interface private @io {
hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
@@ -130,16 +130,16 @@
// Two groups should be created, according to their interfaces.
-// CHECK: hal.executable @linking_linked_vulkan_0 {
-// CHECK-NEXT: hal.interface @io_0 {
-// CHECK-NEXT: hal.interface.binding @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+// CHECK: hal.executable private @linking_linked_vulkan_0 {
+// CHECK-NEXT: hal.interface public @io_0 {
+// CHECK-NEXT: hal.interface.binding public @s0b0_ro_constant, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @s0b1_ro_external, set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @s0b2_xw_external, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb {
-// CHECK-NEXT: hal.executable.entry_point @call_dispatch_1 attributes {interface = @io_0, ordinal = 0 : index}
-// CHECK-NEXT: hal.executable.entry_point @call_dispatch_3 attributes {interface = @io_0, ordinal = 1 : index}
-// CHECK-NEXT: hal.executable.entry_point @call_dispatch_4 attributes {interface = @io_0, ordinal = 2 : index}
+// CHECK-NEXT: hal.executable.variant public @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb {
+// CHECK-NEXT: hal.executable.entry_point public @call_dispatch_1 attributes {interface = @io_0, ordinal = 0 : index}
+// CHECK-NEXT: hal.executable.entry_point public @call_dispatch_3 attributes {interface = @io_0, ordinal = 1 : index}
+// CHECK-NEXT: hal.executable.entry_point public @call_dispatch_4 attributes {interface = @io_0, ordinal = 2 : index}
// CHECK-NEXT: module {
// CHECK-NEXT: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
// CHECK-NEXT: spv.func @call_dispatch_1() "None" {
@@ -162,14 +162,14 @@
// CHECK-NEXT: }
// CHECK-NEXT: }
-// CHECK: hal.executable @linking_linked_vulkan {
-// CHECK-NEXT: hal.interface @io_0 {
-// CHECK-NEXT: hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
+// CHECK: hal.executable private @linking_linked_vulkan {
+// CHECK-NEXT: hal.interface public @io_0 {
+// CHECK-NEXT: hal.interface.binding public @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @s0b1_rw_external, set=0, binding=1, type="StorageBuffer", access="Read|Write"
// CHECK-NEXT: }
-// CHECK-NEXT: hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb {
-// CHECK-NEXT: hal.executable.entry_point @call_dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
-// CHECK-NEXT: hal.executable.entry_point @call_dispatch_2 attributes {interface = @io_0, ordinal = 1 : index}
+// CHECK-NEXT: hal.executable.variant public @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb {
+// CHECK-NEXT: hal.executable.entry_point public @call_dispatch_0 attributes {interface = @io_0, ordinal = 0 : index}
+// CHECK-NEXT: hal.executable.entry_point public @call_dispatch_2 attributes {interface = @io_0, ordinal = 1 : index}
// CHECK-NEXT: module {
// CHECK-NEXT: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
// CHECK-NEXT: spv.func @call_dispatch_0() "None" {
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
index 6590136..a02179b 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
@@ -36,6 +36,6 @@
}
-// CHECK: hal.executable.binary @vulkan_spirv_fb attributes
+// CHECK: hal.executable.binary public @vulkan_spirv_fb attributes
// CHECK-SAME: data = dense
// CHECK-SAME: format = "vulkan-spirv-fb"
diff --git a/iree/compiler/Dialect/HAL/Transforms/PackConstantPoolStorage.cpp b/iree/compiler/Dialect/HAL/Transforms/PackConstantPoolStorage.cpp
index e83234c..b1d4e4b 100644
--- a/iree/compiler/Dialect/HAL/Transforms/PackConstantPoolStorage.cpp
+++ b/iree/compiler/Dialect/HAL/Transforms/PackConstantPoolStorage.cpp
@@ -78,8 +78,7 @@
builder.setInsertionPointAfter(splatValueOp);
auto splatOp = builder.create<ConstantPoolSplatOp>(
splatValueOp.getLoc(), splatValueOp.getName(),
- splatValueOp.value().cast<SplatElementsAttr>(), SymbolRefAttr{},
- IREE::Util::ByteRangeAttr{});
+ splatValueOp.value().cast<SplatElementsAttr>());
splatOp.setNested();
splatValueOp.erase();
}
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir b/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
index 7a0566e..9fe57a8 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
@@ -10,13 +10,13 @@
hal.device.targets = [#device_target_cpu, #device_target_gpu]
} {
-// CHECK: hal.constant_pool @_const_pool attributes
+// CHECK: hal.constant_pool private @_const_pool attributes
// CHECK-SAME: buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824, min_buffer_offset_alignment = 256, max_buffer_range = 134217728, min_buffer_range_alignment = 16>
-// CHECK-NEXT: hal.constant_pool.value @cst0 = dense<1.000000e+00> : tensor<1xf32>
+// CHECK-NEXT: hal.constant_pool.value nested @cst0 = dense<1.000000e+00> : tensor<1xf32>
util.global private @cst0 = dense<1.000000e+00> : tensor<1xf32>
-// CHECK-NEXT: hal.constant_pool.value @cst1 = dense<[2.100000e+00, 3.200000e+00, 4.300000e+00, 5.400000e+00]> : tensor<4xf32>
+// CHECK-NEXT: hal.constant_pool.value nested @cst1 = dense<[2.100000e+00, 3.200000e+00, 4.300000e+00, 5.400000e+00]> : tensor<4xf32>
util.global private @cst1 = dense<[2.1, 3.2, 4.3, 5.4]> : tensor<4xf32>
-// CHECK-NEXT: hal.constant_pool.value @cst2 = dense<[6, 7, 8]> : tensor<3xi8>
+// CHECK-NEXT: hal.constant_pool.value nested @cst2 = dense<[6, 7, 8]> : tensor<3xi8>
util.global private @cst2 = dense<[6, 7, 8]> : tensor<3xi8>
// CHECK-LABEL: func @immutable_variables
@@ -34,8 +34,8 @@
// -----
-// CHECK: hal.constant_pool @_const_pool_init
-// CHECK-NEXT: hal.constant_pool.value @variable_0 = dense<3.000000e+00> : tensor<128xf32>
+// CHECK: hal.constant_pool private @_const_pool_init
+// CHECK-NEXT: hal.constant_pool.value nested @variable_0 = dense<3.000000e+00> : tensor<128xf32>
// CHECK: util.global private mutable @variable_0
util.global private mutable @variable_0 = dense<3.0> : tensor<128xf32>
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
index d9cb28c..627341d 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -iree-hal-materialize-constant-pool-buffers %s | IreeFileCheck %s
-// CHECK-LABEL: hal.constant_pool @dense_variable_init
+// CHECK-LABEL: hal.constant_pool public @dense_variable_init
hal.constant_pool @dense_variable_init attributes {buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824, min_buffer_offset_alignment = 32, max_buffer_range = 134217728, min_buffer_range_alignment = 4>} {
// CHECK-NEXT: @cst0 {{.+}} -> @dense_variable_init_storage_buffer[#util.byte_range<0, 512>]
hal.constant_pool.span @cst0 : tensor<128xf32> = @_storage[#util.byte_range<0, 512>]
@@ -18,7 +18,7 @@
// -----
-// CHECK-LABEL: hal.constant_pool @splat_variable_init
+// CHECK-LABEL: hal.constant_pool public @splat_variable_init
hal.constant_pool @splat_variable_init attributes {buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824, min_buffer_offset_alignment = 32, max_buffer_range = 134217728, min_buffer_range_alignment = 4>} {
// CHECK-NEXT: @cst0 {{.+}} -> @splat_variable_init_splats[#util.byte_range<0, 4>]
hal.constant_pool.splat @cst0 = dense<1.000000e+00> : tensor<1xf32>
@@ -40,7 +40,7 @@
// -----
-// CHECK-LABEL: hal.constant_pool @pool
+// CHECK-LABEL: hal.constant_pool public @pool
hal.constant_pool @pool attributes {buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824, min_buffer_offset_alignment = 32, max_buffer_range = 134217728, min_buffer_range_alignment = 4>} {
// CHECK-NEXT: @cst0 {{.+}} -> @pool_storage0_buffer[#util.byte_range<0, 16>]
hal.constant_pool.span @cst0 : tensor<4xf32> = @_storage0[#util.byte_range<0, 16>]
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
index 89b8a75..6f04508 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
@@ -4,14 +4,14 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @static_tiled_dispatch
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @static_tiled_dispatch
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @static_tiled_dispatch {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry attributes {
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @[[IO]],
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
@@ -64,14 +64,14 @@
hal.device.targets = [#device_target_vmvx, #device_target_cuda]
} {
-// CHECK-LABEL: hal.executable @multi_target_ex
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @multi_target_ex
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @multi_target_ex {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry attributes {
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @[[IO]],
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
@@ -95,8 +95,8 @@
return
}
}
- // CHECK: hal.executable.variant @cuda_nvptx_fb, target = #executable_target_cuda_nvptx_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry attributes {
+ // CHECK: hal.executable.variant public @cuda_nvptx_fb, target = #executable_target_cuda_nvptx_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @[[IO]],
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
@@ -123,14 +123,14 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @dynamic_tiled_dispatch
-// CHECK-NEXT: hal.interface @[[IO:.+]] attributes {push_constants = 4 : index} {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @dynamic_tiled_dispatch
+// CHECK-NEXT: hal.interface public @[[IO:.+]] attributes {push_constants = 4 : index} {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @dynamic_tiled_dispatch {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry attributes {
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @[[IO]],
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
@@ -199,14 +199,14 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @workgroup_infos
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @workgroup_infos
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @workgroup_infos {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry attributes {
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry attributes {
// CHECK-SAME: interface = @[[IO]],
// CHECK-SAME: ordinal = 0 : index
// CHECK-SAME: }
@@ -244,14 +244,14 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @static_tied_result
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Read|Write"
+// CHECK-LABEL: hal.executable public @static_tied_result
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Read|Write"
// CHECK-NEXT: }
flow.executable @static_tied_result {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry
flow.dispatch.entry @entry attributes {
workgroup_rank = 2 : index
}
@@ -295,17 +295,17 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @constant_dispatch
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B2:.+]], set=0, binding=2, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B3:.+]], set=0, binding=3, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B4:.+]], set=0, binding=4, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @constant_dispatch
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B2:.+]], set=0, binding=2, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B3:.+]], set=0, binding=3, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B4:.+]], set=0, binding=4, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @constant_dispatch {
- // CHECK-NEXT: hal.executable.variant @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
- // CHECK-NEXT: hal.executable.entry_point @entry
+ // CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = #executable_target_vmvx_bytecode_fb {
+ // CHECK-NEXT: hal.executable.entry_point public @entry
flow.dispatch.entry @entry attributes {
workgroup_rank = 2 : index
}
@@ -390,10 +390,10 @@
executable_targets = [#hal.executable.target<"vmvx", "vmvx-bytecode-fb">]
}>]} {
-// CHECK-LABEL: hal.executable @unsued_arg
-// CHECK-NEXT: hal.interface @[[IO:.+]] {
-// CHECK-NEXT: hal.interface.binding @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
-// CHECK-NEXT: hal.interface.binding @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
+// CHECK-LABEL: hal.executable public @unsued_arg
+// CHECK-NEXT: hal.interface public @[[IO:.+]] {
+// CHECK-NEXT: hal.interface.binding public @[[S0B0:.+]], set=0, binding=0, type="StorageBuffer", access="Read"
+// CHECK-NEXT: hal.interface.binding public @[[S0B1:.+]], set=0, binding=1, type="StorageBuffer", access="Write|Discard"
// CHECK-NEXT: }
flow.executable @unsued_arg {
flow.dispatch.entry @entry attributes {workgroup_rank = 2 : index}
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/pack_constant_pool_storage.mlir b/iree/compiler/Dialect/HAL/Transforms/test/pack_constant_pool_storage.mlir
index 62d2924..368b911 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/pack_constant_pool_storage.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/pack_constant_pool_storage.mlir
@@ -1,20 +1,20 @@
// RUN: iree-opt -split-input-file -iree-hal-pack-constant-pool-storage -mlir-print-local-scope %s | IreeFileCheck %s
-// CHECK-LABEL: hal.constant_pool @pool
+// CHECK-LABEL: hal.constant_pool public @pool
hal.constant_pool @pool attributes {
buffer_constraints = #hal.buffer_constraints<max_allocation_size = 1073741824,
min_buffer_offset_alignment = 8,
max_buffer_range = 134217728,
min_buffer_range_alignment = 4>
} {
- // CHECK-DAG: hal.constant_pool.span @cst0 : tensor<3xi8> = @_storage[#util.byte_range<0, 3>]
+ // CHECK-DAG: hal.constant_pool.span nested @cst0 : tensor<3xi8> = @_storage[#util.byte_range<0, 3>]
hal.constant_pool.value @cst0 = dense<[6, 7, 8]> : tensor<3xi8>
- // CHECK-DAG: hal.constant_pool.splat @cst1 = dense<1.000000e+00> : tensor<1xf32>
+ // CHECK-DAG: hal.constant_pool.splat nested @cst1 = dense<1.000000e+00> : tensor<1xf32>
hal.constant_pool.value @cst1 = dense<1.000000e+00> : tensor<1xf32>
- // CHECK-DAG: hal.constant_pool.span @cst2 : tensor<4xf32> = @_storage[#util.byte_range<8, 16>]
+ // CHECK-DAG: hal.constant_pool.span nested @cst2 : tensor<4xf32> = @_storage[#util.byte_range<8, 16>]
hal.constant_pool.value @cst2 = dense<[2.1, 3.2, 4.3, 5.4]> : tensor<4xf32>
- // CHECK: hal.constant_storage @_storage = #util.composite<24xi8, [
+ // CHECK: hal.constant_storage nested @_storage = #util.composite<24xi8, [
// CHECK-NEXT: dense<[6, 7, 8]> : tensor<3xi8>,
// CHECK-NEXT: dense<0> : vector<5xi8>,
// CHECK-NEXT: dense<[2.100000e+00, 3.200000e+00, 4.300000e+00, 5.400000e+00]> : tensor<4xf32>,
@@ -23,22 +23,22 @@
// -----
-// CHECK-LABEL: hal.constant_pool @multi_storage
+// CHECK-LABEL: hal.constant_pool public @multi_storage
hal.constant_pool @multi_storage attributes {
buffer_constraints = #hal.buffer_constraints<max_allocation_size = 18,
min_buffer_offset_alignment = 1,
max_buffer_range = 134217728,
min_buffer_range_alignment = 1>
} {
- // CHECK-DAG: hal.constant_pool.span @cst0 : tensor<4xf32> = @_storage[#util.byte_range<0, 16>]
+ // CHECK-DAG: hal.constant_pool.span nested @cst0 : tensor<4xf32> = @_storage[#util.byte_range<0, 16>]
hal.constant_pool.value @cst0 = dense<[2.1, 3.2, 4.3, 5.4]> : tensor<4xf32>
- // CHECK-DAG: hal.constant_pool.span @cst1 : tensor<3xi8> = @_storage_0[#util.byte_range<0, 3>]
+ // CHECK-DAG: hal.constant_pool.span nested @cst1 : tensor<3xi8> = @_storage_0[#util.byte_range<0, 3>]
hal.constant_pool.value @cst1 = dense<[6, 7, 8]> : tensor<3xi8>
- // CHECK-NEXT: hal.constant_storage @_storage = #util.composite<16xi8, [
+ // CHECK-NEXT: hal.constant_storage nested @_storage = #util.composite<16xi8, [
// CHECK-NEXT: dense<[2.100000e+00, 3.200000e+00, 4.300000e+00, 5.400000e+00]> : tensor<4xf32>,
// CHECK-NEXT: ]>
- // CHECK-NEXT: hal.constant_storage @_storage_0 = #util.composite<3xi8, [
+ // CHECK-NEXT: hal.constant_storage nested @_storage_0 = #util.composite<3xi8, [
// CHECK-NEXT: dense<[6, 7, 8]> : tensor<3xi8>,
// CHECK-NEXT: ]>
}
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir
index f7a67df..51ecc83 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir
+++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -iree-vmvx-conversion -canonicalize %s | IreeFileCheck %s
-hal.interface @io attributes {sym_visibility = "private"} {
+hal.interface private @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @s0b1_xw_external, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
