| // RUN: iree-opt --split-input-file --iree-gpu-test-target=gfx942 \ |
| // RUN: --pass-pipeline="builtin.module(hal.executable(hal.executable.variant(iree-hal-configure-target-executable-variants{target=rocm})))" \ |
| // RUN: --iree-codegen-tuning-spec-path=%p/tuning_spec_mmt_tile_and_fuse.mlir \ |
| // RUN: --iree-codegen-notify-transform-strategy-application \ |
| // RUN: --verify-diagnostics %s | FileCheck %s |
| |
| // RUN: iree-opt --split-input-file --iree-gpu-test-target=gfx942 \ |
| // RUN: --pass-pipeline="builtin.module(hal.executable(hal.executable.variant(iree-hal-configure-target-executable-variants{target=rocm})))" \ |
| // RUN: --iree-codegen-tuning-spec-path=%p/tuning_spec_mmt_tile_and_fuse.mlir \ |
| // RUN: --iree-codegen-enable-default-tuning-specs \ |
| // RUN: --iree-codegen-notify-transform-strategy-application \ |
| // RUN: --verify-diagnostics %s | FileCheck %s |
| |
| // Make sure we can apply the lowering strategy from the specified tuning spec. |
| |
| // CHECK: #translation = #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse workgroup_size = [128, 1, 1] subgroup_size = 64> |
| // CHECK: func.func @matmul_transpose_b |
| // CHECK-SAME: translation_info = #translation |
| // CHECK: linalg.generic |
| // CHECK-SAME: __custom_tuning_spec_applied__ |
| // CHECK-SAME: lowering_config = #iree_gpu.lowering_config< |
| |
| #pipeline_layout = #hal.pipeline.layout<bindings = [ |
| #hal.pipeline.binding<storage_buffer>, |
| #hal.pipeline.binding<storage_buffer>, |
| #hal.pipeline.binding<storage_buffer> |
| ]> |
| hal.executable public @main { |
| hal.executable.variant public @rocm_hsaco_fb target(<"rocm", "rocm-hsaco-fb">) { |
| hal.executable.export public @matmul_transpose_b ordinal(0) layout(#pipeline_layout) { |
| ^bb0(%arg0: !hal.device): |
| %x, %y, %z = flow.dispatch.workgroup_count_from_slice |
| hal.return %x, %y, %z : index, index, index |
| } |
| builtin.module { |
| // expected-remark@+1 {{Applied transform configuration strategy @iree_linked_tuning_spec::@__kernel_config}} |
| func.func @matmul_transpose_b() { |
| %cst = arith.constant 0.000000e+00 : f16 |
| %c0 = arith.constant 0 : index |
| %0 = hal.interface.binding.subspan layout(#pipeline_layout) binding(0) alignment(64) offset(%c0) flags(ReadOnly) : !flow.dispatch.tensor<readonly:tensor<2048x1280xf16>> |
| %1 = hal.interface.binding.subspan layout(#pipeline_layout) binding(1) alignment(64) offset(%c0) flags(ReadOnly) : !flow.dispatch.tensor<readonly:tensor<10240x1280xf16>> |
| %2 = hal.interface.binding.subspan layout(#pipeline_layout) binding(2) alignment(64) offset(%c0) : !flow.dispatch.tensor<writeonly:tensor<2048x10240xf32>> |
| %3 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [2048, 1280], strides = [1, 1] : !flow.dispatch.tensor<readonly:tensor<2048x1280xf16>> -> tensor<2048x1280xf16> |
| %4 = flow.dispatch.tensor.load %1, offsets = [0, 0], sizes = [10240, 1280], strides = [1, 1] : !flow.dispatch.tensor<readonly:tensor<10240x1280xf16>> -> tensor<10240x1280xf16> |
| %5 = tensor.empty() : tensor<2048x10240xf32> |
| %6 = linalg.fill ins(%cst : f16) outs(%5 : tensor<2048x10240xf32>) -> tensor<2048x10240xf32> |
| %7 = linalg.matmul_transpose_b |
| ins(%3, %4 : tensor<2048x1280xf16>, tensor<10240x1280xf16>) |
| outs(%6 : tensor<2048x10240xf32>) -> tensor<2048x10240xf32> |
| flow.dispatch.tensor.store %7, %2, offsets = [0, 0], sizes = [2048, 10240], strides = [1, 1] : tensor<2048x10240xf32> -> !flow.dispatch.tensor<writeonly:tensor<2048x10240xf32>> |
| return |
| } |
| } |
| } |
| } |
