diff --git a/mlir/lib/Transforms/Utils/DialectConversion.cpp b/mlir/lib/Transforms/Utils/DialectConversion.cpp index 2f0efe1b1e454e..1c0cb128aeabec 100644 --- a/mlir/lib/Transforms/Utils/DialectConversion.cpp +++ b/mlir/lib/Transforms/Utils/DialectConversion.cpp @@ -2857,13 +2857,7 @@ static LogicalResult legalizeUnresolvedMaterialization( // Try to materialize the conversion. if (const TypeConverter *converter = mat.getConverter()) { - // FIXME: Determine a suitable insertion location when there are multiple - // inputs. - if (inputOperands.size() == 1) - rewriter.setInsertionPointAfterValue(inputOperands.front()); - else - rewriter.setInsertionPoint(op); - + rewriter.setInsertionPoint(op); Value newMaterialization; switch (mat.getMaterializationKind()) { case MaterializationKind::Argument: diff --git a/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir b/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir index 56ae930e6d6273..d3bdbe89a54876 100644 --- a/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir +++ b/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir @@ -478,9 +478,10 @@ func.func @mului_extended_vector1d(%arg0: vector<3xi64>, %arg1: vector<3xi64>) - // ----- // CHECK-LABEL: func @cmpf_2dvector( +// CHECK-SAME: %[[OARG0:.*]]: vector<4x3xf32>, %[[OARG1:.*]]: vector<4x3xf32>) func.func @cmpf_2dvector(%arg0 : vector<4x3xf32>, %arg1 : vector<4x3xf32>) { - // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast - // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast + // CHECK-DAG: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[OARG0]] + // CHECK-DAG: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[OARG1]] // CHECK: %[[EXTRACT1:.*]] = llvm.extractvalue %[[ARG0]][0] : !llvm.array<4 x vector<3xf32>> // CHECK: %[[EXTRACT2:.*]] = llvm.extractvalue %[[ARG1]][0] : !llvm.array<4 x vector<3xf32>> // CHECK: %[[CMP:.*]] = llvm.fcmp "olt" %[[EXTRACT1]], %[[EXTRACT2]] : vector<3xf32> @@ -492,9 +493,10 @@ func.func @cmpf_2dvector(%arg0 : vector<4x3xf32>, %arg1 : vector<4x3xf32>) { // ----- // CHECK-LABEL: func @cmpi_0dvector( +// CHECK-SAME: %[[OARG0:.*]]: vector, %[[OARG1:.*]]: vector) func.func @cmpi_0dvector(%arg0 : vector, %arg1 : vector) { - // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast - // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast + // CHECK-DAG: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[OARG0]] + // CHECK-DAG: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[OARG1]] // CHECK: %[[CMP:.*]] = llvm.icmp "ult" %[[ARG0]], %[[ARG1]] : vector<1xi32> %0 = arith.cmpi ult, %arg0, %arg1 : vector func.return @@ -503,9 +505,10 @@ func.func @cmpi_0dvector(%arg0 : vector, %arg1 : vector) { // ----- // CHECK-LABEL: func @cmpi_2dvector( +// CHECK-SAME: %[[OARG0:.*]]: vector<4x3xi32>, %[[OARG1:.*]]: vector<4x3xi32>) func.func @cmpi_2dvector(%arg0 : vector<4x3xi32>, %arg1 : vector<4x3xi32>) { - // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast - // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast + // CHECK-DAG: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[OARG0]] + // CHECK-DAG: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[OARG1]] // CHECK: %[[EXTRACT1:.*]] = llvm.extractvalue %[[ARG0]][0] : !llvm.array<4 x vector<3xi32>> // CHECK: %[[EXTRACT2:.*]] = llvm.extractvalue %[[ARG1]][0] : !llvm.array<4 x vector<3xi32>> // CHECK: %[[CMP:.*]] = llvm.icmp "ult" %[[EXTRACT1]], %[[EXTRACT2]] : vector<3xi32> diff --git a/mlir/test/Conversion/ArithToLLVM/convert-nd-vector-to-llvmir.mlir b/mlir/test/Conversion/ArithToLLVM/convert-nd-vector-to-llvmir.mlir index 63989347567b56..b234cbbb35f327 100644 --- a/mlir/test/Conversion/ArithToLLVM/convert-nd-vector-to-llvmir.mlir +++ b/mlir/test/Conversion/ArithToLLVM/convert-nd-vector-to-llvmir.mlir @@ -199,9 +199,9 @@ func.func @bitcast_2d(%arg0: vector<2x4xf32>) { // CHECK-LABEL: func @select_2d( func.func @select_2d(%arg0 : vector<4x3xi1>, %arg1 : vector<4x3xi32>, %arg2 : vector<4x3xi32>) { - // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %arg0 - // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %arg1 - // CHECK: %[[ARG2:.*]] = builtin.unrealized_conversion_cast %arg2 + // CHECK-DAG: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %arg0 + // CHECK-DAG: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %arg1 + // CHECK-DAG: %[[ARG2:.*]] = builtin.unrealized_conversion_cast %arg2 // CHECK: %[[EXTRACT1:.*]] = llvm.extractvalue %[[ARG0]][0] : !llvm.array<4 x vector<3xi1>> // CHECK: %[[EXTRACT2:.*]] = llvm.extractvalue %[[ARG1]][0] : !llvm.array<4 x vector<3xi32>> // CHECK: %[[EXTRACT3:.*]] = llvm.extractvalue %[[ARG2]][0] : !llvm.array<4 x vector<3xi32>> diff --git a/mlir/test/Conversion/ArithToSPIRV/arith-to-spirv.mlir b/mlir/test/Conversion/ArithToSPIRV/arith-to-spirv.mlir index ae47ae36ca51cd..beb2c8d2d242c0 100644 --- a/mlir/test/Conversion/ArithToSPIRV/arith-to-spirv.mlir +++ b/mlir/test/Conversion/ArithToSPIRV/arith-to-spirv.mlir @@ -60,8 +60,8 @@ func.func @index_scalar(%lhs: index, %rhs: index) { // CHECK-LABEL: @index_scalar_srem // CHECK-SAME: (%[[A:.+]]: index, %[[B:.+]]: index) func.func @index_scalar_srem(%lhs: index, %rhs: index) { - // CHECK: %[[LHS:.+]] = builtin.unrealized_conversion_cast %[[A]] : index to i32 - // CHECK: %[[RHS:.+]] = builtin.unrealized_conversion_cast %[[B]] : index to i32 + // CHECK-DAG: %[[LHS:.+]] = builtin.unrealized_conversion_cast %[[A]] : index to i32 + // CHECK-DAG: %[[RHS:.+]] = builtin.unrealized_conversion_cast %[[B]] : index to i32 // CHECK: %[[LABS:.+]] = spirv.GL.SAbs %[[LHS]] : i32 // CHECK: %[[RABS:.+]] = spirv.GL.SAbs %[[RHS]] : i32 // CHECK: %[[ABS:.+]] = spirv.UMod %[[LABS]], %[[RABS]] : i32 diff --git a/mlir/test/Conversion/IndexToLLVM/index-to-llvm.mlir b/mlir/test/Conversion/IndexToLLVM/index-to-llvm.mlir index 1b13ebb38dc9e9..26abb3bdc23a16 100644 --- a/mlir/test/Conversion/IndexToLLVM/index-to-llvm.mlir +++ b/mlir/test/Conversion/IndexToLLVM/index-to-llvm.mlir @@ -50,8 +50,8 @@ func.func @trivial_ops(%a: index, %b: index) { // CHECK-LABEL: @ceildivs // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @ceildivs(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : // CHECK: %[[POS_ONE:.*]] = llvm.mlir.constant(1 : // CHECK: %[[NEG_ONE:.*]] = llvm.mlir.constant(-1 : @@ -82,8 +82,8 @@ func.func @ceildivs(%n: index, %m: index) -> index { // CHECK-LABEL: @ceildivu // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @ceildivu(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : // CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : @@ -103,11 +103,11 @@ func.func @ceildivu(%n: index, %m: index) -> index { // CHECK-LABEL: @floordivs // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @floordivs(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] - // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : - // CHECK: %[[POS_ONE:.*]] = llvm.mlir.constant(1 : - // CHECK: %[[NEG_ONE:.*]] = llvm.mlir.constant(-1 : + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[ZERO:.*]] = llvm.mlir.constant(0 : + // CHECK-DAG: %[[POS_ONE:.*]] = llvm.mlir.constant(1 : + // CHECK-DAG: %[[NEG_ONE:.*]] = llvm.mlir.constant(-1 : // CHECK: %[[M_NEG:.*]] = llvm.icmp "slt" %[[M]], %[[ZERO]] // CHECK: %[[X:.*]] = llvm.select %[[M_NEG]], %[[POS_ONE]], %[[NEG_ONE]] diff --git a/mlir/test/Conversion/IndexToSPIRV/index-to-spirv.mlir b/mlir/test/Conversion/IndexToSPIRV/index-to-spirv.mlir index 53dc896e98c7d6..7b26f4fc136965 100644 --- a/mlir/test/Conversion/IndexToSPIRV/index-to-spirv.mlir +++ b/mlir/test/Conversion/IndexToSPIRV/index-to-spirv.mlir @@ -67,8 +67,8 @@ func.func @constant_ops() { // CHECK-LABEL: @ceildivs // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @ceildivs(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] // CHECK: %[[ZERO:.*]] = spirv.Constant 0 // CHECK: %[[POS_ONE:.*]] = spirv.Constant 1 // CHECK: %[[NEG_ONE:.*]] = spirv.Constant -1 @@ -99,8 +99,8 @@ func.func @ceildivs(%n: index, %m: index) -> index { // CHECK-LABEL: @ceildivu // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @ceildivu(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] // CHECK: %[[ZERO:.*]] = spirv.Constant 0 // CHECK: %[[ONE:.*]] = spirv.Constant 1 @@ -120,8 +120,8 @@ func.func @ceildivu(%n: index, %m: index) -> index { // CHECK-LABEL: @floordivs // CHECK-SAME: %[[NI:.*]]: index, %[[MI:.*]]: index func.func @floordivs(%n: index, %m: index) -> index { - // CHECK: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] - // CHECK: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] + // CHECK-DAG: %[[N:.*]] = builtin.unrealized_conversion_cast %[[NI]] + // CHECK-DAG: %[[M:.*]] = builtin.unrealized_conversion_cast %[[MI]] // CHECK: %[[ZERO:.*]] = spirv.Constant 0 // CHECK: %[[POS_ONE:.*]] = spirv.Constant 1 // CHECK: %[[NEG_ONE:.*]] = spirv.Constant -1 diff --git a/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir b/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir index f0119afa42f690..586460fd581805 100644 --- a/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir +++ b/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir @@ -78,8 +78,8 @@ func.func @copy_sign_vector_0D(%value: vector<1xf16>, %sign: vector<1xf16>) -> v // CHECK-LABEL: func @copy_sign_vector_0D // CHECK-SAME: (%[[VALUE:.+]]: vector<1xf16>, %[[SIGN:.+]]: vector<1xf16>) -// CHECK: %[[CASTVAL:.+]] = builtin.unrealized_conversion_cast %[[VALUE]] : vector<1xf16> to f16 -// CHECK: %[[CASTSIGN:.+]] = builtin.unrealized_conversion_cast %[[SIGN]] : vector<1xf16> to f16 +// CHECK-DAG: %[[CASTVAL:.+]] = builtin.unrealized_conversion_cast %[[VALUE]] : vector<1xf16> to f16 +// CHECK-DAG: %[[CASTSIGN:.+]] = builtin.unrealized_conversion_cast %[[SIGN]] : vector<1xf16> to f16 // CHECK: %[[SMASK:.+]] = spirv.Constant -32768 : i16 // CHECK: %[[VMASK:.+]] = spirv.Constant 32767 : i16 // CHECK: %[[VCAST:.+]] = spirv.Bitcast %[[CASTVAL]] : f16 to i16 diff --git a/mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir b/mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir index 9d8f4266adf27f..ebfc3b95d6effd 100644 --- a/mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir +++ b/mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir @@ -506,14 +506,15 @@ func.func @memref_reinterpret_cast_unranked_to_dynamic_shape(%offset: index, // ----- -// CHECK-LABEL: @memref_reshape +// CHECK-LABEL: @memref_reshape( +// CHECK-SAME: %[[ARG0:.*]]: memref<2x3xf32>, %[[ARG1:.*]]: memref) func.func @memref_reshape(%input : memref<2x3xf32>, %shape : memref) { %output = memref.reshape %input(%shape) : (memref<2x3xf32>, memref) -> memref<*xf32> return } -// CHECK: [[INPUT:%.*]] = builtin.unrealized_conversion_cast %{{.*}} to [[INPUT_TY:!.*]] -// CHECK: [[SHAPE:%.*]] = builtin.unrealized_conversion_cast %{{.*}} to [[SHAPE_TY:!.*]] +// CHECK-DAG: [[INPUT:%.*]] = builtin.unrealized_conversion_cast %[[ARG0]] : {{.*}} to [[INPUT_TY:!.*]] +// CHECK-DAG: [[SHAPE:%.*]] = builtin.unrealized_conversion_cast %[[ARG1]] : {{.*}} to [[SHAPE_TY:!.*]] // CHECK: [[RANK:%.*]] = llvm.extractvalue [[SHAPE]][3, 0] : [[SHAPE_TY]] // CHECK: [[UNRANKED_OUT_O:%.*]] = llvm.mlir.undef : !llvm.struct<(i64, ptr)> // CHECK: [[UNRANKED_OUT_1:%.*]] = llvm.insertvalue [[RANK]], [[UNRANKED_OUT_O]][0] : !llvm.struct<(i64, ptr)> diff --git a/mlir/test/Conversion/MemRefToLLVM/convert-static-memref-ops.mlir b/mlir/test/Conversion/MemRefToLLVM/convert-static-memref-ops.mlir index f1600d43e7bfb3..96cf2264e9c2f3 100644 --- a/mlir/test/Conversion/MemRefToLLVM/convert-static-memref-ops.mlir +++ b/mlir/test/Conversion/MemRefToLLVM/convert-static-memref-ops.mlir @@ -115,13 +115,11 @@ func.func @zero_d_load(%arg0: memref) -> f32 { // ----- -// CHECK-LABEL: func @static_load -// CHECK: %[[MEMREF:.*]]: memref<10x42xf32>, -// CHECK: %[[I:.*]]: index, -// CHECK: %[[J:.*]]: index) +// CHECK-LABEL: func @static_load( +// CHECK-SAME: %[[MEMREF:.*]]: memref<10x42xf32>, %[[I:.*]]: index, %[[J:.*]]: index) func.func @static_load(%static : memref<10x42xf32>, %i : index, %j : index) { -// CHECK: %[[II:.*]] = builtin.unrealized_conversion_cast %[[I]] -// CHECK: %[[JJ:.*]] = builtin.unrealized_conversion_cast %[[J]] +// CHECK-DAG: %[[II:.*]] = builtin.unrealized_conversion_cast %[[I]] +// CHECK-DAG: %[[JJ:.*]] = builtin.unrealized_conversion_cast %[[J]] // CHECK: %[[ptr:.*]] = llvm.extractvalue %{{.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> // CHECK: %[[st0:.*]] = llvm.mlir.constant(42 : index) : i64 // CHECK: %[[offI:.*]] = llvm.mul %[[II]], %[[st0]] : i64 @@ -148,8 +146,8 @@ func.func @zero_d_store(%arg0: memref, %arg1: f32) { // CHECK: %[[MEMREF:.*]]: memref<10x42xf32>, // CHECK-SAME: %[[I:.*]]: index, %[[J:.*]]: index, func.func @static_store(%static : memref<10x42xf32>, %i : index, %j : index, %val : f32) { -// CHECK: %[[II:.*]] = builtin.unrealized_conversion_cast %[[I]] -// CHECK: %[[JJ:.*]] = builtin.unrealized_conversion_cast %[[J]] +// CHECK-DAG: %[[II:.*]] = builtin.unrealized_conversion_cast %[[I]] +// CHECK-DAG: %[[JJ:.*]] = builtin.unrealized_conversion_cast %[[J]] // CHECK: %[[ptr:.*]] = llvm.extractvalue %{{.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> // CHECK: %[[st0:.*]] = llvm.mlir.constant(42 : index) : i64 // CHECK: %[[offI:.*]] = llvm.mul %[[II]], %[[st0]] : i64 @@ -205,7 +203,7 @@ module attributes { dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry> } { func.func @address() { %c1 = arith.constant 1 : index %0 = memref.alloc(%c1) : memref> - // CHECK: %[[CST_S:.*]] = arith.constant 1 : index + // CHECK-DAG: %[[CST_S:.*]] = arith.constant 1 : index // CHECK: %[[CST:.*]] = builtin.unrealized_conversion_cast // CHECK: llvm.mlir.zero // CHECK: llvm.getelementptr %{{.*}}[[CST]] @@ -269,8 +267,8 @@ func.func @memref.reshape(%arg0: memref<4x5x6xf32>) -> memref<2x6x20xf32> { // CHECK-LABEL: func @memref.reshape.dynamic.dim // CHECK-SAME: %[[arg:.*]]: memref, %[[shape:.*]]: memref<4xi64>) -> memref func.func @memref.reshape.dynamic.dim(%arg: memref, %shape: memref<4xi64>) -> memref { - // CHECK: %[[arg_cast:.*]] = builtin.unrealized_conversion_cast %[[arg]] : memref to !llvm.struct<(ptr, ptr, i64, array<3 x i64>, array<3 x i64>)> - // CHECK: %[[shape_cast:.*]] = builtin.unrealized_conversion_cast %[[shape]] : memref<4xi64> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> + // CHECK-DAG: %[[arg_cast:.*]] = builtin.unrealized_conversion_cast %[[arg]] : memref to !llvm.struct<(ptr, ptr, i64, array<3 x i64>, array<3 x i64>)> + // CHECK-DAG: %[[shape_cast:.*]] = builtin.unrealized_conversion_cast %[[shape]] : memref<4xi64> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> // CHECK: %[[undef:.*]] = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<4 x i64>, array<4 x i64>)> // CHECK: %[[alloc_ptr:.*]] = llvm.extractvalue %[[arg_cast]][0] : !llvm.struct<(ptr, ptr, i64, array<3 x i64>, array<3 x i64>)> // CHECK: %[[align_ptr:.*]] = llvm.extractvalue %[[arg_cast]][1] : !llvm.struct<(ptr, ptr, i64, array<3 x i64>, array<3 x i64>)> @@ -318,8 +316,8 @@ func.func @memref.reshape.dynamic.dim(%arg: memref, %shape: memref<4x // CHECK-LABEL: func @memref.reshape_index // CHECK-SAME: %[[arg:.*]]: memref, %[[shape:.*]]: memref<1xindex> func.func @memref.reshape_index(%arg0: memref, %shape: memref<1xindex>) -> memref { - // CHECK: %[[arg_cast:.*]] = builtin.unrealized_conversion_cast %[[arg]] : memref to !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> - // CHECK: %[[shape_cast:.*]] = builtin.unrealized_conversion_cast %[[shape]] : memref<1xindex> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> + // CHECK-DAG: %[[arg_cast:.*]] = builtin.unrealized_conversion_cast %[[arg]] : memref to !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> + // CHECK-DAG: %[[shape_cast:.*]] = builtin.unrealized_conversion_cast %[[shape]] : memref<1xindex> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> // CHECK: %[[undef:.*]] = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> // CHECK: %[[alloc_ptr:.*]] = llvm.extractvalue %[[arg_cast]][0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> // CHECK: %[[align_ptr:.*]] = llvm.extractvalue %[[arg_cast]][1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)> diff --git a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir index 882804132e66dd..9dc22abf143bf0 100644 --- a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir +++ b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir @@ -10,9 +10,9 @@ // CHECK-LABEL: func @view( // CHECK: %[[ARG0F:.*]]: index, %[[ARG1F:.*]]: index, %[[ARG2F:.*]]: index func.func @view(%arg0 : index, %arg1 : index, %arg2 : index) { - // CHECK: %[[ARG2:.*]] = builtin.unrealized_conversion_cast %[[ARG2F:.*]] - // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[ARG0F:.*]] - // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[ARG1F:.*]] + // CHECK-DAG: %[[ARG2:.*]] = builtin.unrealized_conversion_cast %[[ARG2F]] + // CHECK-DAG: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[ARG0F]] + // CHECK-DAG: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[ARG1F]] // CHECK: llvm.mlir.constant(2048 : index) : i64 // CHECK: llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> %0 = memref.alloc() : memref<2048xi8> @@ -408,8 +408,8 @@ func.func @atomic_rmw_with_offset(%I : memref<10xi32, strided<[1], offset: 5>>, // CHECK-SAME: %[[ARG0:.+]]: memref<10xi32, strided<[1], offset: 5>> // CHECK-SAME: %[[ARG1:.+]]: i32 // CHECK-SAME: %[[ARG2:.+]]: index -// CHECK: %[[MEMREF_STRUCT:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : memref<10xi32, strided<[1], offset: 5>> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> -// CHECK: %[[INDEX:.+]] = builtin.unrealized_conversion_cast %[[ARG2]] : index to i64 +// CHECK-DAG: %[[MEMREF_STRUCT:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : memref<10xi32, strided<[1], offset: 5>> to !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> +// CHECK-DAG: %[[INDEX:.+]] = builtin.unrealized_conversion_cast %[[ARG2]] : index to i64 // CHECK: %[[BASE_PTR:.+]] = llvm.extractvalue %[[MEMREF_STRUCT]][1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> // CHECK: %[[OFFSET:.+]] = llvm.mlir.constant(5 : index) : i64 // CHECK: %[[OFFSET_PTR:.+]] = llvm.getelementptr %[[BASE_PTR]][%[[OFFSET]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32 diff --git a/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir b/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir index 52ed14e8cce233..bb003e11993f5f 100644 --- a/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir +++ b/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir @@ -196,8 +196,8 @@ func.func @load_i4(%arg0: memref>, %i: // CHECK-LABEL: @store_i4 func.func @store_i4(%arg0: memref>, %value: i4, %i: index) { - // CHECK: %[[VAL:.+]] = builtin.unrealized_conversion_cast %{{.+}} : i4 to i32 - // CHECK: %[[INDEX:.+]] = builtin.unrealized_conversion_cast %{{.+}} : index to i32 + // CHECK-DAG: %[[VAL:.+]] = builtin.unrealized_conversion_cast %{{.+}} : i4 to i32 + // CHECK-DAG: %[[INDEX:.+]] = builtin.unrealized_conversion_cast %{{.+}} : index to i32 // CHECK: %[[ZERO:.+]] = spirv.Constant 0 : i32 // CHECK: %[[EIGHT:.+]] = spirv.Constant 8 : i32 // CHECK: %[[FOUR:.+]] = spirv.Constant 4 : i32 diff --git a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir index 10c03a270005f1..6dd5b1988e2a2f 100644 --- a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir +++ b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir @@ -16,10 +16,11 @@ module attributes { #spirv.resource_limits<>> } { -// CHECK-LABEL: @load_store_zero_rank_float +// CHECK-LABEL: @load_store_zero_rank_float( +// CHECK-SAME: %[[OARG0:.*]]: memref{{.*}}, %[[OARG1:.*]]: memref func.func @load_store_zero_rank_float(%arg0: memref>, %arg1: memref>) { - // CHECK: [[ARG0:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> - // CHECK: [[ARG1:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr [0])>, StorageBuffer> // CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32 // CHECK: spirv.AccessChain [[ARG0]][ // CHECK-SAME: [[ZERO]], [[ZERO]] @@ -35,9 +36,10 @@ func.func @load_store_zero_rank_float(%arg0: memref>, %arg1: memref>) { - // CHECK: [[ARG0:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> - // CHECK: [[ARG1:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr [0])>, StorageBuffer> // CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32 // CHECK: spirv.AccessChain [[ARG0]][ // CHECK-SAME: [[ZERO]], [[ZERO]] @@ -53,9 +55,10 @@ func.func @load_store_zero_rank_int(%arg0: memref>, %dest: memref>) { - // CHECK: %[[SRC:.+]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> - // CHECK: %[[DST:.+]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: %[[SRC:.+]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr [0])>, StorageBuffer> + // CHECK-DAG: %[[DST:.+]] = builtin.unrealized_conversion_cast %[[OARG2]] : memref> to !spirv.ptr [0])>, StorageBuffer> // CHECK: %[[AC0:.+]] = spirv.AccessChain %[[SRC]] // CHECK: spirv.Load "StorageBuffer" %[[AC0]] %0 = memref.load %source[%i] : memref> @@ -173,9 +176,10 @@ module attributes { } { // CHECK-LABEL: @load_store_zero_rank_float +// CHECK-SAME: %[[OARG0:.*]]: memref{{.*}}, %[[OARG1:.*]]: memref func.func @load_store_zero_rank_float(%arg0: memref>, %arg1: memref>) { - // CHECK: [[ARG0:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr, CrossWorkgroup> - // CHECK: [[ARG1:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr, CrossWorkgroup> + // CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref> to !spirv.ptr, CrossWorkgroup> + // CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr, CrossWorkgroup> // CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32 // CHECK: spirv.AccessChain [[ARG0]][ // CHECK-SAME: [[ZERO]] @@ -191,9 +195,10 @@ func.func @load_store_zero_rank_float(%arg0: memref>, %arg1: memref>) { - // CHECK: [[ARG0:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr, CrossWorkgroup> - // CHECK: [[ARG1:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr, CrossWorkgroup> + // CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref> to !spirv.ptr, CrossWorkgroup> + // CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr, CrossWorkgroup> // CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32 // CHECK: spirv.AccessChain [[ARG0]][ // CHECK-SAME: [[ZERO]] @@ -209,9 +214,10 @@ func.func @load_store_zero_rank_int(%arg0: memref>, %dest: memref>) { - // CHECK: %[[SRC:.+]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr - // CHECK: %[[DST:.+]] = builtin.unrealized_conversion_cast {{.+}} : memref> to !spirv.ptr + // CHECK-DAG: %[[SRC:.+]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref> to !spirv.ptr + // CHECK-DAG: %[[DST:.+]] = builtin.unrealized_conversion_cast %[[OARG2]] : memref> to !spirv.ptr // CHECK: %[[AC0:.+]] = spirv.PtrAccessChain %[[SRC]] // CHECK: spirv.Load "CrossWorkgroup" %[[AC0]] %0 = memref.load %source[%i] : memref> @@ -328,8 +334,8 @@ module attributes { // CHECK-LABEL: func.func @reinterpret_cast // CHECK-SAME: (%[[MEM:.*]]: memref>, %[[OFF:.*]]: index) func.func @reinterpret_cast(%arg: memref>, %arg1: index) -> memref, #spirv.storage_class> { -// CHECK: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref> to !spirv.ptr -// CHECK: %[[OFF1:.*]] = builtin.unrealized_conversion_cast %[[OFF]] : index to i32 +// CHECK-DAG: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref> to !spirv.ptr +// CHECK-DAG: %[[OFF1:.*]] = builtin.unrealized_conversion_cast %[[OFF]] : index to i32 // CHECK: %[[RET:.*]] = spirv.InBoundsPtrAccessChain %[[MEM1]][%[[OFF1]]] : !spirv.ptr, i32 // CHECK: %[[RET1:.*]] = builtin.unrealized_conversion_cast %[[RET]] : !spirv.ptr to memref, #spirv.storage_class> // CHECK: return %[[RET1]] @@ -340,8 +346,8 @@ func.func @reinterpret_cast(%arg: memref>) func.func @reinterpret_cast_0(%arg: memref>) -> memref, #spirv.storage_class> { -// CHECK: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref> to !spirv.ptr -// CHECK: %[[RET:.*]] = builtin.unrealized_conversion_cast %[[MEM1]] : !spirv.ptr to memref, #spirv.storage_class> +// CHECK-DAG: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref> to !spirv.ptr +// CHECK-DAG: %[[RET:.*]] = builtin.unrealized_conversion_cast %[[MEM1]] : !spirv.ptr to memref, #spirv.storage_class> // CHECK: return %[[RET]] %ret = memref.reinterpret_cast %arg to offset: [0], sizes: [10], strides: [1] : memref> to memref, #spirv.storage_class> return %ret : memref, #spirv.storage_class> @@ -375,8 +381,8 @@ module attributes { // CHECK-LABEL: func.func @cast // CHECK-SAME: (%[[MEM:.*]]: memref<4x?xf32, #spirv.storage_class>) func.func @cast(%arg: memref<4x?xf32, #spirv.storage_class>) -> memref> { -// CHECK: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref<4x?xf32, #spirv.storage_class> to !spirv.ptr -// CHECK: %[[MEM2:.*]] = builtin.unrealized_conversion_cast %[[MEM1]] : !spirv.ptr to memref> +// CHECK-DAG: %[[MEM1:.*]] = builtin.unrealized_conversion_cast %[[MEM]] : memref<4x?xf32, #spirv.storage_class> to !spirv.ptr +// CHECK-DAG: %[[MEM2:.*]] = builtin.unrealized_conversion_cast %[[MEM1]] : !spirv.ptr to memref> // CHECK: return %[[MEM2]] %ret = memref.cast %arg : memref<4x?xf32, #spirv.storage_class> to memref> return %ret : memref> diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir index 73d23679152848..86a552c03a4739 100644 --- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir +++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir @@ -275,8 +275,8 @@ func.func @async_cp_i4( // CHECK-SAME: %[[IDX:[a-zA-Z0-9_]+]]: index, %[[SRCELEMENTS:[a-zA-Z0-9_]+]]: index func.func @async_cp_zfill_f32_align4( %src: memref<128x128xf32>, %dst: memref<3x16x128xf32, 3>, %i : index, %srcElements : index) { - // CHECK: %[[IDX1:.*]] = builtin.unrealized_conversion_cast %[[IDX]] : index to i64 - // CHECK: %[[SRC1:.*]] = builtin.unrealized_conversion_cast %[[SRCELEMENTS]] : index to i64 + // CHECK-DAG: %[[IDX1:.*]] = builtin.unrealized_conversion_cast %[[IDX]] : index to i64 + // CHECK-DAG: %[[SRC1:.*]] = builtin.unrealized_conversion_cast %[[SRCELEMENTS]] : index to i64 // CHECK-DAG: %[[BASEDST:.*]] = llvm.extractvalue %{{.*}}[1] : !llvm.struct<(ptr<3>, ptr<3>, i64, array<3 x i64>, array<3 x i64>)> // CHECK-DAG: %[[S2048:.*]] = llvm.mlir.constant(2048 : index) : i64 // CHECK-DAG: %[[LI1:.*]] = llvm.mul %[[IDX1]], %[[S2048]] : i64 @@ -310,8 +310,8 @@ func.func @async_cp_zfill_f32_align4( // CHECK-SAME: %[[IDX:[a-zA-Z0-9_]+]]: index, %[[SRCELEMENTS:[a-zA-Z0-9_]+]]: index) func.func @async_cp_zfill_f32_align1( %src: memref<128x128xf32>, %dst: memref<3x16x128xf32, 3>, %i : index, %srcElements : index) { - // CHECK: %[[IDX1:.*]] = builtin.unrealized_conversion_cast %[[IDX]] : index to i64 - // CHECK: %[[SRC1:.*]] = builtin.unrealized_conversion_cast %[[SRCELEMENTS]] : index to i64 + // CHECK-DAG: %[[IDX1:.*]] = builtin.unrealized_conversion_cast %[[IDX]] : index to i64 + // CHECK-DAG: %[[SRC1:.*]] = builtin.unrealized_conversion_cast %[[SRCELEMENTS]] : index to i64 // CHECK-DAG: %[[BASEDST:.*]] = llvm.extractvalue %{{.*}}[1] : !llvm.struct<(ptr<3>, ptr<3>, i64, array<3 x i64>, array<3 x i64>)> // CHECK-DAG: %[[S2048:.*]] = llvm.mlir.constant(2048 : index) : i64 // CHECK-DAG: %[[LI1:.*]] = llvm.mul %[[IDX1]], %[[S2048]] : i64 @@ -532,8 +532,11 @@ func.func @mbarrier_nocomplete() { func.return } -// CHECK-LABEL: func @mbarrier_wait +// CHECK-LABEL: func @mbarrier_wait( +// CHECK-SAME: %[[ARG0:.*]]: !nvgpu.mbarrier.group{{.*}}, %[[ARG1:.*]]: !nvgpu.mbarrier.token) func.func @mbarrier_wait(%barriers : !nvgpu.mbarrier.group, num_barriers = 5>, %token : !tokenType) { +// CHECK-DAG: %[[CARG0:.*]] = builtin.unrealized_conversion_cast %[[ARG0]] +// CHECK-DAG: %[[CARG1:.*]] = builtin.unrealized_conversion_cast %[[ARG1]] %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %n = arith.constant 100 : index @@ -545,8 +548,9 @@ func.func @mbarrier_wait(%barriers : !nvgpu.mbarrier.group, ptr<3>, i64, array<1 x i64>, array<1 x i64>)> +// CHECK: %[[S4:.+]] = llvm.extractvalue %[[CARG0]][1] : !llvm.struct<(ptr<3>, ptr<3>, i64, array<1 x i64>, array<1 x i64>)> // CHECK: %[[S5:.+]] = llvm.getelementptr %[[S4]][%[[S3]]] : (!llvm.ptr<3>, i64) -> !llvm.ptr<3>, i64 +// CHECK: nvvm.mbarrier.test.wait.shared {{.*}}, %[[CARG1]] %mbarId = arith.remui %i, %numBarriers : index %isDone = nvgpu.mbarrier.test.wait %barriers[%mbarId], %token : !nvgpu.mbarrier.group, num_barriers = 5>, !tokenType } @@ -871,9 +875,9 @@ func.func @warpgroup_mma_128_128_64( %descB: !nvgpu.warpgroup.descriptor>, %acc: !nvgpu.warpgroup.accumulator>) { -// CHECK: %[[S0:.+]] = builtin.unrealized_conversion_cast %[[arg0]] : !nvgpu.warpgroup.descriptor> to i64 -// CHECK: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor> to i64 -// CHECK: %[[ARG:.+]] = builtin.unrealized_conversion_cast %[[arg2]] : !nvgpu.warpgroup.accumulator> to !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> +// CHECK-DAG: %[[S0:.+]] = builtin.unrealized_conversion_cast %[[arg0]] : !nvgpu.warpgroup.descriptor> to i64 +// CHECK-DAG: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor> to i64 +// CHECK-DAG: %[[ARG:.+]] = builtin.unrealized_conversion_cast %[[arg2]] : !nvgpu.warpgroup.accumulator> to !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: nvvm.wgmma.fence.aligned // CHECK: %[[UD:.+]] = llvm.mlir.undef : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: %[[S2:.+]] = llvm.extractvalue %[[ARG]][0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> @@ -1280,9 +1284,9 @@ func.func @warpgroup_matrix_multiply_m128n128k64( to memref<128x128xf32,3> -// CHECK: %[[S0:.+]] = builtin.unrealized_conversion_cast %[[arg0]] : !nvgpu.warpgroup.descriptor> to i64 -// CHECK: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor> to i64 -// CHECK: %[[S2:.+]] = builtin.unrealized_conversion_cast %[[arg2]] : memref<128x128xf32, 3> to !llvm.struct<(ptr<3>, ptr<3>, i64, array<2 x i64>, array<2 x i64>)> +// CHECK-DAG: %[[S0:.+]] = builtin.unrealized_conversion_cast %[[arg0]] : !nvgpu.warpgroup.descriptor> to i64 +// CHECK-DAG: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor> to i64 +// CHECK-DAG: %[[S2:.+]] = builtin.unrealized_conversion_cast %[[arg2]] : memref<128x128xf32, 3> to !llvm.struct<(ptr<3>, ptr<3>, i64, array<2 x i64>, array<2 x i64>)> // CHECK: %[[S3:.+]] = llvm.mlir.constant(0.000000e+00 : f32) : f32 // CHECK: %[[S4:.+]] = llvm.mlir.undef : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: %[[S5:.+]] = llvm.extractvalue %[[S4]][0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> @@ -1296,7 +1300,7 @@ func.func @warpgroup_matrix_multiply_m128n128k64( // CHECK: nvvm.wgmma.fence.aligned // CHECK: %[[S137:.+]] = llvm.mlir.undef : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: %[[S138:.+]] = llvm.extractvalue %136[0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> -// CHECK: %[[S139:.+]] = nvvm.wgmma.mma_async %0, %1, %[[S138]], , D[, , ], A[, , ], B[, , ] : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> +// CHECK: %[[S139:.+]] = nvvm.wgmma.mma_async %[[S0]], %1, %[[S138]], , D[, , ], A[, , ], B[, , ] : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> // CHECK: nvvm.wgmma.mma_async // CHECK: nvvm.wgmma.mma_async // CHECK: %[[S154:.+]] = nvvm.wgmma.mma_async diff --git a/mlir/test/Conversion/VectorToLLVM/vector-scalable-memcpy.mlir b/mlir/test/Conversion/VectorToLLVM/vector-scalable-memcpy.mlir index 811b10721bf284..80e6caa05db5e5 100644 --- a/mlir/test/Conversion/VectorToLLVM/vector-scalable-memcpy.mlir +++ b/mlir/test/Conversion/VectorToLLVM/vector-scalable-memcpy.mlir @@ -6,8 +6,8 @@ func.func @vector_scalable_memcopy(%src : memref, %dst : memref, % %c4 = arith.constant 4 : index %vs = vector.vscale %step = arith.muli %c4, %vs : index - // CHECK: [[SRCMRS:%[0-9]+]] = builtin.unrealized_conversion_cast [[SRC]] : memref to !llvm.struct<(ptr - // CHECK: [[DSTMRS:%[0-9]+]] = builtin.unrealized_conversion_cast [[DST]] : memref to !llvm.struct<(ptr + // CHECK-DAG: [[SRCMRS:%[0-9]+]] = builtin.unrealized_conversion_cast [[SRC]] : memref to !llvm.struct<(ptr + // CHECK-DAG: [[DSTMRS:%[0-9]+]] = builtin.unrealized_conversion_cast [[DST]] : memref to !llvm.struct<(ptr // CHECK: scf.for [[LOOPIDX:%arg[0-9]+]] = {{.*}} scf.for %i0 = %c0 to %size step %step { // CHECK: [[DATAIDX:%[0-9]+]] = builtin.unrealized_conversion_cast [[LOOPIDX]] : index to i64 diff --git a/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir b/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir index 54dcf07053906f..bf4281ebcdec94 100644 --- a/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir +++ b/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir @@ -1262,8 +1262,8 @@ func.func @insert_strided_slice_scalable(%arg0 : vector<1x1x[4]xi32>, %arg1: vec // CHECK-SAME: %[[ARG_0:.*]]: vector<1x1x[4]xi32>, // CHECK-SAME: %[[ARG_1:.*]]: vector<1x4x[4]xi32>) -> vector<1x4x[4]xi32> { -// CHECK: %[[CAST_1:.*]] = builtin.unrealized_conversion_cast %[[ARG_0]] : vector<1x1x[4]xi32> to !llvm.array<1 x array<1 x vector<[4]xi32>>> -// CHECK: %[[CAST_2:.*]] = builtin.unrealized_conversion_cast %[[ARG_1]] : vector<1x4x[4]xi32> to !llvm.array<1 x array<4 x vector<[4]xi32>>> +// CHECK-DAG: %[[CAST_1:.*]] = builtin.unrealized_conversion_cast %[[ARG_0]] : vector<1x1x[4]xi32> to !llvm.array<1 x array<1 x vector<[4]xi32>>> +// CHECK-DAG: %[[CAST_2:.*]] = builtin.unrealized_conversion_cast %[[ARG_1]] : vector<1x4x[4]xi32> to !llvm.array<1 x array<4 x vector<[4]xi32>>> // CHECK: %[[EXT_1:.*]] = llvm.extractvalue %[[CAST_2]][0] : !llvm.array<1 x array<4 x vector<[4]xi32>>> // CHECK: %[[EXT_2:.*]] = llvm.extractvalue %[[CAST_1]][0, 0] : !llvm.array<1 x array<1 x vector<[4]xi32>>> @@ -2491,8 +2491,8 @@ func.func @make_fixed_vector_of_scalable_vector(%f : f64) -> vector<3x[2]xf64> // CHECK-LABEL: @vector_interleave_0d // CHECK-SAME: %[[LHS:.*]]: vector, %[[RHS:.*]]: vector) func.func @vector_interleave_0d(%a: vector, %b: vector) -> vector<2xi8> { - // CHECK: %[[LHS_RANK1:.*]] = builtin.unrealized_conversion_cast %[[LHS]] : vector to vector<1xi8> - // CHECK: %[[RHS_RANK1:.*]] = builtin.unrealized_conversion_cast %[[RHS]] : vector to vector<1xi8> + // CHECK-DAG: %[[LHS_RANK1:.*]] = builtin.unrealized_conversion_cast %[[LHS]] : vector to vector<1xi8> + // CHECK-DAG: %[[RHS_RANK1:.*]] = builtin.unrealized_conversion_cast %[[RHS]] : vector to vector<1xi8> // CHECK: %[[ZIP:.*]] = llvm.shufflevector %[[LHS_RANK1]], %[[RHS_RANK1]] [0, 1] : vector<1xi8> // CHECK: return %[[ZIP]] %0 = vector.interleave %a, %b : vector -> vector<2xi8> diff --git a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir index 6c6a9a1d0c6c5a..0d67851dfe41de 100644 --- a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir +++ b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir @@ -401,8 +401,8 @@ func.func @splat_size1_vector(%f : f32) -> vector<1xf32> { // CHECK-LABEL: func @shuffle // CHECK-SAME: %[[ARG0:.+]]: vector<1xf32>, %[[ARG1:.+]]: vector<1xf32> -// CHECK: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] -// CHECK: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] +// CHECK-DAG: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] +// CHECK-DAG: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] // CHECK: spirv.CompositeConstruct %[[V0]], %[[V1]], %[[V1]], %[[V0]] : (f32, f32, f32, f32) -> vector<4xf32> func.func @shuffle(%v0 : vector<1xf32>, %v1: vector<1xf32>) -> vector<4xf32> { %shuffle = vector.shuffle %v0, %v1 [0, 1, 1, 0] : vector<1xf32>, vector<1xf32> @@ -413,8 +413,8 @@ func.func @shuffle(%v0 : vector<1xf32>, %v1: vector<1xf32>) -> vector<4xf32> { // CHECK-LABEL: func @shuffle_index_vector // CHECK-SAME: %[[ARG0:.+]]: vector<1xindex>, %[[ARG1:.+]]: vector<1xindex> -// CHECK: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] -// CHECK: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] +// CHECK-DAG: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] +// CHECK-DAG: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] // CHECK: spirv.CompositeConstruct %[[V0]], %[[V1]], %[[V1]], %[[V0]] : (i32, i32, i32, i32) -> vector<4xi32> func.func @shuffle_index_vector(%v0 : vector<1xindex>, %v1: vector<1xindex>) -> vector<4xindex> { %shuffle = vector.shuffle %v0, %v1 [0, 1, 1, 0] : vector<1xindex>, vector<1xindex> @@ -472,8 +472,8 @@ func.func @shuffle(%v0 : vector<3xi32>, %v1: vector<1xi32>) -> vector<3xi32> { // CHECK-LABEL: func @shuffle // CHECK-SAME: %[[ARG0:.+]]: vector<1xi32>, %[[ARG1:.+]]: vector<1xi32> -// CHECK: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : vector<1xi32> to i32 -// CHECK: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] : vector<1xi32> to i32 +// CHECK-DAG: %[[V0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : vector<1xi32> to i32 +// CHECK-DAG: %[[V1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] : vector<1xi32> to i32 // CHECK: %[[RES:.+]] = spirv.CompositeConstruct %[[V0]], %[[V1]] : (i32, i32) -> vector<2xi32> // CHECK: return %[[RES]] func.func @shuffle(%v0 : vector<1xi32>, %v1: vector<1xi32>) -> vector<2xi32> { @@ -496,8 +496,8 @@ func.func @interleave(%a: vector<2xf32>, %b: vector<2xf32>) -> vector<4xf32> { // CHECK-LABEL: func @interleave_size1 // CHECK-SAME: (%[[ARG0:.+]]: vector<1xf32>, %[[ARG1:.+]]: vector<1xf32>) -// CHECK: %[[V0:.*]] = builtin.unrealized_conversion_cast %[[ARG0]] : vector<1xf32> to f32 -// CHECK: %[[V1:.*]] = builtin.unrealized_conversion_cast %[[ARG1]] : vector<1xf32> to f32 +// CHECK-DAG: %[[V0:.*]] = builtin.unrealized_conversion_cast %[[ARG0]] : vector<1xf32> to f32 +// CHECK-DAG: %[[V1:.*]] = builtin.unrealized_conversion_cast %[[ARG1]] : vector<1xf32> to f32 // CHECK: %[[RES:.*]] = spirv.CompositeConstruct %[[V0]], %[[V1]] : (f32, f32) -> vector<2xf32> // CHECK: return %[[RES]] func.func @interleave_size1(%a: vector<1xf32>, %b: vector<1xf32>) -> vector<2xf32> { diff --git a/mlir/test/Dialect/SCF/bufferize.mlir b/mlir/test/Dialect/SCF/bufferize.mlir index 6193101e9264d7..ff1612310255a0 100644 --- a/mlir/test/Dialect/SCF/bufferize.mlir +++ b/mlir/test/Dialect/SCF/bufferize.mlir @@ -4,8 +4,8 @@ // CHECK-SAME: %[[PRED:.*]]: i1, // CHECK-SAME: %[[TRUE_TENSOR:.*]]: tensor, // CHECK-SAME: %[[FALSE_TENSOR:.*]]: tensor) -> tensor { -// CHECK: %[[TRUE_MEMREF:.*]] = bufferization.to_memref %[[TRUE_TENSOR]] : memref -// CHECK: %[[FALSE_MEMREF:.*]] = bufferization.to_memref %[[FALSE_TENSOR]] : memref +// CHECK-DAG: %[[TRUE_MEMREF:.*]] = bufferization.to_memref %[[TRUE_TENSOR]] : memref +// CHECK-DAG: %[[FALSE_MEMREF:.*]] = bufferization.to_memref %[[FALSE_TENSOR]] : memref // CHECK: %[[RESULT_MEMREF:.*]] = scf.if %[[PRED]] -> (memref) { // CHECK: scf.yield %[[TRUE_MEMREF]] : memref // CHECK: } else { diff --git a/mlir/test/Dialect/Vector/linearize.mlir b/mlir/test/Dialect/Vector/linearize.mlir index 31a59b809a74ba..ec3806ca0f204f 100644 --- a/mlir/test/Dialect/Vector/linearize.mlir +++ b/mlir/test/Dialect/Vector/linearize.mlir @@ -35,8 +35,8 @@ func.func @test_linearize(%arg0: vector<2x2xf32>) -> vector<2x2xf32> { // ALL-LABEL: test_partial_linearize // ALL-SAME: (%[[ORIG_ARG:.*]]: vector<2x2xf32>, %[[ORIG_ARG2:.*]]: vector<4x4xf32>) func.func @test_partial_linearize(%arg0: vector<2x2xf32>, %arg1: vector<4x4xf32>) -> vector<2x2xf32> { - // DEFAULT: %[[ARG:.*]] = vector.shape_cast %[[ORIG_ARG]] : vector<2x2xf32> to vector<4xf32> - // DEFAULT: %[[ARG2:.*]] = vector.shape_cast %[[ORIG_ARG2]] : vector<4x4xf32> to vector<16xf32> + // DEFAULT-DAG: %[[ARG:.*]] = vector.shape_cast %[[ORIG_ARG]] : vector<2x2xf32> to vector<4xf32> + // DEFAULT-DAG: %[[ARG2:.*]] = vector.shape_cast %[[ORIG_ARG2]] : vector<4x4xf32> to vector<16xf32> // DEFAULT: %[[CST:.*]] = arith.constant dense<{{.*}}> : vector<4xf32> // DEFAULT: %[[RES:.*]] = vector.shape_cast %[[CST]] : vector<4xf32> to vector<2x2xf32> @@ -204,15 +204,15 @@ func.func @test_extract_strided_slice_2(%arg0 : vector<2x8x2xf32>) -> vector<1x4 // ALL-LABEL: test_vector_shuffle // ALL-SAME: (%[[ORIG_ARG0:.*]]: vector<4x2xf32>, %[[ORIG_ARG1:.*]]: vector<4x2xf32>) -> vector<8x2xf32> { func.func @test_vector_shuffle(%arg0: vector<4x2xf32>, %arg1: vector<4x2xf32>) -> vector<8x2xf32> { - // DEFAULT: %[[ARG0:.*]] = vector.shape_cast %[[ORIG_ARG0]] : vector<4x2xf32> to vector<8xf32> - // DEFAULT: %[[ARG1:.*]] = vector.shape_cast %[[ORIG_ARG1]] : vector<4x2xf32> to vector<8xf32> + // DEFAULT-DAG: %[[ARG0:.*]] = vector.shape_cast %[[ORIG_ARG0]] : vector<4x2xf32> to vector<8xf32> + // DEFAULT-DAG: %[[ARG1:.*]] = vector.shape_cast %[[ORIG_ARG1]] : vector<4x2xf32> to vector<8xf32> // DEFAULT: %[[SHUFFLE:.*]] = vector.shuffle %[[ARG0]], %[[ARG1]] // DEFAULT-SAME: [0, 1, 8, 9, 2, 3, 10, 11, 4, 5, 12, 13, 6, 7, 14, 15] : vector<8xf32>, vector<8xf32> // DEFAULT: %[[RES:.*]] = vector.shape_cast %[[SHUFFLE]] : vector<16xf32> to vector<8x2xf32> // DEFAULT: return %[[RES]] : vector<8x2xf32> - // BW-128: %[[ARG0:.*]] = vector.shape_cast %[[ORIG_ARG0]] : vector<4x2xf32> to vector<8xf32> - // BW-128: %[[ARG1:.*]] = vector.shape_cast %[[ORIG_ARG1]] : vector<4x2xf32> to vector<8xf32> + // BW-128-DAG: %[[ARG0:.*]] = vector.shape_cast %[[ORIG_ARG0]] : vector<4x2xf32> to vector<8xf32> + // BW-128-DAG: %[[ARG1:.*]] = vector.shape_cast %[[ORIG_ARG1]] : vector<4x2xf32> to vector<8xf32> // BW-128: %[[SHUFFLE:.*]] = vector.shuffle %[[ARG0]], %[[ARG1]] // BW-128-SAME: [0, 1, 8, 9, 2, 3, 10, 11, 4, 5, 12, 13, 6, 7, 14, 15] : vector<8xf32>, vector<8xf32> // BW-128: %[[RES:.*]] = vector.shape_cast %[[SHUFFLE]] : vector<16xf32> to vector<8x2xf32> @@ -250,8 +250,8 @@ func.func @test_vector_extract(%arg0: vector<2x8x2xf32>) -> vector<8x2xf32> { // ALL-LABEL: test_vector_insert // ALL-SAME: (%[[DEST:.*]]: vector<2x8x4xf32>, %[[SRC:.*]]: vector<8x4xf32>) -> vector<2x8x4xf32> { func.func @test_vector_insert(%arg0: vector<2x8x4xf32>, %arg1: vector<8x4xf32>) -> vector<2x8x4xf32> { - // DEFAULT: %[[ARG_SRC:.*]] = vector.shape_cast %[[SRC]] : vector<8x4xf32> to vector<32xf32> - // DEFAULT: %[[ARG_DEST:.*]] = vector.shape_cast %[[DEST]] : vector<2x8x4xf32> to vector<64xf32> + // DEFAULT-DAG: %[[ARG_SRC:.*]] = vector.shape_cast %[[SRC]] : vector<8x4xf32> to vector<32xf32> + // DEFAULT-DAG: %[[ARG_DEST:.*]] = vector.shape_cast %[[DEST]] : vector<2x8x4xf32> to vector<64xf32> // DEFAULT: %[[SHUFFLE:.*]] = vector.shuffle %[[ARG_DEST]], %[[ARG_SRC]] // DEFAULT-SAME: [64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, // DEFAULT-SAME: 88, 89, 90, 91, 92, 93, 94, 95, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, @@ -259,8 +259,8 @@ func.func @test_vector_insert(%arg0: vector<2x8x4xf32>, %arg1: vector<8x4xf32>) // DEFAULT: %[[RES:.*]] = vector.shape_cast %[[SHUFFLE]] : vector<64xf32> to vector<2x8x4xf32> // DEFAULT: return %[[RES]] : vector<2x8x4xf32> - // BW-128: %[[ARG_SRC:.*]] = vector.shape_cast %[[SRC]] : vector<8x4xf32> to vector<32xf32> - // BW-128: %[[ARG_DEST:.*]] = vector.shape_cast %[[DEST]] : vector<2x8x4xf32> to vector<64xf32> + // BW-128-DAG: %[[ARG_SRC:.*]] = vector.shape_cast %[[SRC]] : vector<8x4xf32> to vector<32xf32> + // BW-128-DAG: %[[ARG_DEST:.*]] = vector.shape_cast %[[DEST]] : vector<2x8x4xf32> to vector<64xf32> // BW-128: %[[SHUFFLE:.*]] = vector.shuffle %[[ARG_DEST]], %[[ARG_SRC]] // BW-128-SAME: [64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, // BW-128-SAME: 88, 89, 90, 91, 92, 93, 94, 95, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48,