Enable compilation tests for microbenchmarks (#8904)
This creates a microbenchmark suite and the adds the compilation to CI bots.
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 261b43e..e6bd799 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -39,6 +39,7 @@
option(IREE_BUILD_COMPILER "Builds the IREE compiler." ON)
option(IREE_BUILD_TESTS "Builds IREE unit tests." ON)
option(IREE_BUILD_BENCHMARKS "Builds IREE benchmark suites." OFF)
+option(IREE_BUILD_MICROBENCHMARKS "Builds IREE microbenchmark suites." OFF)
option(IREE_BUILD_DOCS "Builds IREE docs." OFF)
option(IREE_BUILD_SAMPLES "Builds IREE sample projects." ON)
option(IREE_BUILD_TRACY "Builds tracy server tools." OFF)
@@ -239,6 +240,7 @@
include(iree_native_test)
include(iree_cc_binary_benchmark)
include(iree_benchmark_suite)
+include(iree_microbenchmark_suite)
include(iree_hal_cts_test_suite)
set(DEFAULT_CMAKE_BUILD_TYPE "Release")
@@ -574,6 +576,11 @@
add_custom_target(iree-benchmark-suites)
endif()
+if(${IREE_BUILD_MICROBENCHMARKS})
+ # Add top-level custom targets to drive generating microbenchmark suites.
+ add_custom_target(iree-microbenchmark-suites)
+endif()
+
if(${IREE_BUILD_DOCS})
# Add a top-level custom target to drive generating all documentation.
# Register it to the default target given that IREE_BUILD_DOCS is explicitly
diff --git a/build_tools/cmake/build_android_benchmark.sh b/build_tools/cmake/build_android_benchmark.sh
index c712fc3..46b219d 100755
--- a/build_tools/cmake/build_android_benchmark.sh
+++ b/build_tools/cmake/build_android_benchmark.sh
@@ -64,11 +64,13 @@
-DIREE_BUILD_COMPILER=ON \
-DIREE_BUILD_TESTS=OFF \
-DIREE_BUILD_BENCHMARKS=ON \
+ -DIREE_BUILD_MICROBENCHMARKS=ON \
-DIREE_BUILD_SAMPLES=OFF
"${CMAKE_BIN}" --build . --target install -- -k 0
# Also generate artifacts for benchmarking on Android.
"${CMAKE_BIN}" --build . --target iree-benchmark-suites -- -k 0
+"${CMAKE_BIN}" --build . --target iree-microbenchmark-suites -- -k 0
# --------------------------------------------------------------------------- #
# --------------------------------------------------------------------------- #
diff --git a/build_tools/cmake/iree_microbenchmark_suite.cmake b/build_tools/cmake/iree_microbenchmark_suite.cmake
new file mode 100644
index 0000000..4ac1021
--- /dev/null
+++ b/build_tools/cmake/iree_microbenchmark_suite.cmake
@@ -0,0 +1,61 @@
+# Copyright 2022 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+# iree_microbenchmark_suite()
+#
+# Generates microbenchmark suites for MLIR input modules.
+# Parameters:
+# NAME: Name of target.
+# SRCS: Source files to compile into a bytecode module (list of strings).
+# FLAGS: Flags to pass to the translation tool (list of strings).
+
+function(iree_microbenchmark_suite)
+ if(NOT IREE_BUILD_MICROBENCHMARKS)
+ return()
+ endif()
+
+ cmake_parse_arguments(
+ _RULE
+ ""
+ "NAME"
+ "SRCS;FLAGS"
+ ${ARGN}
+ )
+
+ iree_package_name(PACKAGE_NAME)
+
+ foreach(_SRC IN LISTS _RULE_SRCS)
+ set(_TRANSLATE_TOOL "iree-compile")
+ set(_TRANSLATE_SRC "${_SRC}")
+ set(_MODULE_FILE_NAME "${_RULE_NAME}_${_SRC}.vmfb")
+ set(_TARGET_NAME "${PACKAGE_NAME}_${_MODULE_FILE_NAME}")
+ iree_get_executable_path(_TRANSLATE_TOOL_EXECUTABLE "${_TRANSLATE_TOOL}")
+ set(_ARGS "${_RULE_FLAGS}")
+ get_filename_component(_TRANSLATE_SRC_PATH "${_TRANSLATE_SRC}" REALPATH)
+ list(APPEND _ARGS "${_TRANSLATE_SRC_PATH}")
+ list(APPEND _ARGS "-o")
+ list(APPEND _ARGS "${_MODULE_FILE_NAME}")
+
+ add_custom_command(
+ OUTPUT
+ "${_MODULE_FILE_NAME}"
+ COMMAND
+ "${_TRANSLATE_TOOL_EXECUTABLE}"
+ ${_ARGS}
+ # Changes to either the translation tool or the input source should
+ # trigger rebuilding.
+ DEPENDS
+ "${_TRANSLATE_TOOL_EXECUTABLE}"
+ "${_TRANSLATE_SRC}"
+ VERBATIM
+ )
+ add_custom_target("${_TARGET_NAME}"
+ DEPENDS
+ "${_MODULE_FILE_NAME}"
+ )
+ add_dependencies(iree-microbenchmark-suites "${_TARGET_NAME}")
+ endforeach(_SRC IN LISTS _SRCS)
+endfunction(iree_microbenchmark_suite)
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-swiftshader-asan/build.sh b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-swiftshader-asan/build.sh
index b6bdca6..ae4b10c 100755
--- a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-swiftshader-asan/build.sh
+++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-swiftshader-asan/build.sh
@@ -37,6 +37,9 @@
"-DIREE_ENABLE_ASAN=ON"
"-B" "${CMAKE_BUILD_DIR?}"
+ # Also check if microbenchmarks are buildable.
+ "-DIREE_BUILD_MICROBENCHMARKS=ON"
+
# Enable CUDA compiler and runtime builds unconditionally. Our CI images all
# have enough deps to at least build CUDA support and compile CUDA binaries
# (but not necessarily test on real hardware).
@@ -55,6 +58,10 @@
echo "------------------"
"${CMAKE_BIN?}" --build "${CMAKE_BUILD_DIR?}" --target iree-test-deps -- -k 0
+echo "Building microbenchmark suites"
+echo "------------------"
+"${CMAKE_BIN?}" --build "${CMAKE_BUILD_DIR?}" --target iree-microbenchmark-suites -- -k 0
+
# Respect the user setting, but default to as many jobs as we have cores.
export CTEST_PARALLEL_LEVEL=${CTEST_PARALLEL_LEVEL:-$(nproc)}
diff --git a/iree/test/microbenchmarks/BUILD b/iree/test/microbenchmarks/BUILD
deleted file mode 100644
index bcc8816..0000000
--- a/iree/test/microbenchmarks/BUILD
+++ /dev/null
@@ -1,41 +0,0 @@
-# Copyright 2022 The IREE Authors
-#
-# Licensed under the Apache License v2.0 with LLVM Exceptions.
-# See https://llvm.org/LICENSE.txt for license information.
-# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-# Tests compilation of microbenchmarks.
-
-load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
-load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-iree_lit_test_suite(
- name = "lit",
- size = "medium",
- srcs = enforce_glob(
- [
- "linalg_transpose.mlir",
- ],
- include = ["*.mlir"],
- exclude =
- [
- "linalg_mmt4d.mlir",
- "mhlo_conv.mlir",
- "mhlo_dot_general.mlir",
- "mhlo_dot.mlir",
- "mhlo_fft_abs.mlir",
- ],
- ),
- tags = [
- "hostonly",
- ],
- tools = [
- "//iree/tools:iree-run-mlir",
- ],
-)
diff --git a/iree/test/microbenchmarks/CMakeLists.txt b/iree/test/microbenchmarks/CMakeLists.txt
index e90f491..7dfa93c 100644
--- a/iree/test/microbenchmarks/CMakeLists.txt
+++ b/iree/test/microbenchmarks/CMakeLists.txt
@@ -1,24 +1,19 @@
-################################################################################
-# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/test/microbenchmarks/BUILD #
-# #
-# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
-# CMake-only content. #
-# #
-# To disable autogeneration for this file entirely, delete this header. #
-################################################################################
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
iree_add_all_subdirs()
-iree_lit_test_suite(
+iree_microbenchmark_suite(
NAME
- lit
+ "microbenchmark"
SRCS
+ "linalg_mmt4d.mlir"
"linalg_transpose.mlir"
- TOOLS
- iree::tools::iree-run-mlir
- LABELS
- "hostonly"
+ "mhlo_conv.mlir"
+ "mhlo_dot.mlir"
+ "mhlo_dot_general.mlir"
+ "mhlo_fft_abs.mlir"
+ FLAGS
+ "--iree-hal-target-backends=dylib-llvm-aot"
+ "--iree-input-type=mhlo"
+ "--iree-llvm-target-cpu-features=host"
)
-
-### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/iree/test/microbenchmarks/mhlo_conv.mlir b/iree/test/microbenchmarks/mhlo_conv.mlir
index f9c37c6..f0678ba 100644
--- a/iree/test/microbenchmarks/mhlo_conv.mlir
+++ b/iree/test/microbenchmarks/mhlo_conv.mlir
@@ -13,17 +13,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<3x3x3x32xf32>
%0 = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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, 1], [0, 1]]> : tensor<2x2xi64>,
rhs_dilation = dense<1> : tensor<2xi64>,
@@ -37,17 +37,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<1x1x32x64xf32>
%0 = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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>,
@@ -61,17 +61,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<1x1x1024x1024xf32>
%0 = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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>,
@@ -96,17 +96,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<15x1x1x1024xf32>
%res = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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 = 1024 : i64,
padding = dense<0> : tensor<2x2xi64>,
rhs_dilation = dense<1> : tensor<2xi64>,
@@ -120,17 +120,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<15x1x1x512xf32>
%res = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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 = 512 : i64,
padding = dense<0> : tensor<2x2xi64>,
rhs_dilation = dense<1> : tensor<2xi64>,
@@ -144,17 +144,17 @@
%filter = util.unfoldable_constant dense<1.0> : tensor<15x1x1x512xf32>
%res = "mhlo.convolution"(%input, %filter) {
batch_group_count = 1 : i64,
- dimension_numbers = {
- input_batch_dimension = 0 : i64,
- input_feature_dimension = 3 : i64,
- input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>,
- kernel_input_feature_dimension = 2 : i64,
- kernel_output_feature_dimension = 3 : i64,
- kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>,
- output_batch_dimension = 0 : i64,
- output_feature_dimension = 3 : i64,
- output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>
- },
+ 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 = 512 : i64,
padding = dense<0> : tensor<2x2xi64>,
rhs_dilation = dense<1> : tensor<2xi64>,
diff --git a/iree/test/microbenchmarks/mhlo_dot_general.mlir b/iree/test/microbenchmarks/mhlo_dot_general.mlir
index 4ca8999..343ab45 100644
--- a/iree/test/microbenchmarks/mhlo_dot_general.mlir
+++ b/iree/test/microbenchmarks/mhlo_dot_general.mlir
@@ -5,12 +5,13 @@
%lhs = util.unfoldable_constant dense<1.0> : tensor<4x384x32xf32>
%rhs = util.unfoldable_constant dense<1.0> : tensor<4x32x384xf32>
%0 = "mhlo.dot_general"(%lhs, %rhs) {
- dot_dimension_numbers = {
- lhs_batching_dimensions = dense<0> : tensor<1xi64>,
- lhs_contracting_dimensions = dense<2> : tensor<1xi64>,
- rhs_batching_dimensions = dense<0> : tensor<1xi64>,
- rhs_contracting_dimensions = dense<1> : tensor<1xi64>
- }
+ dot_dimension_numbers = #mhlo.dot<
+ lhs_batching_dimensions = [0],
+ lhs_contracting_dimensions = [2],
+ rhs_batching_dimensions = [0],
+ rhs_contracting_dimensions = [1],
+ >,
+ precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]
} : (tensor<4x384x32xf32>, tensor<4x32x384xf32>) -> tensor<4x384x384xf32>
return %0 : tensor<4x384x384xf32>
}
@@ -19,12 +20,13 @@
%lhs = util.unfoldable_constant dense<1.0> : tensor<4x384x384xf32>
%rhs = util.unfoldable_constant dense<1.0> : tensor<4x384x32xf32>
%0 = "mhlo.dot_general"(%lhs, %rhs) {
- dot_dimension_numbers = {
- lhs_batching_dimensions = dense<0> : tensor<1xi64>,
- lhs_contracting_dimensions = dense<2> : tensor<1xi64>,
- rhs_batching_dimensions = dense<0> : tensor<1xi64>,
- rhs_contracting_dimensions = dense<1> : tensor<1xi64>
- }
+ dot_dimension_numbers = #mhlo.dot<
+ lhs_batching_dimensions = [0],
+ lhs_contracting_dimensions = [2],
+ rhs_batching_dimensions = [0],
+ rhs_contracting_dimensions = [1],
+ >,
+ precision_config = [#mhlo<"precision DEFAULT">, #mhlo<"precision DEFAULT">]
} : (tensor<4x384x384xf32>, tensor<4x384x32xf32>) -> tensor<4x384x32xf32>
return %0 : tensor<4x384x32xf32>
}
diff --git a/iree/test/microbenchmarks/mhlo_fft_abs.mlir b/iree/test/microbenchmarks/mhlo_fft_abs.mlir
index 83d326c..29b95b4 100644
--- a/iree/test/microbenchmarks/mhlo_fft_abs.mlir
+++ b/iree/test/microbenchmarks/mhlo_fft_abs.mlir
@@ -6,7 +6,7 @@
%input = util.unfoldable_constant dense<1.0> : tensor<6x1024xf32>
%0 = "mhlo.fft"(%input) {
fft_length = dense<1024> : tensor<1xi64>,
- fft_type = "RFFT"
+ fft_type = #mhlo<"fft_type RFFT">
} : (tensor<6x1024xf32>) -> tensor<6x513xcomplex<f32>>
%1 = "mhlo.abs"(%0) : (tensor<6x513xcomplex<f32>>) -> tensor<6x513xf32>
return %1: tensor<6x513xf32>
