[spirv] NFC: replace PlaceholderOp with InterfaceSubspan op (#5479)
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
index 2ba624f..b7bae3a 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
@@ -20,12 +20,10 @@
// {max_compute_workgroup_invocations = 128 : i32,
// max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
// func @parallel_4D() {
-// %arg0 = iree.placeholder for "interace buffer"
-// {binding = @io::@arg0, operand_result_index = 4 : i32} : memref<?x?x?x?xf32>
-// %arg1 = iree.placeholder for "interace buffer"
-// {binding = @io::@arg1, operand_result_index = 9 : i32} : memref<?x?x?x?xf32>
-// %arg2 = iree.placeholder for "interace buffer"
-// {binding = @io::@ret0, operand_result_index = 10 : i32} : memref<?x?x?x?xf32>
+// %c0 = constant 0 : index
+// %arg0 = hal.interface.subspan @io::@arg0[%c0] : memref<?x?x?x?xf32>
+// %arg1 = hal.interface.subspan @io::@arg1[%c0] : memref<?x?x?x?xf32>
+// %arg2 = hal.interface.subspan @io::@ret0[%c0] : memref<?x?x?x?xf32>
// linalg.generic {
// indexing_maps = [#map0, #map0, #map0],
// iterator_types = ["parallel", "parallel", "parallel", "parallel"]}
@@ -99,12 +97,10 @@
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
func @parallel_4D_static() {
- %arg0 = iree.placeholder for "interace buffer"
- {binding = @io::@arg0, operand_result_index = 0 : i32} : memref<3x4x5x6xf32>
- %arg1 = iree.placeholder for "interace buffer"
- {binding = @io::@arg1, operand_result_index = 1 : i32} : memref<3x4x5x6xf32>
- %arg2 = iree.placeholder for "interace buffer"
- {binding = @io::@ret0, operand_result_index = 2 : i32} : memref<3x4x5x6xf32>
+ %c0 = constant 0 : index
+ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<3x4x5x6xf32>
+ %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x4x5x6xf32>
+ %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<3x4x5x6xf32>
linalg.generic {
indexing_maps = [#map0, #map0, #map0],
iterator_types = ["parallel", "parallel", "parallel", "parallel"]}
@@ -178,12 +174,10 @@
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
func @scalar_add() attributes {hal.num_workgroups_fn = @scalar_add__num_workgroups__} {
- %arg0 = iree.placeholder for "interace buffer"
- {binding = @io::@arg0, operand_result_index = 0 : i32} : memref<f32>
- %arg1 = iree.placeholder for "interace buffer"
- {binding = @io::@arg1, operand_result_index = 1 : i32} : memref<f32>
- %arg2 = iree.placeholder for "interace buffer"
- {binding = @io::@ret0, operand_result_index = 2 : i32} : memref<f32>
+ %c0 = constant 0 : index
+ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<f32>
+ %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<f32>
+ %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<f32>
linalg.generic #trait
ins(%arg0, %arg1 : memref<f32>, memref<f32>)
outs(%arg2 : memref<f32>) {
@@ -227,12 +221,10 @@
!flow.dispatch.tensor<writeonly:40xf32>) -> ()}
module {
func @reduce_sum() {
- %arg0 = iree.placeholder for "interace buffer"
- {binding = @io::@arg0, operand_result_index = 0 : i32} : memref<40x50x75xf32>
- %arg1 = iree.placeholder for "interace buffer"
- {binding = @io::@arg1, operand_result_index = 1 : i32} : memref<f32>
- %arg2 = iree.placeholder for "interace buffer"
- {binding = @io::@ret0, operand_result_index = 2 : i32} : memref<40xf32>
+ %c0 = constant 0 : index
+ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<40x50x75xf32>
+ %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<f32>
+ %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<40xf32>
linalg.indexed_generic {
indexing_maps = [affine_map<(d0, d1, d2) -> (d0, d1, d2)>,
affine_map<(d0, d1, d2) -> ()>,
@@ -242,9 +234,9 @@
outs(%arg2 : memref<40xf32>) {
^bb0(%arg3: index, %arg4: index, %arg5: index,
%arg6: f32, %arg7: f32, %arg8: f32): // no predecessors
- %c0 = constant 0 : index
- %0 = cmpi eq, %arg5, %c0 : index
- %1 = cmpi eq, %arg4, %c0 : index
+ %zero = constant 0 : index
+ %0 = cmpi eq, %arg5, %zero : index
+ %1 = cmpi eq, %arg4, %zero : index
%2 = and %0, %1 : i1
%3 = select %2, %arg7, %arg8 : f32
%4 = addf %arg6, %3 : f32
@@ -304,11 +296,11 @@
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
func @matmul() {
- %arg0 = iree.placeholder for "interace buffer" {binding = @io::@arg0} : memref<?x?xf32>
- %arg1 = iree.placeholder for "interace buffer" {binding = @io::@arg1} : memref<?x?xf32>
- %arg2 = iree.placeholder for "interace buffer" {binding = @io::@ret0} : memref<?x?xf32>
- %c4 = constant 4 : index
%c0 = constant 0 : index
+ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32>
+ %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<?x?xf32>
+ %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<?x?xf32>
+ %c4 = constant 4 : index
%c1 = constant 1 : index
%0 = memref.dim %arg0, %c1 : memref<?x?xf32>
%1 = "gpu.block_id"() {dimension = "x"} : () -> index
@@ -371,9 +363,10 @@
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} {
func @conv_1d() attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}} {
%cst = constant 0.000000e+00 : f32
- %0 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<3x6x1xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<3x8x1xf32>
- %2 = iree.placeholder for "interface buffer" {binding = @io::@arg1} : memref<3x1x1xf32>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<3x6x1xf32>
+ %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<3x8x1xf32>
+ %2 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x1x1xf32>
%3 = "gpu.block_id"() {dimension = "x"} : () -> index
%4 = "gpu.block_id"() {dimension = "y"} : () -> index
%5 = "gpu.block_id"() {dimension = "z"} : () -> index
@@ -435,11 +428,11 @@
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
func @conv_no_padding() {
- %arg0 = iree.placeholder for "interace buffer" {binding = @io::@arg0} : memref<?x?x?x?xf32>
- %arg1 = iree.placeholder for "interace buffer" {binding = @io::@arg1} : memref<?x?x?x?xf32>
- %arg2 = iree.placeholder for "interace buffer" {binding = @io::@ret0} : memref<?x?x?x?xf32>
- %c2 = constant 2 : index
%c0 = constant 0 : index
+ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?x?x?xf32>
+ %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<?x?x?x?xf32>
+ %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<?x?x?x?xf32>
+ %c2 = constant 2 : index
%c3 = constant 3 : index
%c1 = constant 1 : index
%0 = memref.dim %arg0, %c0 : memref<?x?x?x?xf32>
@@ -494,10 +487,9 @@
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 * 4)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 * 32)>
// CHECK: func @conv_no_padding
-// CHECK-DAG: %[[ARG0:.+]] = iree.placeholder for "interace buffer" {binding = @io::@arg0}
-// CHECK-DAG: %[[ARG1:.+]] = iree.placeholder for "interace buffer" {binding = @io::@arg1}
-// CHECK-DAG: %[[RET0:.+]] = iree.placeholder for "interace buffer" {binding = @io::@ret0}
-// CHECK-DAG: %[[C0:.+]] = constant 0
+// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0
+// CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1
+// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
// CHECK-DAG: %[[C1:.+]] = constant 1
// CHECK-DAG: %[[C2:.+]] = constant 2
// CHECK-DAG: %[[N:.+]] = memref.dim %[[ARG1]], %[[C0]]
@@ -546,9 +538,10 @@
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} {
func @conv_3d() attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}} {
%cst = constant 0.000000e+00 : f32
- %0 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<2x7x7x7x2xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<2x8x8x8x3xf32>
- %2 = iree.placeholder for "interface buffer" {binding = @io::@arg1} : memref<2x2x2x3x2xf32>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<2x7x7x7x2xf32>
+ %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<2x8x8x8x3xf32>
+ %2 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<2x2x2x3x2xf32>
%3 = "gpu.block_id"() {dimension = "x"} : () -> index
%4 = "gpu.block_id"() {dimension = "y"} : () -> index
%5 = "gpu.block_id"() {dimension = "z"} : () -> index
@@ -611,9 +604,10 @@
}
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
func @pooling_nhwc_max() attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}} {
- %0 = iree.placeholder for "interace buffer" {binding = @io::@arg0, operand_result_index = 0 : i32} : memref<2x16x16x6xf32>
- %1 = iree.placeholder for "interace buffer" {binding = @io::@arg1, operand_result_index = 1 : i32} : memref<3x4xf32>
- %2 = iree.placeholder for "interace buffer" {binding = @io::@ret0, operand_result_index = 2 : i32} : memref<2x14x13x6xf32>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<2x16x16x6xf32>
+ %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x4xf32>
+ %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<2x14x13x6xf32>
%3 = "gpu.block_id"() {dimension = "x"} : () -> index
%4 = "gpu.block_id"() {dimension = "y"} : () -> index
%5 = affine.apply #map0()[%4]
@@ -640,9 +634,9 @@
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 * 4)>
// CHECK-DAG: #[[MAP2:.+]] = affine_map<()[s0] -> (s0 * 32)>
// CHECK: func @pooling_nhwc_max
-// CHECK-DAG: %[[ARG0:.+]] = iree.placeholder for "interace buffer" {binding = @io::@arg0, operand_result_index = 0 : i32}
-// CHECK-DAG: %[[ARG1:.+]] = iree.placeholder for "interace buffer" {binding = @io::@arg1, operand_result_index = 1 : i32}
-// CHECK-DAG: %[[RET0:.+]] = iree.placeholder for "interace buffer" {binding = @io::@ret0, operand_result_index = 2 : i32}
+// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0
+// CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1
+// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
// CHECK-DAG: %[[BIDX:.+]] = "gpu.block_id"() {dimension = "x"}
// CHECK-DAG: %[[BIDY:.+]] = "gpu.block_id"() {dimension = "y"}
// CHECK: %[[IV1:.+]] = affine.apply #[[MAP0]]()[%[[BIDY]]]
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
index cd6cca7..3d60285 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -iree-codegen-convert-to-spirv %s | IreeFileCheck %s
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
// CHECK-LABEL: spv.module
// CHECK: spv.GlobalVariable @__push_constant_var__ : !spv.ptr<!spv.struct<(!spv.array<5 x i32, stride=4> [0])>, PushConstant>
// CHECK: spv.func @push_constant()
@@ -22,16 +22,17 @@
// -----
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
// CHECK-LABEL: spv.module
// CHECK: spv.GlobalVariable @[[RET:.+]] bind(3, 4) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: spv.GlobalVariable @[[ARG0:.+]] bind(1, 2) {aliased} : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: spv.GlobalVariable @[[ARG1:.+]] bind(1, 2) {aliased} : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: spv.func @resource_bindings_in_same_entry_func()
func @resource_bindings_in_same_entry_func() {
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4x4xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4x4xf32>
- %2 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4x4xf32>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32>
+ %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32>
+ %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4x4xf32>
return
}
@@ -43,7 +44,7 @@
// -----
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
// CHECK-LABEL: spv.module
// CHECK: spv.GlobalVariable @[[FUNC2_RET:.+]] bind(3, 4) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: spv.GlobalVariable @[[FUNC2_ARG:.+]] bind(1, 2) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer>
@@ -54,8 +55,9 @@
func @resource_bindings_in_entry_func1() {
// CHECK: spv.mlir.addressof @[[FUNC1_ARG:.+]]
// CHECK: spv.mlir.addressof @[[FUNC1_RET:.+]]
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4x4xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4xvector<4xf32>>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4xvector<4xf32>>
return
}
@@ -63,8 +65,9 @@
func @resource_bindings_in_entry_func2() {
// CHECK: spv.mlir.addressof @[[FUNC2_ARG]]
// CHECK: spv.mlir.addressof @[[FUNC2_RET]]
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4x4xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4x4xf32>
+ %c0 = constant 0 : index
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4x4xf32>
return
}
@@ -181,7 +184,7 @@
// -----
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} {
func @interface_binding() {
%c0 = constant 0 : index
%0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<8x5xf32>
@@ -206,7 +209,7 @@
// -----
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} {
func @interface_wg_id() {
%0 = hal.interface.workgroup.id[0] : index
%1 = hal.interface.workgroup.id[1] : index
@@ -230,7 +233,7 @@
// -----
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} {
func @interface_wg_count() {
%0 = hal.interface.workgroup.count[0] : index
%1 = hal.interface.workgroup.count[1] : index
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/memref_vecrotization.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/memref_vecrotization.mlir
index 956f3b3..78ec31e 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/memref_vecrotization.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/memref_vecrotization.mlir
@@ -36,8 +36,8 @@
// -----
// CHECK-LABEL: func @resource_copy
-// CHECK: %[[A:.+]] = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x1024xvector<4xf32>>
-// CHECK: %[[B:.+]] = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x1024xvector<4xf32>>
+// CHECK: %[[A:.+]] = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x1024xvector<4xf32>>
+// CHECK: %[[B:.+]] = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x1024xvector<4xf32>>
// CHECK: %[[V:.+]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<4096x1024xvector<4xf32>>
// CHECK: memref.store %[[V]], %[[B]][%{{.*}}, %{{.*}}] : memref<4096x1024xvector<4xf32>>
// CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x1024xvector<4xf32>>, vector<32x8xf32>
@@ -45,8 +45,8 @@
func @resource_copy() {
%cst = constant 0.000000e+00 : f32
%c0 = constant 0 : index
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x4096xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x4096xf32>
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf32>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf32>
%v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf32>, vector<1x4xf32>
vector.transfer_write %v, %1[%c0, %c0] : vector<1x4xf32>, memref<4096x4096xf32>
%mat = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf32>, vector<32x8xf32>
@@ -62,8 +62,8 @@
// -----
// CHECK-LABEL: func @resource_copy_f16
-// CHECK: %[[A:.+]] = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x1024xvector<4xf16>>
-// CHECK: %[[B:.+]] = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x1024xvector<4xf16>>
+// CHECK: %[[A:.+]] = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x1024xvector<4xf16>>
+// CHECK: %[[B:.+]] = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x1024xvector<4xf16>>
// CHECK: %[[V:.+]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<4096x1024xvector<4xf16>>
// CHECK: memref.store %[[V]], %[[B]][%{{.*}}, %{{.*}}] : memref<4096x1024xvector<4xf16>>
// CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x1024xvector<4xf16>>, vector<32x8xf16>
@@ -71,8 +71,8 @@
func @resource_copy_f16() {
%cst = constant 0.000000e+00 : f16
%c0 = constant 0 : index
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x4096xf16>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x4096xf16>
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16>
%v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<1x4xf16>
vector.transfer_write %v, %1[%c0, %c0] : vector<1x4xf16>, memref<4096x4096xf16>
%mat = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<32x8xf16>
@@ -88,8 +88,8 @@
// -----
// CHECK-LABEL: func @resource_copy_8xf16
-// CHECK: %[[A:.+]] = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x512xvector<4xf32>>
-// CHECK: %[[B:.+]] = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x512xvector<4xf32>>
+// CHECK: %[[A:.+]] = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x512xvector<4xf32>>
+// CHECK: %[[B:.+]] = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x512xvector<4xf32>>
// CHECK: %[[V:.+]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<4096x512xvector<4xf32>>
// CHECK: memref.store %[[V]], %[[B]][%{{.*}}, %{{.*}}] : memref<4096x512xvector<4xf32>>
// CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x512xvector<4xf32>>, vector<32x8xf16>
@@ -97,8 +97,8 @@
func @resource_copy_8xf16() {
%cst = constant 0.000000e+00 : f16
%c0 = constant 0 : index
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4096x4096xf16>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4096x4096xf16>
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16>
%v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<1x8xf16>
vector.transfer_write %v, %1[%c0, %c0] : vector<1x8xf16>, memref<4096x4096xf16>
%mat = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<32x8xf16>
@@ -117,12 +117,12 @@
func @do_not_vectorize_odd_vector_size() {
%cst = constant 0.0 : f32
%c0 = constant 0 : index
- // CHECK: iree.placeholder
+ // CHECK: hal.interface.binding.subspan
// CHECK-SAME: memref<4x3xf32>
- %0 = iree.placeholder for "interface buffer" {binding = @io::@arg0} : memref<4x3xf32>
- // CHECK: iree.placeholder
+ %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x3xf32>
+ // CHECK: hal.interface.binding.subspan
// CHECK-SAME: memref<4x3xf32>
- %1 = iree.placeholder for "interface buffer" {binding = @io::@ret0} : memref<4x3xf32>
+ %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4x3xf32>
%v = vector.transfer_read %0[%c0, %c0], %cst : memref<4x3xf32>, vector<3xf32>
vector.transfer_write %v, %1[%c0, %c0] : vector<3xf32>, memref<4x3xf32>
return
