Add ML overlay C++ sample.

.gitignore

@@ -6,3 +6,6 @@ __pycache__
 # CMake artifacts
 build/
 build-*/
+
+# Sample artifacts
+imgui.ini

cpp/.clang-format

@@ -0,0 +1,9 @@
+# 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's runtime follows Google style while its compiler follows mostly LLVM
+# style (for naming/etc.) but using Google formatting.
+BasedOnStyle: Google

cpp/CMakeLists.txt

@@ -19,35 +19,54 @@ set_property(GLOBAL PROPERTY USE_FOLDERS ON)
 # Core project dependency
 #-------------------------------------------------------------------------------
 
-message(STATUS "Fetching core IREE repo (this may take a few minutes)...")
-# Note: for log output, set -DFETCHCONTENT_QUIET=OFF,
-# see https://gitlab.kitware.com/cmake/cmake/-/issues/18238#note_440475
-
-include(FetchContent)
-
-FetchContent_Declare(
-  iree
-  GIT_REPOSITORY https://github.com/iree-org/iree.git
-  GIT_TAG 7a435f0e45b2dbc3988c1a751e6810cd80c2dd83 # 2022-09-12
-  GIT_SUBMODULES_RECURSE OFF
-  GIT_SHALLOW OFF
-  GIT_PROGRESS ON
-  USES_TERMINAL_DOWNLOAD ON
-)
-
-# Extend module path to find MLIR CMake modules.
-list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_BINARY_DIR}/lib/cmake/mlir")
+option(IREE_SAMPLES_FETCH_SOURCE "Fetches IREE source code on-demand using CMake FetchContent" ON)
+# Path to an existing IREE source directory
+set(IREE_SAMPLES_SOURCE_PATH "" CACHE STRING "")
 
 # Disable core project features not needed for these out of tree samples.
 set(IREE_BUILD_TESTS OFF CACHE BOOL "" FORCE)
 set(IREE_BUILD_SAMPLES OFF CACHE BOOL "" FORCE)
+# Extend module path to find MLIR CMake modules.
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_BINARY_DIR}/lib/cmake/mlir")
+
+if(IREE_SAMPLES_FETCH_SOURCE)
+  message(STATUS "Fetching core IREE repo (this may take a few minutes)...")
+
+  # Note: for log output, set -DFETCHCONTENT_QUIET=OFF,
+  # see https://gitlab.kitware.com/cmake/cmake/-/issues/18238#note_440475
+  list(APPEND CMAKE_MESSAGE_INDENT "  ")
+  include(FetchContent)
+  FetchContent_Declare(
+    iree
+    GIT_REPOSITORY https://github.com/iree-org/iree.git
+    GIT_TAG 31a5dcb1910c9014db8339afd953eccaa0dc6076 # 2023-12-04
+    GIT_SUBMODULES_RECURSE OFF
+    GIT_SHALLOW OFF
+    GIT_PROGRESS ON
+    USES_TERMINAL_DOWNLOAD ON
+  )
+  FetchContent_MakeAvailable(iree)
+  FetchContent_GetProperties(iree SOURCE_DIR IREE_SOURCE_DIR)
+  list(POP_BACK CMAKE_MESSAGE_INDENT)
+else()
+  message(STATUS "Using IREE sources at path '${IREE_SAMPLES_SOURCE_PATH}'")
+  list(APPEND CMAKE_MESSAGE_INDENT "  ")
+  add_subdirectory(${IREE_SAMPLES_SOURCE_PATH} "third_party/iree" EXCLUDE_FROM_ALL)
+  list(POP_BACK CMAKE_MESSAGE_INDENT)
+endif()
+
+#-------------------------------------------------------------------------------
+# Common settings for dependencies
+#-------------------------------------------------------------------------------
 
-FetchContent_MakeAvailable(iree)
-FetchContent_GetProperties(iree SOURCE_DIR IREE_SOURCE_DIR)
+set(IMGUI_BUILD_SDL2_RENDERER_BINDING ON)
+set(IMGUI_BUILD_VULKAN_BINDING ON)
+set(IMGUI_BUILD_WIN32_BINDING ON)
 
 #-------------------------------------------------------------------------------
 # Individual samples
 #-------------------------------------------------------------------------------
 
+add_subdirectory(ml_overlay)
 add_subdirectory(vision_inference)
 add_subdirectory(vulkan_gui)

cpp/ml_overlay/CMakeLists.txt

@@ -0,0 +1,74 @@
+# 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
+
+if(NOT IREE_TARGET_BACKEND_VULKAN_SPIRV OR
+   NOT IREE_HAL_DRIVER_VULKAN)
+  message(STATUS "Missing Vulkan backend and/or driver, skipping ml_overlay sample")
+  return()
+endif()
+
+# This target statically links against Vulkan.
+# One way to achieve this is by installing the Vulkan SDK from
+# https://vulkan.lunarg.com/.
+include(FindVulkan)
+if(NOT Vulkan_FOUND)
+  message(STATUS "Could not find Vulkan, skipping ml_overlay sample")
+  return()
+endif()
+
+# vcpkg install imgui[vulkan-binding,sdl2-binding]
+#   tested with versions 1.83 - 1.87#1
+find_package(imgui)
+if(NOT imgui_FOUND)
+  message(STATUS "Could not find Dear ImGui, skipping ml_overlay sample")
+  return()
+endif()
+
+# vcpkg install sdl2[vulkan]
+#   tested with versions 2.0.14#4 - 2.0.22#1
+find_package(SDL2)
+if(NOT SDL2_FOUND)
+  message(STATUS "Could not find SDL2, skipping ml_overlay sample")
+  return()
+endif()
+
+# Define the sample executable.
+set(_NAME "iree-samples-ml-overlay")
+add_executable(${_NAME} "")
+target_sources(${_NAME}
+  PRIVATE
+    overlay.cc
+)
+set_target_properties(${_NAME} PROPERTIES OUTPUT_NAME "ml-overlay")
+target_include_directories(${_NAME} PUBLIC
+    $<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}>
+)
+target_link_libraries(${_NAME}
+  imgui::imgui
+  SDL2::SDL2
+  Vulkan::Vulkan
+  #
+  iree_base_internal_flags
+  iree_hal_drivers_vulkan_registration_registration
+  iree_modules_hal_hal
+  iree_tooling_context_util
+  iree_vm_vm
+)
+if(${CMAKE_SYSTEM_NAME} STREQUAL "Windows")
+  set(_GUI_LINKOPTS "-SUBSYSTEM:WINDOWS")
+  target_link_libraries(${_NAME}
+    d3d11.lib
+    dxgi.lib
+  )
+else()
+  set(_GUI_LINKOPTS "")
+endif()
+target_link_options(${_NAME}
+  PRIVATE
+    ${_GUI_LINKOPTS}
+)
+
+message(STATUS "Configured ml-overlay sample successfully")

cpp/ml_overlay/TODO.md

@@ -0,0 +1,62 @@
+ml-overlay cpp/ml_overlay/filters.vmfb
+(supports dragging vmfb on to exe)
+
+ml-overlay cpp/ml_overlay/filters.vmfb \
+  --filter=filter_add \
+  --always_on_top=false \
+  --overlay=false \
+  --x=100 \
+  --y=200 \
+  --width=1024 \
+  --height=1024 \
+  --imgui_demo_window
+
+iree-compile
+  --iree-execution-model=async-external \
+  --iree-hal-target-backends=vulkan-spirv \
+  --iree-input-type=mhlo \
+  cpp/ml_overlay/filters.mlir
+
+
+
+
+
+
+Window
+
+vkCmdCopyImageToBuffer region
+
+struct Frame {
+  VkImage source_image;
+  VkImageLayout source_layout;
+};
+
+Pipeline
+  ProcessFrame(Frame& frame);
+
+source image:
+  VkImage
+  optional semaphore when available?
+
+target image:
+  VkImage
+
+current results
+  accumulated
+  last produced image
+  last produced outputs vm_list
+
+source, target_storage -> target
+
+command buffer with copy image to buffer
+semaphore signal, import into fence
+submit command buffer
+VkBuffer import to iree_hal_buffer_t
+wrap in buffer view with size
+next fence
+invoke with copy wait semaphore, next fence for ready
+export fence to ready semaphore
+{ or use imgui command buffer?
+  command buffer with copy buffer to image
+  submit command buffer waiting on fence
+}

cpp/ml_overlay/filters.mlir

@@ -0,0 +1,74 @@
+func.func @filter_add(
+  %source: tensor<?x?x4xi8>
+) -> (tensor<?x?x4xi8>) {
+  %0 = arith.addi %source, %source : tensor<?x?x4xi8>
+  return %0 : tensor<?x?x4xi8>
+}
+
+func.func @filter_add_inplace(
+  %source: tensor<?x?x4xi8>,
+  %target_storage: !hal.buffer {iree.abi.output = 0 : index}
+) -> (tensor<?x?x4xi8>) {
+  %0 = arith.addi %source, %source : tensor<?x?x4xi8>
+  return %0 : tensor<?x?x4xi8>
+}
+
+// func.func @filter_invert(
+//   %source: tensor<?x?x4xi8>
+// ) -> (tensor<?x?x4xi8>) {
+//   %0 = arith.constant dense<[[[255, 255, 255, 255]]]> : tensor<1x1x4xi8>
+//   %1 = shape.shape_of %source : tensor<?x?x4xi8> -> tensor<3xindex>
+//   %2 = "mhlo.dynamic_broadcast_in_dim"(%0, %1) {broadcast_dimensions = dense<[0, 1, 2]> : tensor<3xi64>} : (tensor<1x1x4xi8>, tensor<3xindex>) -> tensor<?x?x4xi8>
+
+//   %3 = arith.constant dense<[[[0, 0, 0, 255]]]> : tensor<1x1x4xi8>
+//   %4 = "mhlo.dynamic_broadcast_in_dim"(%3, %1) {broadcast_dimensions = dense<[0, 1, 2]> : tensor<3xi64>} : (tensor<1x1x4xi8>, tensor<3xindex>) -> tensor<?x?x4xi8>
+
+//   %5 = arith.subi %2, %source : tensor<?x?x4xi8>
+//   %out = arith.addi %5, %4 : tensor<?x?x4xi8>
+//   return %out : tensor<?x?x4xi8>
+// }
+
+// func.func @filter_drain_red(
+//   %source: tensor<?x?x4xi8>
+// ) -> (tensor<?x?x4xi8>) {
+//   // bgra
+//   %0 = arith.constant dense<[[[255, 255, 200, 255]]]> : tensor<1x1x4xi8>
+//   %1 = shape.shape_of %source : tensor<?x?x4xi8> -> tensor<3xindex>
+//   %2 = "mhlo.dynamic_broadcast_in_dim"(%0, %1) {broadcast_dimensions = dense<[0, 1, 2]> : tensor<3xi64>} : (tensor<1x1x4xi8>, tensor<3xindex>) -> tensor<?x?x4xi8>
+
+//   %out = arith.minsi %2, %source : tensor<?x?x4xi8>
+//   return %out : tensor<?x?x4xi8>
+// }
+
+// func.func @filter_cst(
+//   %source: tensor<?x?x4xi8>
+// ) -> (tensor<?x?x4xi8>) {
+//   %0 = arith.constant dense<[[[255, 128, 0, 255]]]> : tensor<1x1x4xi8>
+//   %1 = shape.shape_of %source : tensor<?x?x4xi8> -> tensor<3xindex>
+//   %2 = "mhlo.dynamic_broadcast_in_dim"(%0, %1) {broadcast_dimensions = dense<[0, 1, 2]> : tensor<3xi64>} : (tensor<1x1x4xi8>, tensor<3xindex>) -> tensor<?x?x4xi8>
+//   return %2 : tensor<?x?x4xi8>
+// }
+
+func.func @filter_bboxes(
+  %source: tensor<?x?x4xi8>
+) -> (tensor<?x?x4xi8>, tensor<?x16xi32>) {
+  %0 = arith.addi %source, %source : tensor<?x?x4xi8>
+  // tensor<?x16xi32>
+  // id, data0, data1, data2 | x0, y0, x1, y1 | a, b, g, r | x x x x
+  %1 = arith.constant dense<[
+    [1000, 100, 101, 102,
+     20, 40, 50, 60,
+     255, 255, 0, 0,
+     0, 0, 0, 0],
+    [2000, 200, 201, 202,
+     200, 240, 350, 460,
+     255, 0, 0, 255,
+     0, 0, 0, 0],
+    [3000, 300, 301, 302,
+     100, 340, 250, 360,
+     128, 255, 255, 255,
+     0, 0, 0, 0]
+  ]> : tensor<3x16xi32>
+  %2 = tensor.cast %1 : tensor<3x16xi32> to tensor<?x16xi32>
+  return %0, %2 : tensor<?x?x4xi8>, tensor<?x16xi32>
+}