[vulkan] Clean up various Vulkan specific e2e tests (iree-org#9672)

These tests are accumulated over time but we haven't re-check whether they are useful anymore. * `compare.mlir` is identical to the one under `xla_ops/`. * `dot_general.mlir` is a subset of the one under `xla_ops/`. * `vectorized_conv.mlir` is merged into `conv.mlir`. * `log_plus_one.mlir` is identical to the one under `xla_ops/`. * `pw_add_multiwg.mlir` was added during the early time and just checks a pointwise addition; not really useful now. * `reduce.mlir` is merged into the one under `xla_ops/`.
dan-garvey · Jun 29, 2022 · 70503f2 · 70503f2
1 parent 151b405
commit 70503f2
Show file tree

Hide file tree

Showing 10 changed files with 169 additions and 464 deletions.
diff --git a/tests/e2e/vulkan_specific/BUILD b/tests/e2e/vulkan_specific/BUILD
@@ -5,9 +5,7 @@
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
 # Tests for end-to-end IREE support specific to the vulkan-spirv lowering.
-# TODO(ravishankarm): Reorganize these tests.
 
-load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
 load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite")
 
 package(
@@ -16,30 +14,6 @@ package(
     licenses = ["notice"],  # Apache 2.0
 )
 
-iree_check_single_backend_test_suite(
-    name = "check_vulkan-spirv_vulkan",
-    srcs = enforce_glob(
-        [
-            "compare.mlir",
-            "conv.mlir",
-            "dot_general.mlir",
-            "log_plus_one.mlir",
-            "pw_add_multiwg.mlir",
-            "reduce.mlir",
-            "vectorized_conv.mlir",
-        ],
-        include = ["*.mlir"],
-        exclude = [
-            "add_f16.mlir",
-            "dot_f16.mlir",
-            "gemm.mlir",
-        ],
-    ),
-    compiler_flags = ["--iree-input-type=mhlo"],
-    driver = "vulkan",
-    target_backend = "vulkan-spirv",
-)
-
 iree_check_single_backend_test_suite(
     name = "check_vulkan-spirv_vulkan_f16",
     srcs = [
@@ -60,9 +34,9 @@ iree_check_single_backend_test_suite(
 )
 
 iree_check_single_backend_test_suite(
-    name = "check_vulkan-spirv_vulkan_vectorized_conv",
+    name = "check_vulkan-spirv_vulkan_conv",
     srcs = [
-        "vectorized_conv.mlir",
+        "conv.mlir",
     ],
     compiler_flags = [
         "--iree-input-type=mhlo",

diff --git a/tests/e2e/vulkan_specific/CMakeLists.txt b/tests/e2e/vulkan_specific/CMakeLists.txt
@@ -10,25 +10,6 @@
 
 iree_add_all_subdirs()
 
-iree_check_single_backend_test_suite(
-  NAME
-    check_vulkan-spirv_vulkan
-  SRCS
-    "compare.mlir"
-    "conv.mlir"
-    "dot_general.mlir"
-    "log_plus_one.mlir"
-    "pw_add_multiwg.mlir"
-    "reduce.mlir"
-    "vectorized_conv.mlir"
-  TARGET_BACKEND
-    "vulkan-spirv"
-  DRIVER
-    "vulkan"
-  COMPILER_FLAGS
-    "--iree-input-type=mhlo"
-)
-
 iree_check_single_backend_test_suite(
   NAME
     check_vulkan-spirv_vulkan_f16
@@ -50,9 +31,9 @@ iree_check_single_backend_test_suite(
 
 iree_check_single_backend_test_suite(
   NAME
-    check_vulkan-spirv_vulkan_vectorized_conv
+    check_vulkan-spirv_vulkan_conv
   SRCS
-    "vectorized_conv.mlir"
+    "conv.mlir"
   TARGET_BACKEND
     "vulkan-spirv"
   DRIVER

diff --git a/tests/e2e/vulkan_specific/compare.mlir b/tests/e2e/vulkan_specific/compare.mlir
diff --git a/tests/e2e/vulkan_specific/conv.mlir b/tests/e2e/vulkan_specific/conv.mlir
@@ -1,4 +1,4 @@
-func.func @conv() {
+func.func @normal_conv_1() {
   %0 = util.unfoldable_constant dense<
       [[[[0.5       , 0.5212766 ],
          [0.54255319, 0.56382979],
@@ -82,3 +82,117 @@ func.func @conv() {
         : tensor<1x3x4x3xf32>) : tensor<1x3x4x3xf32>
    return
 }
+func.func @normal_conv_2() {
+  %input = util.unfoldable_constant dense<
+     [[[[6.0, 7.5, 0.0, 1.5],
+        [1.5, 3.5, 4.5, 2.0],
+        [3.0, 6.0, 0.5, 3.0]],
+       [[3.5, 7.0, 2.5, 6.5],
+        [4.0, 4.5, 8.0, 2.5],
+        [7.5, 7.5, 0.0, 1.5]],
+       [[7.0, 3.5, 0.0, 0.5],
+        [4.5, 0.0, 5.0, 1.5],
+        [5.5, 1.0, 0.0, 0.0]]]]>
+    : tensor<1x3x3x4xf32>
+  %filter = util.unfoldable_constant dense<
+      [[[[2.0, 2.5, 2.5, 3.0, 4.0, 2.0, 0.5, 2.0, 4.5, 5.0, 5.0, 4.0, 0.5, 0.5, 3.5, 4.5,
+          4.5, 1.5, 3.0, 3.5, 1.0, 0.0, 1.5, 2.5, 4.5, 5.0, 2.0, 2.0, 3.0, 2.0, 2.0, 1.5],
+         [2.0, 2.0, 4.0, 2.0, 1.5, 5.0, 3.5, 2.5, 2.5, 0.0, 0.5, 2.5, 4.5, 1.5, 0.0, 2.5,
+          0.0, 0.5, 1.0, 2.0, 1.0, 0.0, 1.5, 1.0, 5.0, 0.0, 3.5, 2.5, 4.5, 0.0, 5.0, 1.0],
+         [5.0, 3.5, 1.0, 4.5, 1.0, 1.5, 1.5, 1.0, 1.5, 2.0, 0.5, 1.0, 4.5, 5.0, 0.5, 2.0,
+          5.0, 3.0, 4.0, 1.0, 1.5, 0.0, 0.0, 3.0, 0.0, 3.0, 1.5, 5.0, 1.5, 4.0, 4.0, 4.0],
+         [1.0, 1.5, 1.0, 0.0, 4.0, 4.0, 1.5, 4.0, 5.0, 1.0, 4.0, 2.0, 1.5, 0.0, 2.0, 1.5,
+          3.0, 4.5, 4.0, 0.0, 4.0, 2.5, 4.5, 0.0, 4.5, 3.0, 2.5, 1.5, 0.5, 4.0, 0.0, 2.0]],
+        [[4.5, 3.0, 2.5, 3.5, 4.0, 4.0, 4.5, 1.0, 4.0, 3.0, 3.0, 4.5, 0.5, 3.0, 4.0, 4.0,
+          1.5, 1.0, 1.5, 5.0, 3.0, 1.5, 3.0, 2.5, 3.5, 0.0, 4.0, 2.0, 5.0, 3.0, 2.5, 4.0],
+         [1.0, 1.5, 4.5, 3.5, 2.5, 1.5, 2.0, 2.5, 1.5, 1.5, 3.5, 4.5, 4.5, 4.5, 3.5, 1.5,
+          5.0, 1.0, 1.5, 4.5, 5.0, 3.5, 3.5, 2.5, 0.5, 1.0, 1.0, 4.0, 0.5, 2.5, 4.0, 2.0],
+         [0.0, 1.0, 2.5, 2.5, 0.0, 4.0, 0.5, 0.5, 0.0, 1.5, 4.0, 4.0, 2.0, 2.0, 0.0, 4.5,
+          1.5, 3.5, 1.5, 1.0, 0.5, 0.5, 1.0, 0.5, 2.0, 1.0, 2.5, 2.5, 2.5, 1.0, 2.5, 3.5],
+         [3.5, 3.0, 0.5, 3.0, 3.5, 1.0, 1.5, 0.5, 4.5, 2.5, 4.5, 4.5, 1.0, 0.0, 4.5, 0.5,
+          4.5, 5.0, 0.0, 3.0, 0.0, 5.0, 2.0, 4.0, 2.0, 1.5, 1.5, 4.0, 4.0, 3.5, 0.0, 1.5]]],
+       [[[4.0, 3.5, 3.5, 5.0, 0.5, 4.0, 2.0, 3.5, 0.0, 2.0, 4.5, 0.0, 5.0, 3.0, 2.0, 1.0,
+          2.0, 3.0, 1.5, 5.0, 1.5, 3.5, 4.0, 2.5, 0.0, 4.0, 2.5, 2.0, 3.5, 5.0, 5.0, 2.0],
+         [0.5, 1.5, 1.5, 4.5, 1.0, 2.5, 1.0, 1.5, 2.5, 5.0, 3.5, 1.0, 3.5, 0.5, 3.0, 5.0,
+          2.5, 0.0, 0.0, 5.0, 1.5, 5.0, 0.5, 5.0, 4.5, 4.5, 3.0, 3.0, 3.5, 4.0, 4.0, 3.5],
+         [0.0, 4.0, 3.0, 4.0, 4.5, 4.0, 1.5, 3.0, 0.5, 3.5, 2.0, 4.5, 1.0, 0.0, 4.0, 1.0,
+          3.5, 4.0, 2.0, 2.0, 0.5, 3.5, 3.0, 4.5, 2.0, 0.5, 2.5, 4.5, 3.5, 0.5, 1.5, 2.5],
+         [3.5, 1.5, 3.0, 3.0, 3.5, 4.5, 0.5, 4.5, 3.0, 0.0, 1.5, 4.0, 2.0, 0.5, 2.0, 2.5,
+          0.0, 1.5, 5.0, 0.5, 2.0, 2.0, 2.0, 0.0, 0.0, 5.0, 4.0, 2.0, 3.0, 4.5, 1.5, 1.5]],
+        [[1.0, 0.5, 5.0, 1.0, 0.5, 1.5, 2.0, 5.0, 0.5, 0.5, 0.0, 3.5, 4.0, 5.0, 2.0, 1.5,
+          2.5, 3.0, 1.5, 1.0, 4.5, 4.0, 0.5, 2.0, 5.0, 0.0, 4.0, 1.5, 4.5, 2.5, 2.5, 0.5],
+         [3.5, 4.0, 3.0, 2.0, 3.5, 1.5, 2.5, 1.5, 3.0, 2.0, 3.5, 1.5, 0.0, 2.5, 4.5, 1.5,
+          3.5, 2.5, 2.5, 4.0, 0.0, 4.0, 1.5, 3.0, 4.5, 5.0, 1.5, 1.0, 3.5, 0.0, 1.5, 5.0],
+         [0.0, 1.5, 3.0, 0.5, 4.5, 1.0, 4.5, 2.0, 4.5, 0.5, 1.5, 1.0, 2.0, 4.5, 3.5, 2.0,
+          4.5, 2.0, 0.5, 1.0, 3.5, 1.0, 1.5, 4.5, 5.0, 3.5, 5.0, 3.0, 3.0, 1.0, 5.0, 1.5],
+         [3.0, 0.0, 5.0, 4.0, 0.0, 5.0, 3.5, 3.0, 2.5, 4.5, 3.0, 2.5, 1.0, 3.5, 0.5, 4.5,
+          1.0, 1.0, 2.5, 3.0, 2.0, 1.0, 1.0, 0.5, 0.0, 4.5, 0.0, 1.0, 4.0, 1.5, 5.0, 0.0]]]]>
+    : tensor<2x2x4x32xf32>
+
+    %0 = "mhlo.convolution"(%input, %filter) {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<0> : tensor<2x2xi64>, rhs_dilation = dense<1> : tensor<2xi64>, window_strides = dense<1> : tensor<2xi64>} : (tensor<1x3x3x4xf32>, tensor<2x2x4x32xf32>) -> tensor<1x2x2x32xf32>
+
+   check.expect_almost_eq_const(%0, dense<
+     [[[[113.25, 127.0, 198.0, 173.25, 159.5, 190.75, 135.5, 160.0,
+         169.5, 130.0, 173.75, 174.5, 158.5, 136.75, 159.75, 177.75,
+         164.5, 122.25, 116.0, 168.0, 124.75, 144.0, 113.5, 159.0,
+         208.0, 186.5, 190.5, 158.5, 213.75, 140.5, 206.75, 135.25],
+        [129.75, 147.25, 181.25, 181.75, 142.5, 161.75, 117.75, 153.25,
+         119.5, 128.75, 149.25, 171.0, 152.5, 142.5, 166.0, 122.25,
+         177.75, 142.75, 116.5, 170.0, 117.5, 176.75, 116.75, 162.25,
+         161.25, 135.0, 145.5, 163.25, 190.5, 138.25, 162.5, 146.75]],
+       [[111.75, 115.75, 173.5, 158.25, 122.5, 187.25, 129.0, 142.5,
+         142.25, 109.0, 175.75, 158.5, 172.75, 146.25, 122.25, 157.25,
+         157.5, 141.25, 104.25, 151.25, 136.25, 122.0, 127.75, 125.75,
+         180.5, 131.25, 168.75, 151.5, 180.75, 152.75, 193.5, 128.75],
+        [138.25, 133.75, 157.5, 168.5, 131.0, 149.75, 115.25, 130.75,
+         114.5, 107.25, 127.75, 163.75, 153.5, 149.25, 133.5, 114.0,
+         164.75, 120.75, 116.0, 149.5, 127.5, 113.5, 116.0, 129.75,
+         126.75, 94.25, 135.0, 157.75, 158.75, 142.0, 158.75, 126.25]]]]>
+        : tensor<1x2x2x32xf32>) : tensor<1x2x2x32xf32>
+   return
+}
+
+func.func @depthwise_conv() {
+  %input = util.unfoldable_constant dense<
+    [[[[6.0, 7.5, 0.0, 1.5, 1.5, 3.5, 4.5, 2.0, 3.0, 6.0, 0.5, 3.0, 3.5, 7.0, 2.5, 6.5],
+       [4.0, 4.5, 8.0, 2.5, 7.5, 7.5, 0.0, 1.5, 7.0, 3.5, 0.0, 0.5, 4.5, 0.0, 5.0, 1.5],
+       [5.5, 1.0, 0.0, 0.0, 2.0, 2.5, 3.0, 4.0, 7.5, 2.0, 4.5, 5.0, 0.5, 0.5, 3.5, 4.5],
+       [1.5, 3.0, 5.5, 7.0, 0.0, 7.0, 1.5, 6.0, 5.0, 5.5, 2.0, 3.0, 2.0, 7.5, 1.5, 6.0]]]]>
+    : tensor<1x1x4x16xf32>
+  %filter = util.unfoldable_constant dense<
+    [[[[2.0, 2.0, 4.0, 2.0, 1.5, 5.0, 3.5, 2.5, 2.5, 0.0, 0.5, 2.5, 4.5, 1.5, 0.0, 2.5]]]]>
+    : tensor<1x1x1x16xf32>
+
+    %0 = "mhlo.convolution"(%input, %filter) {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 = 16 : i64, padding = dense<0> : tensor<2x2xi64>, rhs_dilation = dense<1> : tensor<2xi64>, window_strides = dense<1> : tensor<2xi64>} : (tensor<1x1x4x16xf32>, tensor<1x1x1x16xf32>) -> tensor<1x1x4x16xf32>
+
+   check.expect_almost_eq_const(%0, dense<
+     [[[[12.0, 15.0, 0.0, 3.0, 2.25, 17.5, 15.75, 5.0, 7.5, 0.0, 0.25, 7.5, 15.75, 10.5, 0.0, 16.25],
+        [8.0, 9.0, 32.0, 5.0, 11.25, 37.5, 0.0, 3.75, 17.5, 0.0, 0.0, 1.25, 20.25, 0.0, 0.0, 3.75],
+        [11.0, 2.0, 0.0, 0.0, 3.0, 12.5, 10.5, 10.0, 18.75, 0.0, 2.25, 12.5, 2.25, 0.75, 0.0, 11.25],
+        [3.0, 6.0, 22.0, 14.0, 0.0, 35.0, 5.25, 15.0, 12.5, 0.0, 1.0, 7.5, 9.0, 11.25, 0.0, 15.0]]]]>
+        : tensor<1x1x4x16xf32>) : tensor<1x1x4x16xf32>
+   return
+}
+