Integrate LLVM to llvm/llvm-project@e46edf0c1fe2 (#24855)
Still carries a temporary revert of llvm/llvm-project@202ece62657c from
`sm-iree-integrates/llvm-20260821`. Should be possible to drop in a
subsequent integrate via llvm/llvm-project#218693.
Aligns LIT CHECKs in a few GPU/ROCDL tests with syntax updates in
https://github.com/llvm/llvm-project/commit/80dfaec65.
Assisted-by: Codex
---------
Signed-off-by: Artem Gindinson <gindinson@roofline.ai>
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/lower_inner_tiled.mlir b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/lower_inner_tiled.mlir
index 07a8ae5..c198b15 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/lower_inner_tiled.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/lower_inner_tiled.mlir
@@ -30,7 +30,7 @@
// CHECK-SAME: %[[RHS:[A-Za-z0-9]+]]: vector<4xf16>
// CHECK-SAME: %[[ACC:[A-Za-z0-9]+]]: vector<4xf32>
// CHECK: amdgpu.mfma 16x16x16 %[[LHS]] * %[[RHS]] + %[[ACC]]
-// CHECK-SAME: blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<4xf32>
// -----
@@ -64,7 +64,7 @@
// CHECK-SAME: %[[RHS:[A-Za-z0-9]+]]: vector<4xf16>
// CHECK-SAME: %[[ACC:[A-Za-z0-9]+]]: vector<16xf32>
// CHECK: amdgpu.mfma 32x32x8 %[[LHS]] * %[[RHS]] + %[[ACC]]
-// CHECK-SAME: blgp = none : vector<4xf16>, vector<4xf16>, vector<16xf32>
+// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<16xf32>
// -----
@@ -98,7 +98,7 @@
// CHECK-SAME: %[[RHS:[A-Za-z0-9]+]]: vector<4xf16>
// CHECK-SAME: %[[ACC:[A-Za-z0-9]+]]: vector<16xf32>
// CHECK: amdgpu.mfma 32x32x8 %[[RHS]] * %[[LHS]] + %[[ACC]]
-// CHECK-SAME: blgp = none : vector<4xf16>, vector<4xf16>, vector<16xf32>
+// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<16xf32>
// -----
@@ -483,7 +483,7 @@
// CHECK-SAME: (%[[ACC_B:.+]] = %[[ACC]]) : (vector<4x1xf32>) -> vector<4xf32>
// CHECK: vector.shape_cast %[[ACC_B]] : vector<4x1xf32> to vector<4xf32>
// CHECK: %[[MMA:.+]] = amdgpu.mfma 16x16x16 %[[LHSCAST]] * %[[RHSCAST]] + %[[ACCCAST]]
-// CHECK-SAME: blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: util.hoistable_conversion "shape_cast_from_intrinsic" inverts("shape_cast_to_intrinsic")
// CHECK-SAME: (%[[MMA_B:.+]] = %[[MMA]]) : (vector<4xf32>) -> vector<4x1xf32>
// CHECK: vector.shape_cast %[[MMA_B]] : vector<4xf32> to vector<4x1xf32>
@@ -618,7 +618,7 @@
// CHECK-DAG: %[[LHS_S:.+]] = vector.extract %[[LHS]][0] : f64 from vector<1xf64>
// CHECK-DAG: %[[RHS_S:.+]] = vector.extract %[[RHS]][0] : f64 from vector<1xf64>
// CHECK-DAG: %[[ACC_S:.+]] = vector.extract %[[ACC]][0] : f64 from vector<1xf64>
-// CHECK: %[[MMA:.+]] = amdgpu.mfma 4x4x4 %[[LHS_S]] * %[[RHS_S]] + %[[ACC_S]] {blocks = 4 : i32} blgp = none : f64, f64, f64
+// CHECK: %[[MMA:.+]] = amdgpu.mfma blocks(4) 4x4x4 %[[LHS_S]] * %[[RHS_S]] + %[[ACC_S]] : f64, f64, f64
// CHECK: vector.broadcast %[[MMA]] : f64 to vector<1xf64>
// -----
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx942.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx942.mlir
index 40b8aa9..173acc4 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx942.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx942.mlir
@@ -36,7 +36,7 @@
// CHECK: scf.for {{.*}} = %c0 to %c256 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_f16_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -69,7 +69,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_f16_f32()
// CHECK: scf.for {{.*}} = %c0 to %c256 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_f16_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -126,7 +126,7 @@
// This has more than 2 iterations. So we have prefetching enabled for this case. Due to
// prefetching, we have one iteration peeled of so upper bound is 2048 - 128 = 1920.
// CHECK: scf.for {{.*}} = %c0 to %c1920 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-32: amdgpu.mfma
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true, true, true]} : vector<1x1x4x1xf32>, memref<2x10x64x64xf32, #amdgpu.address_space<fat_raw_buffer>>
@@ -206,7 +206,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_16x16x32_f8_f32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<4xf32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_16x16x32_f8_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -243,7 +243,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_i8_i32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xi8>, vector<8xi8>, vector<4xi32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xi8>, vector<8xi8>, vector<4xi32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xi32>, memref<256x256xi32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_i8_i32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -280,7 +280,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_32x32x16_f8_f32()
// Each subgroup handles 1 * 1 tiles, and for each tile we accumulate (256/16) = 16 times
// along the K dimension. So in total 16 mfma ops.
-// CHECK-COUNT-16: amdgpu.mfma 32x32x16 {{.*}} blgp = none : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<16xf32>
+// CHECK-COUNT-16: amdgpu.mfma 32x32x16 {{.*}} : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<16xf32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_32x32x16_f8_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -324,7 +324,7 @@
// CHECK-LABEL: func.func @matmul_transpose_b_256x256x256_i8_i32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xi8>, vector<8xi8>, vector<4xi32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xi8>, vector<8xi8>, vector<4xi32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xi32>, memref<256x256xi32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_transpose_b_256x256x256_i8_i32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -356,7 +356,7 @@
// CHECK-LABEL: func.func @conv_nhwc
// CHECK: scf.for {{.*}} = %c0 to %c215 step %c1 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-16: amdgpu.mfma
// CHECK-COUNT-8: vector.transfer_write {{.+}} {in_bounds = [true, true, true, true]} : vector<1x1x4x1xf32>, memref<2x256x512x256xf32, #amdgpu.address_space<fat_raw_buffer>>
@@ -412,7 +412,7 @@
// CHECK: scf.for {{.*}} = %c0 to %c1152 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-32: amdgpu.mfma
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true, true, true]} : vector<1x4x1x1xf32>, memref<2x1024x20x64xf32, #amdgpu.address_space<fat_raw_buffer>>
@@ -463,7 +463,7 @@
// CHECK-DAG: %[[LHS_SHARED:.+]] = memref.alloc() : memref<1x16x132xf16, #gpu.address_space<workgroup>>
// CHECK-DAG: memref.subview %[[LHS_SHARED]][0, 0, 0] [1, 16, 128] [1, 1, 1]
// CHECK: scf.for {{.*}} = %c0 to %c1408 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-16: amdgpu.mfma
// CHECK: iree_codegen.dispatch_config @contract_schedule_considering_read_layout workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -506,11 +506,11 @@
// CHECK: %[[B_CAST:.+]] = vector.shape_cast %{{.+}} : vector<1x8xf16> to vector<8xf16>
// CHECK: %[[A_SLICE_0:.+]] = vector.extract_strided_slice %[[A_CAST]] offsets = [0], sizes = [4], strides = [1] : vector<8xf16> to vector<4xf16>
// CHECK: %[[B_SLICE_0:.+]] = vector.extract_strided_slice %[[B_CAST]] offsets = [0], sizes = [4], strides = [1] : vector<8xf16> to vector<4xf16>
-// CHECK: %[[MFMA_0:.*]] = amdgpu.mfma 32x32x8 %[[A_SLICE_0]] * %[[B_SLICE_0]] + %[[ACC]] blgp = none
+// CHECK: %[[MFMA_0:.*]] = amdgpu.mfma 32x32x8 %[[A_SLICE_0]] * %[[B_SLICE_0]] + %[[ACC]]
// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<16xf32>
// CHECK: %[[A_SLICE_1:.+]] = vector.extract_strided_slice %[[A_CAST]] offsets = [4], sizes = [4], strides = [1] : vector<8xf16> to vector<4xf16>
// CHECK: %[[B_SLICE_1:.+]] = vector.extract_strided_slice %[[B_CAST]] offsets = [4], sizes = [4], strides = [1] : vector<8xf16> to vector<4xf16>
-// CHECK: amdgpu.mfma 32x32x8 %[[A_SLICE_1]] * %[[B_SLICE_1]] + %[[MFMA_0]] blgp = none
+// CHECK: amdgpu.mfma 32x32x8 %[[A_SLICE_1]] * %[[B_SLICE_1]] + %[[MFMA_0]]
// CHECK-SAME: : vector<4xf16>, vector<4xf16>, vector<16xf32>
// Ensure right number of instructions are being generated.
@@ -555,7 +555,7 @@
// CHECK: scf.for {{.*}} = %c0 to %c256 step %c128 iter_args({{.*}}) -> (vector<16xf32>)
// Each subgroup handles 1 * 1 tiles, and for each tile we accumulate (128 / 16) = 8 times
// along the K dimension. So in total 8 mfma ops.
-// CHECK-COUNT-8: amdgpu.mfma 32x32x16 {{.*}} blgp = none : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<16xf32>
+// CHECK-COUNT-8: amdgpu.mfma 32x32x16 {{.*}} : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<16xf32>
// CHECK: scf.yield
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @virtual_intrinsic_256x256x256_32x32x16_f8E4M3FNUZ_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -614,8 +614,7 @@
// CHECK: %[[LHS_CAST:.+]] = vector.shape_cast %[[EXTRACT_LHS]] : vector<1x8xf8E4M3FNUZ> to vector<8xf8E4M3FNUZ>
// CHECK: %[[RHS_CAST:.+]] = vector.shape_cast %[[EXTRACT_RHS]] : vector<1x8xf8E4M3FNUZ> to vector<8xf8E4M3FNUZ>
-// CHECK: amdgpu.mfma 16x16x32 %[[LHS_CAST]] * %[[RHS_CAST]] + %{{.*}} blgp = none
-// CHECK-SAME: : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<4xf32>
+// CHECK: amdgpu.mfma 16x16x32 %[[LHS_CAST]] * %[[RHS_CAST]] + %{{.*}} : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<4xf32>
// Ensure right number of instructions are being generated.
// CHECK-COUNT-3: amdgpu.mfma
@@ -740,7 +739,7 @@
// CHECK: transfer_read
// CHECK: scf.for %{{.*}} = %c0 to %c4096 step %c64
-// CHECK-COUNT-48: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-48: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @attention_20x4096x64x4096x64 workgroup_size = [128, 1, 1] subgroup_size = 64
@@ -801,7 +800,7 @@
// CHECK-LABEL: func.func @attention_multiple_m_transpose()
// CHECK: scf.for %{{.*}} = %c0 to %c4608 step %c64
-// CHECK-COUNT-96: amdgpu.mfma 16x16x16 {{.*}}blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-96: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @attention_multiple_m_transpose workgroup_size = [128, 1, 1] subgroup_size = 64
@@ -862,7 +861,7 @@
// CHECK-LABEL: func.func @attention_mfma_32x32x8()
// CHECK: scf.for %{{.*}} = %c0 to %c4608 step %c32
-// CHECK-COUNT-24: amdgpu.mfma 32x32x8 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<16xf32>
+// CHECK-COUNT-24: amdgpu.mfma 32x32x8 {{.*}} : vector<4xf16>, vector<4xf16>, vector<16xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @attention_mfma_32x32x8 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -932,7 +931,7 @@
// CHECK-LABEL: func.func @online_attention_split_k2()
// CHECK: scf.for %{{.*}} = %c0 to %c256 step %c32
// CHECK-SAME: -> (vector<1x1x1x4x1x1x1x1x1x1x1x4xf32>, vector<1x1x1x1x1x1x1x1x1xf32>, vector<1x1x1x1x1x1x1x1x1xf32>)
-// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @online_attention_split_k2 workgroup_size = [64, 1, 1] subgroup_size = 64
@@ -1079,7 +1078,7 @@
// CHECK: %[[OUTPUT_BINDING_ALIGNED:.+]] = memref.assume_alignment %[[OUTPUT_BINDING]]
// CHECK: %[[OUTPUT_BUFFER:.+]] = amdgpu.fat_raw_buffer_cast %[[OUTPUT_BINDING_ALIGNED]]
// CHECK: %[[FOR_RESULT:.+]]:4 = scf.for {{.*}} = %c0 to %c256 step %c128 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: %[[FLAT_OUTPUT_BUFFER:.+]] = memref.collapse_shape %[[OUTPUT_BUFFER]]
// CHECK-COUNT-4: vector.scatter %[[FLAT_OUTPUT_BUFFER]]{{.*}} : memref<65536xf32, #amdgpu.address_space<fat_raw_buffer>>, vector<4xindex>, vector<4xi1>, vector<4xf32>
// CHECK: iree_codegen.dispatch_config @matmul_map_store workgroup_size = [256, 1, 1] subgroup_size = 64
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx950.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx950.mlir
index ac4b4e8..3b74599 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx950.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx950.mlir
@@ -35,7 +35,7 @@
// CHECK: scf.for {{.*}} = %c0 to %c512 step %c256 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 8 times
// along the K dimension. So in total 32 mfma ops.
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xf16>, vector<8xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xf16>, vector<8xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x512_f16_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -68,7 +68,7 @@
// CHECK-LABEL: func.func @matmul_256x256x512_f16_f32()
// CHECK: scf.for {{.*}} = %c0 to %c512 step %c256 iter_args(%[[ARG:.+]] = {{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xf16>, vector<8xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xf16>, vector<8xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x512_f16_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -128,7 +128,7 @@
// This has more than 2 iterations. So we have prefetching enabled for this case. Due to
// prefetching, we have one iteration peeled of so upper bound is 2048 - 256 = 1792.
// CHECK: scf.for {{.*}} = %c0 to %c1792 step %c256 iter_args({{.*}}) -> (vector<4xf32>, vector<4xf32>, vector<4xf32>, vector<4xf32>)
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xf16>, vector<8xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xf16>, vector<8xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK-COUNT-32: amdgpu.mfma
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true, true, true]} : vector<1x1x4x1xf32>, memref<2x10x64x64xf32, #amdgpu.address_space<fat_raw_buffer>>
@@ -213,7 +213,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_16x16x128_f8_f32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 2 times
// along the K dimension. So in total 8 mfma ops.
-// CHECK-COUNT-8: amdgpu.mfma 16x16x128 {{.*}} blgp = none : vector<32xf8E4M3FN>, vector<32xf8E4M3FN>, vector<4xf32>
+// CHECK-COUNT-8: amdgpu.mfma 16x16x128 {{.*}} : vector<32xf8E4M3FN>, vector<32xf8E4M3FN>, vector<4xf32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_16x16x128_f8_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -250,7 +250,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_i8_i32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 4 times
// along the K dimension. So in total 16 mfma ops.
-// CHECK-COUNT-16: amdgpu.mfma 16x16x64 {{.*}} blgp = none : vector<16xi8>, vector<16xi8>, vector<4xi32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x64 {{.*}} : vector<16xi8>, vector<16xi8>, vector<4xi32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xi32>, memref<256x256xi32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_i8_i32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -287,7 +287,7 @@
// CHECK-LABEL: func.func @matmul_256x256x256_32x32x64_f8_f32()
// Each subgroup handles 1 * 1 tiles, and for each tile we accumulate (256/64) = 4 times
// along the K dimension. So in total 4 mfma ops.
-// CHECK-COUNT-4: amdgpu.mfma 32x32x64 {{.*}} blgp = none : vector<32xf8E4M3FN>, vector<32xf8E4M3FN>, vector<16xf32>
+// CHECK-COUNT-4: amdgpu.mfma 32x32x64 {{.*}} : vector<32xf8E4M3FN>, vector<32xf8E4M3FN>, vector<16xf32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xf32>, memref<256x256xf32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_256x256x256_32x32x64_f8_f32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -330,7 +330,7 @@
// CHECK-LABEL: func.func @matmul_transpose_b_256x256x256_i8_i32()
// Each subgroup handles 2 * 2 tiles, and for each tile we accumulate 4 times
// along the K dimension. So in total 16 mfma ops.
-// CHECK-COUNT-16: amdgpu.mfma 16x16x64 {{.*}} blgp = none : vector<16xi8>, vector<16xi8>, vector<4xi32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x64 {{.*}} : vector<16xi8>, vector<16xi8>, vector<4xi32>
// CHECK-COUNT-4: vector.transfer_write {{.+}} {in_bounds = [true, true]} : vector<4x1xi32>, memref<256x256xi32, #amdgpu.address_space<fat_raw_buffer>>
// CHECK: iree_codegen.dispatch_config @matmul_transpose_b_256x256x256_i8_i32 workgroup_size = [256, 1, 1] subgroup_size = 64
@@ -389,8 +389,8 @@
// CHECK: transfer_read
// CHECK: scf.for %{{.*}} = %c0 to %c4096 step %c128
-// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} blgp = none : vector<8xf16>, vector<8xf16>, vector<4xf32>
-// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<4xf32>
+// CHECK-COUNT-32: amdgpu.mfma 16x16x32 {{.*}} : vector<8xf16>, vector<8xf16>, vector<4xf32>
+// CHECK-COUNT-16: amdgpu.mfma 16x16x16 {{.*}} : vector<4xf16>, vector<4xf16>, vector<4xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @attention_20x4096x64x4096x64 workgroup_size = [128, 1, 1] subgroup_size = 64
@@ -451,8 +451,8 @@
// CHECK-LABEL: func.func @attention_mfma_32x32x16()
// CHECK: scf.for %{{.*}} = %c0 to %c4608 step %c64
-// CHECK-COUNT-16: amdgpu.mfma 32x32x16 {{.*}} blgp = none : vector<8xf16>, vector<8xf16>, vector<16xf32>
-// CHECK-COUNT-8: amdgpu.mfma 32x32x8 {{.*}} blgp = none : vector<4xf16>, vector<4xf16>, vector<16xf32>
+// CHECK-COUNT-16: amdgpu.mfma 32x32x16 {{.*}} : vector<8xf16>, vector<8xf16>, vector<16xf32>
+// CHECK-COUNT-8: amdgpu.mfma 32x32x8 {{.*}} : vector<4xf16>, vector<4xf16>, vector<16xf32>
// CHECK: scf.yield
// CHECK: iree_codegen.dispatch_config @attention_mfma_32x32x16 workgroup_size = [256, 1, 1] subgroup_size = 64
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
index 036f646..73c2e6b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -291,15 +291,9 @@
rocdl.s.setprio 2 { iree_gpu.swap_mfma = 1 }
rocdl.s.setprio 3 { iree_gpu.swap_mfma = 2 }
rocdl.s.setprio 4 { iree_gpu.swap_mfma = 5 }
- %0 = amdgpu.mfma 4x4x1 %in * %in + %out {
- abid = 0 : i32, cbsz = 0 : i32, blocks = 16 : i32
- } blgp = none : f32, f32, vector<4xf32>
- %1 = amdgpu.mfma 4x4x1 %in * %in + %0 {
- abid = 0 : i32, cbsz = 0 : i32, blocks = 16 : i32
- } blgp = none : f32, f32, vector<4xf32>
- %2 = amdgpu.mfma 4x4x1 %in * %in + %1 {
- abid = 0 : i32, cbsz = 0 : i32, blocks = 16 : i32
- } blgp = none : f32, f32, vector<4xf32>
+ %0 = amdgpu.mfma blocks(16) 4x4x1 %in * %in + %out : f32, f32, vector<4xf32>
+ %1 = amdgpu.mfma blocks(16) 4x4x1 %in * %in + %0 : f32, f32, vector<4xf32>
+ %2 = amdgpu.mfma blocks(16) 4x4x1 %in * %in + %1 : f32, f32, vector<4xf32>
call @foo(%2) : (vector<4xf32>) -> ()
return
}
diff --git a/third_party/llvm-project b/third_party/llvm-project
index d98828d..53e39b2 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit d98828d394e4967148f6992afc3f23662a8e17a8
+Subproject commit 53e39b225bd164efeae02c9a0721021469d5bc67
