| #compilation0 = #iree_codegen.compilation_info< |
| lowering_config = <tile_sizes = [[32, 32], [8, 8, 0], [0, 0, 8]]>, |
| translation_info = <CPUDoubleTilingPadExpert>, |
| workgroup_size = []> |
| #compilation1 = #iree_codegen.compilation_info< |
| lowering_config = <tile_sizes = [[64, 64], [4, 4, 0], [0, 0, 4]]>, |
| translation_info = <CPUDoubleTilingPadExpert>, |
| workgroup_size = []> |
| #compilation2 = #iree_codegen.compilation_info< |
| lowering_config = <tile_sizes = [[32, 64], [8, 32, 0], [0, 0, 8]], tile_interchange = [[1, 0], [], []]>, |
| translation_info = <CPUDoubleTilingPadExpert>, |
| workgroup_size = []> |
| |
| func.func @lowering_config_test() { |
| %a = util.unfoldable_constant dense<1.0> : tensor<128x256xf32> |
| %b = util.unfoldable_constant dense<2.0> : tensor<256x512xf32> |
| %c = util.unfoldable_constant dense<2.0> : tensor<256x1024xf32> |
| %0 = "mhlo.dot"(%a, %b) {compilation_info = #compilation0} : (tensor<128x256xf32>, tensor<256x512xf32>) -> tensor<128x512xf32> |
| %1 = "mhlo.dot"(%a, %c) {compilation_info = #compilation1} : (tensor<128x256xf32>, tensor<256x1024xf32>) -> tensor<128x1024xf32> |
| %2 = "mhlo.dot"(%a, %c) {compilation_info = #compilation2} : (tensor<128x256xf32>, tensor<256x1024xf32>) -> tensor<128x1024xf32> |
| check.expect_almost_eq_const(%0, dense<512.0> : tensor<128x512xf32>) : tensor<128x512xf32> |
| check.expect_almost_eq_const(%1, dense<512.0> : tensor<128x1024xf32>) : tensor<128x1024xf32> |
| check.expect_almost_eq_const(%2, dense<512.0> : tensor<128x1024xf32>) : tensor<128x1024xf32> |
| return |
| } |
| |
| // Conv dims: N, OH, OW, OC, KH, KW, (IC) |
| // Remove H |
| #conv_compilation0 = #iree_codegen.compilation_info< |
| lowering_config = <tile_sizes = [[0, 7, 7, 64, 0, 0, 0], [6, 1, 7, 32, 0, 0, 0], [0, 0, 0, 0, 1, 3, 4]]>, |
| translation_info = <CPUConvTileAndDecomposeExpert>, |
| workgroup_size = []> |
| // Remove W |
| #conv_compilation1 = #iree_codegen.compilation_info< |
| lowering_config = <tile_sizes = [[0, 7, 7, 64, 0, 0, 0], [6, 7, 1, 32, 0, 0, 0], [0, 0, 0, 0, 3, 1, 4]]>, |
| translation_info = <CPUConvTileAndDecomposeExpert>, |
| workgroup_size = []> |
| func.func @conv() { |
| %input = util.unfoldable_constant dense<1.0> : tensor<36x7x7x512xf32> |
| %filter = util.unfoldable_constant dense<1.0> : tensor<3x3x512x512xf32> |
| %0 = "mhlo.convolution"(%input, %filter) { |
| compilation_info = #conv_compilation0, |
| batch_group_count = 1 : i64, |
| dimension_numbers = #mhlo.conv<raw input_batch_dimension = 0, input_feature_dimension = 3, input_spatial_dimensions = [1, 2], kernel_input_feature_dimension = 2, kernel_output_feature_dimension = 3, kernel_spatial_dimensions = [0, 1], output_batch_dimension = 0, output_feature_dimension = 3, output_spatial_dimensions = [1, 2]>, |
| feature_group_count = 1 : i64, |
| padding = dense<1> : tensor<2x2xi64>, |
| rhs_dilation = dense<1> : tensor<2xi64>, |
| window_strides = dense<1> : tensor<2xi64> |
| } : (tensor<36x7x7x512xf32>, tensor<3x3x512x512xf32>) -> tensor<36x7x7x512xf32> |
| %1 = "mhlo.convolution"(%input, %filter) { |
| compilation_info = #conv_compilation1, |
| batch_group_count = 1 : i64, |
| dimension_numbers = #mhlo.conv<raw input_batch_dimension = 0, input_feature_dimension = 3, input_spatial_dimensions = [1, 2], kernel_input_feature_dimension = 2, kernel_output_feature_dimension = 3, kernel_spatial_dimensions = [0, 1], output_batch_dimension = 0, output_feature_dimension = 3, output_spatial_dimensions = [1, 2]>, |
| feature_group_count = 1 : i64, |
| padding = dense<1> : tensor<2x2xi64>, |
| rhs_dilation = dense<1> : tensor<2xi64>, |
| window_strides = dense<1> : tensor<2xi64> |
| } : (tensor<36x7x7x512xf32>, tensor<3x3x512x512xf32>) -> tensor<36x7x7x512xf32> |
| check.expect_almost_eq(%0, %1) : tensor<36x7x7x512xf32> |
| return |
| } |
