feature: Conversion of HSOpticalFlow Program to HIP (#453)

* feature: Conversion of HSOpticalFlow Program to HIP

* Remove tracked binary and fix stb_image include path

* fix: README linting errors

* fix: hip detection and compiling

* fix: ppm data can now be found no matter where you are

* fix: address PR review feedback for optical_flow

- Add #include <hip/hip_runtime.h> to derivativesKernel.hip and
  downscaleKernel.hip (reviewer request for self-contained headers)
- Fix LoadImageAsFP32 docstring: loads RGB (3-channel), not 4-channel
- Update copyright year to 2026 in CMakeLists.txt and Makefile
- Add optical_flow to Applications/Makefile EXAMPLES list
- Add optical_flow entry to root README.md
- Remove accidental .venv/ entry from root .gitignore
- Rewrite optical_flow README with accurate build/run instructions:
  correct CMake compiler flag (clang++ not hipcc), ROCM_PATH and
  LD_LIBRARY_PATH env vars, working-directory-independent run instructions

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>

* fix: Add hip_runtime to all needed files

---------

Co-authored-by: Claude Sonnet 4 <noreply@anthropic.com>

Author: mapatel-amd

.gitignore

@@ -8,4 +8,4 @@
 CMakeUserPresets.json
 .cline_storage
 config.mk
-.claude/
+.claude/

Applications/CMakeLists.txt

@@ -42,6 +42,7 @@ add_subdirectory(convolution)
 add_subdirectory(floyd_warshall)
 add_subdirectory(histogram)
 add_subdirectory(prefix_sum)
+add_subdirectory(optical_flow)
 
 if(NOT CXX_FS_HEADER_FOUND)
     message(WARNING "filesystem not found, not building fdtd example")

Applications/Makefile

@@ -26,6 +26,7 @@ EXAMPLES := \
   fdtd \
   floyd_warshall \
   histogram \
+  optical_flow \
   prefix_sum
 
 SKIP_FROM_TEST :=

Applications/optical_flow/.gitignore

@@ -0,0 +1,2 @@
+applications_optical_flow
+*.flo

Applications/optical_flow/CMakeLists.txt

@@ -0,0 +1,53 @@
+# MIT License
+#
+# Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+set(example_name applications_optical_flow)
+
+cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
+project(${example_name} LANGUAGES CXX)
+
+include("${CMAKE_CURRENT_LIST_DIR}/../../Common/HipPlatform.cmake")
+select_gpu_language()
+
+enable_language(${ROCM_EXAMPLES_GPU_LANGUAGE})
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_STANDARD 17)
+set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_EXTENSIONS OFF)
+set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_STANDARD_REQUIRED ON)
+select_hip_platform()
+
+include("${CMAKE_CURRENT_LIST_DIR}/../../Common/ROCmPath.cmake")
+
+add_executable(${example_name} main.hip flowHIP.hip flowGold.cpp)
+add_test(NAME ${example_name} COMMAND ${example_name})
+
+set(include_dirs "../../Common")
+if(ROCM_EXAMPLES_HIP_PLATFORM STREQUAL "nvidia")
+    list(APPEND include_dirs "${ROCM_PATH}/include")
+endif()
+
+target_include_directories(${example_name} PRIVATE ${include_dirs})
+set_source_files_properties(main.hip flowHIP.hip PROPERTIES LANGUAGE ${ROCM_EXAMPLES_GPU_LANGUAGE})
+
+install(TARGETS ${example_name})

Applications/optical_flow/Makefile

@@ -0,0 +1,61 @@
+# MIT License
+#
+# Copyright (c) 2022-2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+EXAMPLE := applications_optical_flow
+COMMON_INCLUDE_DIR := ../../Common
+GPU_RUNTIME ?= HIP
+
+ROCM_PATH ?= /opt/rocm
+HIP_INCLUDE_DIR := $(ROCM_PATH)/include
+
+HIPCXX ?= $(ROCM_PATH)/bin/hipcc
+
+CXX_STD   := c++17
+ICXXFLAGS := -std=$(CXX_STD)
+ICPPFLAGS := -I $(COMMON_INCLUDE_DIR)
+ILDFLAGS  :=
+ILDLIBS   :=
+
+ifeq ($(GPU_RUNTIME), CUDA)
+ICXXFLAGS += -x cu
+ICPPFLAGS += -isystem $(HIP_INCLUDE_DIR)
+else ifeq ($(GPU_RUNTIME), HIP)
+CXXFLAGS ?= -Wall -Wextra
+else
+$(error GPU_RUNTIME is set to "$(GPU_RUNTIME)". GPU_RUNTIME must be either CUDA or HIP)
+endif
+
+ICXXFLAGS += $(CXXFLAGS)
+ICPPFLAGS += $(CPPFLAGS)
+ILDFLAGS  += $(LDFLAGS)
+ILDLIBS   += $(LDLIBS)
+
+$(EXAMPLE): main.hip flowHIP.hip flowGold.cpp $(COMMON_INCLUDE_DIR)/example_utils.hpp
+	$(HIPCXX) $(ICXXFLAGS) $(ICPPFLAGS) $(ILDFLAGS) -o $@ main.hip flowHIP.hip flowGold.cpp $(ILDLIBS)
+
+test: $(EXAMPLE)
+	./$(EXAMPLE) $(TEST_ARGS)
+
+clean:
+	$(RM) $(EXAMPLE)
+
+.PHONY: clean test

Applications/optical_flow/README.md

@@ -0,0 +1,126 @@
+# Optical Flow
+
+## Description
+
+This example implements the **Horn-Schunck variational optical flow** algorithm using HIP. It estimates the motion field (displacement vectors) between two consecutive image frames by minimizing a global energy functional that balances data fidelity and spatial smoothness.
+
+The algorithm operates on a Gaussian image pyramid: flow is computed coarse-to-fine, with each level refining the estimate from the level above. At each pyramid level, image warping aligns the target frame with the source, and a Jacobi iterative solver computes the incremental flow update.
+
+The program computes optical flow on both the CPU (`flowGold`) and GPU (`flowHIP`), compares the results via L1 norm, and writes two `.flo` files (Middlebury format) for inspection.
+
+## Application Flow
+
+1. Load two consecutive frames (`data/frame10.ppm`, `data/frame11.ppm`) as single-channel FP32 images.
+2. Build a Gaussian pyramid with `nLevels` levels by repeatedly downscaling with a 4-tap filter.
+3. At each pyramid level (coarse to fine):
+   - Upscale the flow estimate from the coarser level.
+   - Warp the target image toward the source using the current flow estimate.
+   - Compute image derivatives (Ix, Iy, Iz) via finite differences.
+   - Run `nSolverIters` Jacobi iterations to solve for the incremental flow update.
+   - Repeat for `nWarpIters` warping passes.
+4. Copy GPU results to host and compare against the CPU reference (L1 norm per pixel).
+5. Write `FlowGPU.flo` and `FlowCPU.flo` to the working directory.
+
+## Key APIs and Concepts
+
+| Concept | HIP API |
+|---|---|
+| Texture objects with bilinear filtering | `hipCreateTextureObject`, `hipTextureObject_t` |
+| Pitched 2D texture resource | `hipResourceTypePitch2D`, `hipResourceDesc` |
+| Mirror address mode | `hipAddressModeMirror` |
+| Normalized texture coordinates | `texDescr.normalizedCoords = true` |
+| In-kernel texture fetch | `tex2D<float>(tex, x, y)` |
+| Block synchronization | `cg::this_thread_block()`, `cg::sync()` |
+
+### Pitch Alignment Requirement
+
+ROCm requires `pitchInBytes` for `hipResourceTypePitch2D` to be a multiple of **256 bytes** (64 floats × 4 bytes). The `STRIDE_ALIGNMENT` constant in `common.h` is set to `64` to satisfy this constraint. CUDA only requires 128 bytes (32 floats), so porting code that used `StrideAlignment = 32` will fail at texture creation.
+
+## Prerequisites
+
+- A ROCm-capable AMD GPU
+- ROCm SDK installed ([installation guide](https://rocm.docs.amd.com/en/latest/index.html) or [TheRock releases](https://github.com/ROCm/TheRock/blob/main/RELEASES.md))
+
+## Building
+
+Set `ROCM_PATH` to your ROCm installation root before building. For a standard system install:
+
+```bash
+export ROCM_PATH=/opt/rocm
+```
+
+For a Python venv-based install (e.g. TheRock):
+
+```bash
+export ROCM_PATH=/path/to/venv/lib/python3.12/site-packages/_rocm_sdk_devel
+```
+
+### Make
+
+```bash
+cd Applications/optical_flow
+make ROCM_PATH=$ROCM_PATH
+```
+
+If your ROCm device libraries are not found automatically, pass their path explicitly:
+
+```bash
+make ROCM_PATH=$ROCM_PATH \
+     CXXFLAGS="--rocm-device-lib-path=$ROCM_PATH/lib/llvm/amdgcn/bitcode"
+```
+
+### CMake
+
+CMake 3.28 and later require passing `clang++` directly rather than the `hipcc` wrapper script:
+
+```bash
+cd Applications/optical_flow
+cmake -B build \
+  -DROCM_PATH=$ROCM_PATH \
+  -DCMAKE_HIP_COMPILER=$ROCM_PATH/lib/llvm/bin/clang++
+cmake --build build -j$(nproc)
+```
+
+## Running
+
+If ROCm is not installed to a standard system path, set `LD_LIBRARY_PATH` so the runtime libraries can be found:
+
+```bash
+export LD_LIBRARY_PATH=$ROCM_PATH/lib:$LD_LIBRARY_PATH
+```
+
+The binary locates the sample frames (`data/frame10.ppm`, `data/frame11.ppm`) relative to the source directory automatically, so it can be run from any working directory:
+
+```bash
+# Make build
+./applications_optical_flow
+
+# CMake build
+./build/applications_optical_flow
+```
+
+### Expected Output
+
+```text
+HSOpticalFlow Starting...
+
+Using device: <GPU name>
+Loading "<source-dir>/data/frame10.ppm" ...
+Loading "<source-dir>/data/frame11.ppm" ...
+Computing optical flow on CPU...
+Computing optical flow on GPU...
+L1 error : 0.000xxx
+```
+
+The program exits with `EXIT_SUCCESS` when the L1 error between the GPU and CPU results is below `0.05`. Two output files are written to the current working directory:
+
+- `FlowGPU.flo` — GPU optical flow result
+- `FlowCPU.flo` — CPU reference result
+
+Both files use the [Middlebury `.flo` format](http://vision.middlebury.edu/flow/code/flow-code/README.txt) and can be visualized with tools such as `flowiz` or the Middlebury flow utilities.
+
+## Key Notes
+
+- `STRIDE_ALIGNMENT` is 64 because `hipResourceTypePitch2D` requires `pitchInBytes` to be a multiple of 256 bytes (4 bytes × 64 floats = 256 bytes).
+- CMake 3.28+ does not accept the `hipcc` wrapper as `CMAKE_HIP_COMPILER`. Pass the `clang++` binary from `$ROCM_PATH/lib/llvm/bin/clang++` instead.
+- The binary resolves the input image paths relative to the source file location at compile time (via `__FILE__`/`/proc/self/exe`), so no specific working directory is required at runtime.

Applications/optical_flow/addKernel.hip

@@ -0,0 +1,57 @@
+// MIT License
+//
+// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include "common.h"
+
+///////////////////////////////////////////////////////////////////////////////
+/// \brief add two vectors of size _count_
+///
+/// HIP kernel
+/// \param[in]  op1   term one
+/// \param[in]  op2   term two
+/// \param[in]  count vector size
+/// \param[out] sum   result
+///////////////////////////////////////////////////////////////////////////////
+__global__ void AddKernel(const float *op1, const float *op2, int count, float *sum)
+{
+    const int pos = threadIdx.x + blockIdx.x * blockDim.x;
+
+    if (pos >= count)
+        return;
+
+    sum[pos] = op1[pos] + op2[pos];
+}
+
+///////////////////////////////////////////////////////////////////////////////
+/// \brief add two vectors of size _count_
+/// \param[in]  op1   term one
+/// \param[in]  op2   term two
+/// \param[in]  count vector size
+/// \param[out] sum   result
+///////////////////////////////////////////////////////////////////////////////
+static void Add(const float *op1, const float *op2, int count, float *sum)
+{
+    dim3 threads(256);
+    dim3 blocks(iDivUp(count, threads.x));
+
+    AddKernel<<<blocks, threads>>>(op1, op2, count, sum);
+}

Applications/optical_flow/common.h

@@ -0,0 +1,50 @@
+// MIT License
+//
+// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#pragma once
+
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+// HIP hipResourceTypePitch2D requires pitchInBytes to be a multiple of 256 bytes.
+// 64 floats * 4 bytes = 256 bytes satisfies this
+constexpr int STRIDE_ALIGNMENT = 64;
+
+inline int iAlignUp(int n, int m = STRIDE_ALIGNMENT)
+{
+    int mod = n % m;
+    if (mod)
+        return n + m - mod;
+    else
+        return n;
+}
+
+inline int iDivUp(int n, int m) { return (n + m - 1) / m; }
+
+template <typename T> inline void Swap(T &a, T &b)
+{
+    T t = a;
+    a   = b;
+    b   = t;
+}