Skip to content

Integrate LLVM to llvm/llvm-project@e46edf0c1fe2 #24855

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
AGindinson merged 2 commits into from
Aug 26, 2026
Merged

Integrate LLVM to llvm/llvm-project@e46edf0c1fe2 #24855

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
30d2118
Integrate LLVM to llvm/llvm-project@e46edf0c1fe2
AGindinson Aug 25, 2026
fb47ee3
[fixup] Update IR syntax in GPU/ROCDL tests
AGindinson Aug 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 5 additions & 5 deletions ...ler/src/iree/compiler/Codegen/Dialect/GPU/TransformExtensions/test/lower_inner_tiled.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ module attributes { transform.with_named_sequence } {
// 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>

// -----

Expand Down Expand Up @@ -64,7 +64,7 @@ module attributes { transform.with_named_sequence } {
// 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>

// -----

Expand Down Expand Up @@ -98,7 +98,7 @@ module attributes { transform.with_named_sequence } {
// 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>

// -----

Expand Down Expand Up @@ -483,7 +483,7 @@ module attributes { transform.with_named_sequence } {
// 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>
Expand Down Expand Up @@ -618,7 +618,7 @@ module attributes { transform.with_named_sequence } {
// 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>

// -----
Expand Down
39 changes: 19 additions & 20 deletions compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_vector_distribute_gfx942.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ func.func @matmul_256x256x256_f16_f32() attributes {hal.executable.target = #exe
// 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
Expand Down Expand Up @@ -69,7 +69,7 @@ func.func @matmul_256x256x256_f16_f32() attributes {hal.executable.target = #exe

// 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
Expand Down Expand Up @@ -126,7 +126,7 @@ func.func @expanded_matmul_transpose_b() attributes {hal.executable.target = #ex
// 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>>
Expand Down Expand Up @@ -206,7 +206,7 @@ func.func @matmul_256x256x256_16x16x32_f8_f32() attributes {hal.executable.targe
// 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

Expand Down Expand Up @@ -243,7 +243,7 @@ func.func @matmul_256x256x256_i8_i32() attributes {hal.executable.target = #exec
// 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

Expand Down Expand Up @@ -280,7 +280,7 @@ func.func @matmul_256x256x256_32x32x16_f8_f32() attributes {hal.executable.targe
// 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

Expand Down Expand Up @@ -324,7 +324,7 @@ func.func @matmul_transpose_b_256x256x256_i8_i32() attributes {hal.executable.ta
// 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

Expand Down Expand Up @@ -356,7 +356,7 @@ func.func @conv_nhwc() attributes {hal.executable.target = #executable_target_ro

// 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>>
Expand Down Expand Up @@ -412,7 +412,7 @@ func.func @generic_2x1024x20x64x1280_f16() attributes {hal.executable.target = #
// 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>>
Expand Down Expand Up @@ -463,7 +463,7 @@ func.func @contract_schedule_considering_read_layout() attributes {hal.executabl
// 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
Expand Down Expand Up @@ -506,11 +506,11 @@ func.func @virtual_intrinsic_256x256x256_16x16x32xf8E4M3FNUZ_f32() attributes {h
// 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.
Expand Down Expand Up @@ -555,7 +555,7 @@ func.func @virtual_intrinsic_256x256x256_32x32x16_f8E4M3FNUZ_f32() attributes {h
// 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
Expand Down Expand Up @@ -614,8 +614,7 @@ func.func @virtual_intrinsic_256x256x256_16x16x32xf8E4M3FNUZ_f32() attributes {h

// 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
Expand Down Expand Up @@ -740,7 +739,7 @@ func.func @attention_20x4096x64x4096x64() attributes {hal.executable.target = #e
// 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

Expand Down Expand Up @@ -801,7 +800,7 @@ func.func @attention_multiple_m_transpose() attributes {hal.executable.target =

// 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

Expand Down Expand Up @@ -862,7 +861,7 @@ func.func @attention_mfma_32x32x8() attributes {hal.executable.target = #executa

// 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

Expand Down Expand Up @@ -932,7 +931,7 @@ func.func @online_attention_split_k2() attributes {hal.executable.target = #exec
// 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

Expand Down Expand Up @@ -1079,7 +1078,7 @@ func.func @matmul_map_store() attributes {hal.executable.target = #executable_ta
// 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
Loading
Loading