| // RUN: iree-run-mlir --Xcompiler,iree-input-type=stablehlo --Xcompiler,iree-hal-target-backends=vmvx %s --input=1x128x128x1xf32 | FileCheck %s |
| // RUN: iree-run-mlir --Xcompiler,iree-input-type=stablehlo --Xcompiler,iree-hal-target-backends=llvm-cpu %s --input=1x128x128x1xf32 | FileCheck %s |
| |
| // Image edge detection module generated by. |
| // https://github.com/iree-org/iree/blob/main/samples/colab/tensorflow_edge_detection.ipynb. |
| // |
| // Input : a single 128x128 pixel image as a tensor<1x128x128x1xf32>, with pixels in [0.0, 1.0] |
| // Output: a single image in the same format after running edge detection |
| |
| module { |
| // CHECK-LABEL: EXEC @edge_detect_sobel_operator |
| func.func @edge_detect_sobel_operator(%arg0: tensor<1x128x128x1xf32>) -> tensor<1x128x128x1xf32> { |
| %0 = stablehlo.constant dense<[[[[-1.000000e+00]], [[0.000000e+00]], [[1.000000e+00]]], [[[-2.000000e+00]], [[0.000000e+00]], [[2.000000e+00]]], [[[-1.000000e+00]], [[0.000000e+00]], [[1.000000e+00]]]]> : tensor<3x3x1x1xf32> |
| %1 = stablehlo.constant dense<[[[[1.000000e+00]], [[2.000000e+00]], [[1.000000e+00]]], [[[0.000000e+00]], [[0.000000e+00]], [[0.000000e+00]]], [[[-1.000000e+00]], [[-2.000000e+00]], [[-1.000000e+00]]]]> : tensor<3x3x1x1xf32> |
| %2 = stablehlo.convolution(%arg0, %0) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {stride = [1, 1], pad = [[1, 1], [1, 1]], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x128x128x1xf32>, tensor<3x3x1x1xf32>) -> tensor<1x128x128x1xf32> |
| %3 = stablehlo.multiply %2, %2 : tensor<1x128x128x1xf32> |
| %4 = stablehlo.convolution(%arg0, %1) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {stride = [1, 1], pad = [[1, 1], [1, 1]], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x128x128x1xf32>, tensor<3x3x1x1xf32>) -> tensor<1x128x128x1xf32> |
| %5 = stablehlo.multiply %4, %4 : tensor<1x128x128x1xf32> |
| %6 = stablehlo.add %3, %5 : tensor<1x128x128x1xf32> |
| %7 = stablehlo.sqrt %6 : tensor<1x128x128x1xf32> |
| return %7 : tensor<1x128x128x1xf32> |
| } |
| // CHECK: 1x128x128x1xf32= |
| } |
