diff --git a/conda/environments/all_cuda-129_arch-aarch64.yaml b/conda/environments/all_cuda-129_arch-aarch64.yaml
index 096a6f7bc..4ea32715c 100644
--- a/conda/environments/all_cuda-129_arch-aarch64.yaml
+++ b/conda/environments/all_cuda-129_arch-aarch64.yaml
@@ -17,15 +17,15 @@ dependencies:
 - imagecodecs>=2021.6.8
 - ipython
 - lazy-loader>=0.4
-- libnvimgcodec-dev>=0.8.0,<0.9.0
-- libnvimgcodec>=0.8.0,<0.9.0
+- libnvimgcodec-dev>=0.9.0,<0.10.0
+- libnvimgcodec>=0.9.0,<0.10.0
 - libnvjpeg-dev
 - matplotlib-base>=3.7
 - nbsphinx
 - ninja
 - numpy>=2.0,<3.0
 - numpydoc>=1.7
-- nvimgcodec>=0.8.0,<0.9.0
+- nvimgcodec>=0.9.0,<0.10.0
 - openslide-python>=1.3.0
 - pip
 - pooch>=1.6.0
diff --git a/conda/environments/all_cuda-129_arch-x86_64.yaml b/conda/environments/all_cuda-129_arch-x86_64.yaml
index 6d67658ec..361862de5 100644
--- a/conda/environments/all_cuda-129_arch-x86_64.yaml
+++ b/conda/environments/all_cuda-129_arch-x86_64.yaml
@@ -18,15 +18,15 @@ dependencies:
 - ipython
 - lazy-loader>=0.4
 - libcufile-dev
-- libnvimgcodec-dev>=0.8.0,<0.9.0
-- libnvimgcodec>=0.8.0,<0.9.0
+- libnvimgcodec-dev>=0.9.0,<0.10.0
+- libnvimgcodec>=0.9.0,<0.10.0
 - libnvjpeg-dev
 - matplotlib-base>=3.7
 - nbsphinx
 - ninja
 - numpy>=2.0,<3.0
 - numpydoc>=1.7
-- nvimgcodec>=0.8.0,<0.9.0
+- nvimgcodec>=0.9.0,<0.10.0
 - openslide-python>=1.3.0
 - pip
 - pooch>=1.6.0
diff --git a/conda/environments/all_cuda-133_arch-aarch64.yaml b/conda/environments/all_cuda-133_arch-aarch64.yaml
index eee401cca..949001071 100644
--- a/conda/environments/all_cuda-133_arch-aarch64.yaml
+++ b/conda/environments/all_cuda-133_arch-aarch64.yaml
@@ -17,15 +17,15 @@ dependencies:
 - imagecodecs>=2021.6.8
 - ipython
 - lazy-loader>=0.4
-- libnvimgcodec-dev>=0.8.0,<0.9.0
-- libnvimgcodec>=0.8.0,<0.9.0
+- libnvimgcodec-dev>=0.9.0,<0.10.0
+- libnvimgcodec>=0.9.0,<0.10.0
 - libnvjpeg-dev
 - matplotlib-base>=3.7
 - nbsphinx
 - ninja
 - numpy>=2.0,<3.0
 - numpydoc>=1.7
-- nvimgcodec>=0.8.0,<0.9.0
+- nvimgcodec>=0.9.0,<0.10.0
 - openslide-python>=1.3.0
 - pip
 - pooch>=1.6.0
diff --git a/conda/environments/all_cuda-133_arch-x86_64.yaml b/conda/environments/all_cuda-133_arch-x86_64.yaml
index c0a0d491e..86a1e653a 100644
--- a/conda/environments/all_cuda-133_arch-x86_64.yaml
+++ b/conda/environments/all_cuda-133_arch-x86_64.yaml
@@ -18,15 +18,15 @@ dependencies:
 - ipython
 - lazy-loader>=0.4
 - libcufile-dev
-- libnvimgcodec-dev>=0.8.0,<0.9.0
-- libnvimgcodec>=0.8.0,<0.9.0
+- libnvimgcodec-dev>=0.9.0,<0.10.0
+- libnvimgcodec>=0.9.0,<0.10.0
 - libnvjpeg-dev
 - matplotlib-base>=3.7
 - nbsphinx
 - ninja
 - numpy>=2.0,<3.0
 - numpydoc>=1.7
-- nvimgcodec>=0.8.0,<0.9.0
+- nvimgcodec>=0.9.0,<0.10.0
 - openslide-python>=1.3.0
 - pip
 - pooch>=1.6.0
diff --git a/conda/recipes/libcucim/conda_build_config.yaml b/conda/recipes/libcucim/conda_build_config.yaml
index 189db8c16..2e1a48037 100644
--- a/conda/recipes/libcucim/conda_build_config.yaml
+++ b/conda/recipes/libcucim/conda_build_config.yaml
@@ -17,4 +17,4 @@ cmake_version:
   - ">=4.0"
 
 nvimgcodec_version:
-  - ">=0.8.0,<0.9.0"
+  - ">=0.9.0,<0.10.0"
diff --git a/cpp/cmake/deps/nvimgcodec.cmake b/cpp/cmake/deps/nvimgcodec.cmake
index 38e5689f9..9176fb14d 100644
--- a/cpp/cmake/deps/nvimgcodec.cmake
+++ b/cpp/cmake/deps/nvimgcodec.cmake
@@ -1,19 +1,83 @@
 #
 # cmake-format: off
-# SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 # cmake-format: on
 #
 
 if (NOT TARGET deps::nvimgcodec)
+    # Packaging pins nvImageCodec to >=0.9.0,<0.10.0 and the vendored headers
+    # declare 0.9, so only 0.9 is supported here.  The unversioned entries are
+    # kept for wheel and system layouts that ship just the symlinks.
+    set(NVIMGCODEC_SONAME_CANDIDATES
+        "libnvimgcodec.so.0.9.0"
+        "libnvimgcodec.so.0"
+        "libnvimgcodec.so")
+
+    function(_cucim_find_nvimgcodec_library out_var root_dir)
+        foreach(_soname IN LISTS NVIMGCODEC_SONAME_CANDIDATES)
+            if(EXISTS "${root_dir}/${_soname}")
+                set(${out_var} "${root_dir}/${_soname}" PARENT_SCOPE)
+                return()
+            endif()
+        endforeach()
+        set(${out_var} "" PARENT_SCOPE)
+    endfunction()
+
+    # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above
+    # resolve to whatever major-0 build is installed, so a version below the
+    # required minimum would otherwise be picked up silently.  That API is not
+    # compatible and there are no compile-time version guards, so fail at
+    # configure time instead of at runtime.
+    function(_cucim_check_nvimgcodec_version include_dir)
+        set(_version_header "${include_dir}/nvimgcodec_version.h")
+        if(NOT EXISTS "${_version_header}")
+            message(WARNING
+                "nvImageCodec: nvimgcodec_version.h not found under '${include_dir}' - "
+                "skipping version check. cuCIM requires >=0.9.0,<0.10.0.")
+            return()
+        endif()
+
+        file(STRINGS "${_version_header}" _major_line
+             REGEX "^#define[ \t]+NVIMGCODEC_VER_MAJOR[ \t]+[0-9]+")
+        file(STRINGS "${_version_header}" _minor_line
+             REGEX "^#define[ \t]+NVIMGCODEC_VER_MINOR[ \t]+[0-9]+")
+        string(REGEX REPLACE ".*NVIMGCODEC_VER_MAJOR[ \t]+([0-9]+).*" "\\1" _major "${_major_line}")
+        string(REGEX REPLACE ".*NVIMGCODEC_VER_MINOR[ \t]+([0-9]+).*" "\\1" _minor "${_minor_line}")
+
+        if(NOT _major MATCHES "^[0-9]+$" OR NOT _minor MATCHES "^[0-9]+$")
+            message(WARNING
+                "nvImageCodec: could not parse version from '${_version_header}' - "
+                "skipping version check.")
+            return()
+        endif()
+
+        if(_major EQUAL 0 AND _minor LESS 9)
+            message(FATAL_ERROR
+                "nvImageCodec ${_major}.${_minor} found at '${include_dir}', but cuCIM "
+                "requires >=0.9.0 (packaging pins >=0.9.0,<0.10.0). The 0.9 API is not "
+                "backward compatible and this build has no compile-time version guards, "
+                "so it would fail at runtime. Install libnvimgcodec-dev >=0.9.")
+        elseif(_major GREATER 0 OR _minor GREATER_EQUAL 10)
+            message(WARNING
+                "nvImageCodec ${_major}.${_minor} is outside the pinned range "
+                "(>=0.9.0,<0.10.0) and is untested with this cuCIM build.")
+        endif()
+    endfunction()
+
     # First try to find it as a package
     find_package(nvimgcodec QUIET)
 
     if(nvimgcodec_FOUND)
         # Use the found package
+        if(DEFINED nvimgcodec_VERSION AND nvimgcodec_VERSION VERSION_LESS "0.9.0")
+            message(FATAL_ERROR
+                "nvImageCodec ${nvimgcodec_VERSION} found via find_package, but cuCIM "
+                "requires >=0.9.0 (packaging pins >=0.9.0,<0.10.0).")
+        endif()
         add_library(deps::nvimgcodec INTERFACE IMPORTED GLOBAL)
         target_link_libraries(deps::nvimgcodec INTERFACE nvimgcodec::nvimgcodec)
-        message(STATUS "✓ nvImageCodec found via find_package")
+        message(STATUS "✓ nvImageCodec found via find_package (version: ${nvimgcodec_VERSION})")
     else()
         # Manual detection in various environments
         set(NVIMGCODEC_LIB_PATH "")
@@ -22,31 +86,23 @@ if (NOT TARGET deps::nvimgcodec)
         # Try conda environment detection (both Python packages and native packages)
         if(DEFINED ENV{CONDA_BUILD})
             # Conda build environment
-            set(NVIMGCODEC_LIB_PATH "$ENV{PREFIX}/lib/libnvimgcodec.so.0")
+            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "$ENV{PREFIX}/lib")
             set(NVIMGCODEC_INCLUDE_PATH "$ENV{PREFIX}/include/")
-            if(NOT EXISTS "${NVIMGCODEC_LIB_PATH}")
-                set(NVIMGCODEC_LIB_PATH "$ENV{PREFIX}/lib/libnvimgcodec.so")
-            endif()
         elseif(DEFINED ENV{CONDA_PREFIX})
             # Active conda environment - try native package first
             set(CONDA_NATIVE_ROOT "$ENV{CONDA_PREFIX}")
             if(EXISTS "${CONDA_NATIVE_ROOT}/include/nvimgcodec.h")
                 set(NVIMGCODEC_INCLUDE_PATH "${CONDA_NATIVE_ROOT}/include/")
-                if(EXISTS "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so.0")
-                    set(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so.0")
-                elseif(EXISTS "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so")
-                    set(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so")
-                endif()
+                _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib")
             else()
                 # Fallback: try Python site-packages in conda environment
                 foreach(PY_VER "3.13" "3.12" "3.11" "3.10" "3.9")
                     set(CONDA_PYTHON_ROOT "$ENV{CONDA_PREFIX}/lib/python${PY_VER}/site-packages/nvidia/nvimgcodec")
                     if(EXISTS "${CONDA_PYTHON_ROOT}/include/nvimgcodec.h")
                         set(NVIMGCODEC_INCLUDE_PATH "${CONDA_PYTHON_ROOT}/include/")
-                        if(EXISTS "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                        elseif(EXISTS "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so")
+                        _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib")
+                        if(NOT NVIMGCODEC_LIB_PATH)
+                            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}")
                         endif()
                         break()
                     endif()
@@ -67,10 +123,9 @@ if (NOT TARGET deps::nvimgcodec)
                     set(NVIMGCODEC_PYTHON_ROOT "${PYTHON_SITE_PACKAGES}/nvidia/nvimgcodec")
                     if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/include/nvimgcodec.h")
                         set(NVIMGCODEC_INCLUDE_PATH "${NVIMGCODEC_PYTHON_ROOT}/include/")
-                        if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
+                        _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib")
+                        if(NOT NVIMGCODEC_LIB_PATH)
+                            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}")
                         endif()
                     endif()
                 endif()
@@ -78,14 +133,18 @@ if (NOT TARGET deps::nvimgcodec)
 
             # System-wide installation fallback
             if(NOT NVIMGCODEC_LIB_PATH)
-                if(EXISTS /usr/lib/x86_64-linux-gnu/libnvimgcodec.so.0)
-                    set(NVIMGCODEC_LIB_PATH /usr/lib/x86_64-linux-gnu/libnvimgcodec.so.0)
+                _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib/x86_64-linux-gnu")
+                if(NVIMGCODEC_LIB_PATH)
                     set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
-                elseif(EXISTS /usr/lib/aarch64-linux-gnu/libnvimgcodec.so.0)
-                    set(NVIMGCODEC_LIB_PATH /usr/lib/aarch64-linux-gnu/libnvimgcodec.so.0)
+                else()
+                    _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib/aarch64-linux-gnu")
+                endif()
+                if(NVIMGCODEC_LIB_PATH)
                     set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
-                elseif(EXISTS /usr/lib64/libnvimgcodec.so.0) # CentOS (x86_64)
-                    set(NVIMGCODEC_LIB_PATH /usr/lib64/libnvimgcodec.so.0)
+                else()
+                    _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib64") # CentOS (x86_64)
+                endif()
+                if(NVIMGCODEC_LIB_PATH)
                     set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
                 endif()
             endif()
@@ -93,6 +152,7 @@ if (NOT TARGET deps::nvimgcodec)
 
         # Create the target if we found the library
         if(NVIMGCODEC_LIB_PATH AND EXISTS "${NVIMGCODEC_LIB_PATH}")
+            _cucim_check_nvimgcodec_version("${NVIMGCODEC_INCLUDE_PATH}")
             add_library(deps::nvimgcodec SHARED IMPORTED GLOBAL)
             set_target_properties(deps::nvimgcodec PROPERTIES
                 IMPORTED_LOCATION "${NVIMGCODEC_LIB_PATH}"
diff --git a/cpp/plugins/cucim.kit.cuslide2/CMakeLists.txt b/cpp/plugins/cucim.kit.cuslide2/CMakeLists.txt
index e1aae4073..4f6e85f21 100644
--- a/cpp/plugins/cucim.kit.cuslide2/CMakeLists.txt
+++ b/cpp/plugins/cucim.kit.cuslide2/CMakeLists.txt
@@ -153,6 +153,8 @@ set(CUSLIDE2_CORE_SOURCES
     # TIFF structure management (uses nvImageCodec for parsing)
     ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/ifd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/ifd.h
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/ome_xml.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/ome_xml.h
     ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/tiff.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/tiff.h
     ${CMAKE_CURRENT_SOURCE_DIR}/src/cuslide/tiff/types.h
diff --git a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake
index e0f780513..ac8943ea5 100644
--- a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake
+++ b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake
@@ -1,6 +1,6 @@
 #
 # cmake-format: off
-# SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 # cmake-format: on
 #
@@ -9,13 +9,77 @@ if (NOT TARGET deps::nvimgcodec)
     set(NVIMGCODEC_LIB_PATH "")
     set(NVIMGCODEC_INCLUDE_PATH "")
 
+    # Packaging pins nvImageCodec to >=0.9.0,<0.10.0 and the vendored headers
+    # declare 0.9, so only 0.9 is supported here.  The unversioned entries are
+    # kept for wheel and system layouts that ship just the symlinks.
+    set(NVIMGCODEC_SONAME_CANDIDATES
+        "libnvimgcodec.so.0.9.0"
+        "libnvimgcodec.so.0"
+        "libnvimgcodec.so")
+
+    function(_cucim_find_nvimgcodec_library out_var root_dir)
+        foreach(_soname IN LISTS NVIMGCODEC_SONAME_CANDIDATES)
+            if(EXISTS "${root_dir}/${_soname}")
+                set(${out_var} "${root_dir}/${_soname}" PARENT_SCOPE)
+                return()
+            endif()
+        endforeach()
+        set(${out_var} "" PARENT_SCOPE)
+    endfunction()
+
+    # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above
+    # resolve to whatever major-0 build is installed, so a version below the
+    # required minimum would otherwise be picked up silently.  That API is not
+    # compatible and there are no compile-time version guards, so fail at
+    # configure time instead of at runtime.
+    function(_cucim_check_nvimgcodec_version include_dir)
+        set(_version_header "${include_dir}/nvimgcodec_version.h")
+        if(NOT EXISTS "${_version_header}")
+            message(WARNING
+                "nvImageCodec: nvimgcodec_version.h not found under '${include_dir}' - "
+                "skipping version check. cuCIM requires >=0.9.0,<0.10.0.")
+            return()
+        endif()
+
+        file(STRINGS "${_version_header}" _major_line
+             REGEX "^#define[ \t]+NVIMGCODEC_VER_MAJOR[ \t]+[0-9]+")
+        file(STRINGS "${_version_header}" _minor_line
+             REGEX "^#define[ \t]+NVIMGCODEC_VER_MINOR[ \t]+[0-9]+")
+        string(REGEX REPLACE ".*NVIMGCODEC_VER_MAJOR[ \t]+([0-9]+).*" "\\1" _major "${_major_line}")
+        string(REGEX REPLACE ".*NVIMGCODEC_VER_MINOR[ \t]+([0-9]+).*" "\\1" _minor "${_minor_line}")
+
+        if(NOT _major MATCHES "^[0-9]+$" OR NOT _minor MATCHES "^[0-9]+$")
+            message(WARNING
+                "nvImageCodec: could not parse version from '${_version_header}' - "
+                "skipping version check.")
+            return()
+        endif()
+
+        if(_major EQUAL 0 AND _minor LESS 9)
+            message(FATAL_ERROR
+                "nvImageCodec ${_major}.${_minor} found at '${include_dir}', but cuCIM "
+                "requires >=0.9.0 (packaging pins >=0.9.0,<0.10.0). The 0.9 API is not "
+                "backward compatible and this build has no compile-time version guards, "
+                "so it would fail at runtime. Install libnvimgcodec-dev >=0.9.")
+        elseif(_major GREATER 0 OR _minor GREATER_EQUAL 10)
+            message(WARNING
+                "nvImageCodec ${_major}.${_minor} is outside the pinned range "
+                "(>=0.9.0,<0.10.0) and is untested with this cuCIM build.")
+        endif()
+    endfunction()
+
     # Try find_package first (works in conda and system installations)
     find_package(nvimgcodec QUIET)
 
     if (nvimgcodec_FOUND)
+        if(DEFINED nvimgcodec_VERSION AND nvimgcodec_VERSION VERSION_LESS "0.9.0")
+            message(FATAL_ERROR
+                "nvImageCodec ${nvimgcodec_VERSION} found via find_package, but cuCIM "
+                "requires >=0.9.0 (packaging pins >=0.9.0,<0.10.0).")
+        endif()
         add_library(deps::nvimgcodec INTERFACE IMPORTED GLOBAL)
         target_link_libraries(deps::nvimgcodec INTERFACE nvimgcodec::nvimgcodec)
-        message(STATUS "✓ nvImageCodec found via find_package")
+        message(STATUS "✓ nvImageCodec found via find_package (version: ${nvimgcodec_VERSION})")
     else()
         # Manual detection: try conda environment, Python site-packages, and system paths
         if (DEFINED ENV{CONDA_PREFIX})
@@ -23,11 +87,7 @@ if (NOT TARGET deps::nvimgcodec)
             set(CONDA_NATIVE_ROOT "$ENV{CONDA_PREFIX}")
             if(EXISTS "${CONDA_NATIVE_ROOT}/include/nvimgcodec.h")
                 set(NVIMGCODEC_INCLUDE_PATH "${CONDA_NATIVE_ROOT}/include/")
-                if(EXISTS "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so.0")
-                    set(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so.0")
-                elseif(EXISTS "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so")
-                    set(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib/libnvimgcodec.so")
-                endif()
+                _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_NATIVE_ROOT}/lib")
             endif()
 
             # Fallback: try Python site-packages in conda environment
@@ -36,16 +96,10 @@ if (NOT TARGET deps::nvimgcodec)
                     set(CONDA_PYTHON_ROOT "$ENV{CONDA_PREFIX}/lib/python${PY_VER}/site-packages/nvidia/nvimgcodec")
                     if(EXISTS "${CONDA_PYTHON_ROOT}/include/nvimgcodec.h")
                         set(NVIMGCODEC_INCLUDE_PATH "${CONDA_PYTHON_ROOT}/include/")
-                        # Check for library in lib/ subdirectory first (conda package structure)
-                        if(EXISTS "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                        elseif(EXISTS "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                        # Check for library directly in nvimgcodec directory (pip package structure)
-                        elseif(EXISTS "${CONDA_PYTHON_ROOT}/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/libnvimgcodec.so.0")
-                        elseif(EXISTS "${CONDA_PYTHON_ROOT}/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/libnvimgcodec.so")
+                        # Check conda package layout first, then pip-style layout.
+                        _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}/lib")
+                        if(NOT NVIMGCODEC_LIB_PATH)
+                            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${CONDA_PYTHON_ROOT}")
                         endif()
                         if(NVIMGCODEC_LIB_PATH)
                             break()
@@ -80,14 +134,9 @@ if (NOT TARGET deps::nvimgcodec)
                     set(NVIMGCODEC_PYTHON_ROOT "${PYTHON_USER_SITE_PACKAGES}/nvidia/nvimgcodec")
                     if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/include/nvimgcodec.h")
                         set(NVIMGCODEC_INCLUDE_PATH "${NVIMGCODEC_PYTHON_ROOT}/include/")
-                        if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so.0")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so")
+                        _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib")
+                        if(NOT NVIMGCODEC_LIB_PATH)
+                            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}")
                         endif()
                     endif()
                 endif()
@@ -97,14 +146,9 @@ if (NOT TARGET deps::nvimgcodec)
                     set(NVIMGCODEC_PYTHON_ROOT "${PYTHON_SITE_PACKAGES}/nvidia/nvimgcodec")
                     if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/include/nvimgcodec.h")
                         set(NVIMGCODEC_INCLUDE_PATH "${NVIMGCODEC_PYTHON_ROOT}/include/")
-                        if(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so.0")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib/libnvimgcodec.so")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so.0")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so.0")
-                        elseif(EXISTS "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so")
-                            set(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/libnvimgcodec.so")
+                        _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}/lib")
+                        if(NOT NVIMGCODEC_LIB_PATH)
+                            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "${NVIMGCODEC_PYTHON_ROOT}")
                         endif()
                     endif()
                 endif()
@@ -113,20 +157,25 @@ if (NOT TARGET deps::nvimgcodec)
 
         # System-wide installation fallback
         if(NOT NVIMGCODEC_LIB_PATH)
-            if(EXISTS /usr/lib/x86_64-linux-gnu/libnvimgcodec.so.0)
-                set(NVIMGCODEC_LIB_PATH /usr/lib/x86_64-linux-gnu/libnvimgcodec.so.0)
+            _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib/x86_64-linux-gnu")
+            if(NVIMGCODEC_LIB_PATH)
                 set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
-            elseif(EXISTS /usr/lib/aarch64-linux-gnu/libnvimgcodec.so.0)
-                set(NVIMGCODEC_LIB_PATH /usr/lib/aarch64-linux-gnu/libnvimgcodec.so.0)
+            else()
+                _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib/aarch64-linux-gnu")
+            endif()
+            if(NVIMGCODEC_LIB_PATH)
                 set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
-            elseif(EXISTS /usr/lib64/libnvimgcodec.so.0)
-                set(NVIMGCODEC_LIB_PATH /usr/lib64/libnvimgcodec.so.0)
+            else()
+                _cucim_find_nvimgcodec_library(NVIMGCODEC_LIB_PATH "/usr/lib64")
+            endif()
+            if(NVIMGCODEC_LIB_PATH)
                 set(NVIMGCODEC_INCLUDE_PATH "/usr/include/")
             endif()
         endif()
 
         # Create target only if nvImageCodec was found
         if(NVIMGCODEC_LIB_PATH AND EXISTS "${NVIMGCODEC_LIB_PATH}")
+            _cucim_check_nvimgcodec_version("${NVIMGCODEC_INCLUDE_PATH}")
             add_library(deps::nvimgcodec SHARED IMPORTED GLOBAL)
             set_target_properties(deps::nvimgcodec PROPERTIES
                 IMPORTED_LOCATION "${NVIMGCODEC_LIB_PATH}"
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp
index 790609a4e..d3c73e74b 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp
@@ -20,6 +20,7 @@
 #include <cstring>
 #include <fcntl.h>
 #include <filesystem>
+#include <memory_resource>
 #include <memory>
 
 using json = nlohmann::json;
@@ -104,21 +105,21 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
     // Detect if this is a Philips TIFF file
     // Philips TIFF also has multiple SubfileType=0 (by design)
     bool is_philips_tiff = (tif->tiff_type() == cuslide::tiff::TiffType::Philips);
+    const std::string& first_desc = tif->ifd(0)->image_description();
+    bool looks_like_ome = (first_desc.find("<OME") != std::string::npos ||
+                           first_desc.find(":OME") != std::string::npos);
 
-    // Fallback: detect unrecognised pyramidal TIFFs that were not caught by
-    // ImageDescription or nvImageCodec metadata-blob detection above.
-    // If the file has >=3 IFDs forming a resolution pyramid, assume it is a
-    // valid multi-resolution WSI so it is not rejected by the SubfileType=0
-    // single-main-image constraint below.
-    // TODO: revisit once all known pyramid formats are covered by the
-    //       is_known_multi_ifd_format check — this heuristic may become unnecessary.
+    // Fallback for files with no usable ImageDescription: check for multiple
+    // resolution levels. Aperio SVS files typically have 3-6 IFDs with multiple
+    // resolution levels, so treat a pyramid-shaped IFD chain as Aperio/SVS.
     if (!is_aperio_svs && ifd_count >= 3 && level_count >= 3)
     {
+        // Check if IFDs form a pyramid structure (decreasing sizes)
         bool is_pyramid = true;
         for (size_t i = 1; i < std::min(size_t(3), level_count); ++i)
         {
             auto ifd_curr = tif->level_ifd(i);
-            auto ifd_prev = tif->level_ifd(i - 1);
+            auto ifd_prev = tif->level_ifd(i-1);
             if (ifd_curr->width() >= ifd_prev->width())
             {
                 is_pyramid = false;
@@ -129,7 +130,7 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
         if (is_pyramid)
         {
             #ifdef DEBUG
-            fmt::print("ℹ️  Detected pyramid structure → treating as multi-resolution TIFF\n");
+            fmt::print("ℹ️  Detected pyramid structure → treating as Aperio SVS/multi-resolution TIFF\n");
             #endif // DEBUG
             is_aperio_svs = true;
         }
@@ -138,13 +139,13 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
     // WSI formats commonly have multiple SubfileType=0 IFDs for pyramid levels.
     // Only enforce the single-main-image constraint for unrecognized Generic TIFFs.
     bool is_known_multi_ifd_format =
-        is_aperio_svs || is_philips_tiff ||
+        is_aperio_svs || is_philips_tiff || looks_like_ome ||
         tif->tiff_type() == cuslide::tiff::TiffType::Hamamatsu ||
         tif->tiff_type() == cuslide::tiff::TiffType::Leica ||
         tif->tiff_type() == cuslide::tiff::TiffType::Ventana ||
         tif->tiff_type() == cuslide::tiff::TiffType::Trestle ||
         tif->tiff_type() == cuslide::tiff::TiffType::OmeTiff ||
-        tif->tiff_type() == cuslide::tiff::TiffType::QpTiff;
+        tif->tiff_type() == cuslide::tiff::TiffType::Qptiff;
 
     if (!is_known_multi_ifd_format)
     {
@@ -159,11 +160,14 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
             }
         }
 
-        // Assume that the image has only one main (high resolution) image.
-        if (main_ifd_list.size() != 1)
+        // Legacy TIFF path expected exactly one main image, but modern
+        // multiplex/stacked TIFF variants may legitimately carry multiple
+        // SubfileType=0 IFDs. We only reject the clearly invalid case with
+        // no main image at all.
+        if (main_ifd_list.empty())
         {
             throw std::runtime_error(
-                fmt::format("This format has more than one image with Subfile Type 0 so cannot be loaded!"));
+                fmt::format("No main image (Subfile Type 0) found in TIFF file."));
         }
     }
 
@@ -181,31 +185,110 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
     std::pmr::vector<int64_t> shape(
         { level0_ifd->height(), level0_ifd->width(), level0_ifd->samples_per_pixel() }, &resource);
 
-    DLDataType dtype{ kDLUInt, 8, 1 };
+    const std::string& json_str = tif->metadata();
+    json parsed_metadata{};
+    bool has_ome_metadata = false;
+    std::vector<std::string> ome_channel_names;
+    int64_t ome_size_c = -1;
+    double ome_physical_size_x = 0.0;
+    double ome_physical_size_y = 0.0;
+    bool has_ome_physical_size_x = false;
+    bool has_ome_physical_size_y = false;
+    try
+    {
+        parsed_metadata = json::parse(json_str);
+        if (parsed_metadata.contains("ome"))
+        {
+            const auto& ome = parsed_metadata["ome"];
+            if (ome.contains("size_c"))
+            {
+                ome_size_c = ome["size_c"].get<int64_t>();
+            }
+            if (ome.contains("channel_names") && ome["channel_names"].is_array())
+            {
+                for (const auto& ch : ome["channel_names"])
+                {
+                    if (ch.is_string())
+                    {
+                        ome_channel_names.emplace_back(ch.get<std::string>());
+                    }
+                }
+            }
+            if (ome.contains("physical_size_x") && ome["physical_size_x"].is_number())
+            {
+                ome_physical_size_x = ome["physical_size_x"].get<double>();
+                has_ome_physical_size_x = true;
+            }
+            if (ome.contains("physical_size_y") && ome["physical_size_y"].is_number())
+            {
+                ome_physical_size_y = ome["physical_size_y"].get<double>();
+                has_ome_physical_size_y = true;
+            }
+            has_ome_metadata = true;
+        }
+    }
+    catch (const std::exception&)
+    {
+        // Metadata is best-effort; continue with robust TIFF defaults.
+    }
+
+    DLDataType dtype{ kDLUInt, static_cast<uint8_t>(std::max<uint32_t>(8, level0_ifd->bits_per_sample())), 1 };
 
     // TODO: Fill correct values for cucim::io::format::ImageMetadataDesc
-    uint8_t n_ch = level0_ifd->samples_per_pixel();
-    if (n_ch != 3)
+    uint8_t n_ch = static_cast<uint8_t>(level0_ifd->samples_per_pixel());
+    if (has_ome_metadata && ome_size_c > 0)
     {
-        // Image loaded by a slow-path(libtiff) always will have 4 channel
-        // (by TIFFRGBAImageGet() method in libtiff)
-        n_ch = 4;
-        shape[2] = 4;
+        n_ch = static_cast<uint8_t>(ome_size_c);
+        shape[2] = n_ch;
     }
-    std::pmr::vector<std::string_view> channel_names(&resource);
-    channel_names.reserve(n_ch);
-    if (n_ch == 3)
+    else
     {
-        channel_names.emplace_back(std::string_view{ "R" });
-        channel_names.emplace_back(std::string_view{ "G" });
-        channel_names.emplace_back(std::string_view{ "B" });
+        shape[2] = n_ch;
+    }
+
+    // Own the names as strings while building. ImageMetadata::channel_names
+    // takes string_views and may retain pointers into the metadata buffer, so
+    // we copy into `resource` once at the end rather than keeping views into
+    // this temporary storage.
+    std::vector<std::string> channel_name_storage;
+    channel_name_storage.reserve(n_ch);
+    if (!ome_channel_names.empty())
+    {
+        for (size_t i = 0; i < static_cast<size_t>(n_ch); ++i)
+        {
+            if (i < ome_channel_names.size())
+            {
+                channel_name_storage.push_back(ome_channel_names[i]);
+            }
+            else
+            {
+                channel_name_storage.push_back(fmt::format("C{}", i));
+            }
+        }
+    }
+    else if (n_ch == 3)
+    {
+        channel_name_storage = { "R", "G", "B" };
+    }
+    else if (n_ch == 4)
+    {
+        channel_name_storage = { "R", "G", "B", "A" };
     }
     else
     {
-        channel_names.emplace_back(std::string_view{ "R" });
-        channel_names.emplace_back(std::string_view{ "G" });
-        channel_names.emplace_back(std::string_view{ "B" });
-        channel_names.emplace_back(std::string_view{ "A" });
+        for (uint8_t i = 0; i < n_ch; ++i)
+        {
+            channel_name_storage.push_back(fmt::format("C{}", i));
+        }
+    }
+
+    std::pmr::vector<std::string_view> channel_names(&resource);
+    channel_names.reserve(channel_name_storage.size());
+    for (const auto& name : channel_name_storage)
+    {
+        char* buf = static_cast<char*>(resource.allocate(name.size() + 1, alignof(char)));
+        std::memcpy(buf, name.c_str(), name.size() + 1);
+        channel_names.emplace_back(buf, name.size());
     }
 
     // Spacing units
@@ -222,7 +305,15 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
     const auto x_resolution = level0_ifd->x_resolution();
     const auto y_resolution = level0_ifd->y_resolution();
 
-    switch (resolution_unit)
+    if (has_ome_physical_size_x && has_ome_physical_size_y)
+    {
+        spacing.emplace_back(static_cast<float>(ome_physical_size_y));
+        spacing.emplace_back(static_cast<float>(ome_physical_size_x));
+        spacing.emplace_back(1.0f);
+        spacing_units.emplace_back(micrometer_unit);
+        spacing_units.emplace_back(micrometer_unit);
+    }
+    else switch (resolution_unit)
     {
     case 1: // no absolute unit of measurement
         spacing.emplace_back(y_resolution);
@@ -305,7 +396,6 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo
     std::string_view raw_data{ image_description.empty() ? "" : image_description.c_str() };
 
     // Dynamically allocate memory for json_data (need to be freed manually);
-    const std::string& json_str = tif->metadata();
     char* json_data_ptr = static_cast<char*>(cucim_malloc(json_str.size() + 1));
     memcpy(json_data_ptr, json_str.data(), json_str.size() + 1);
     std::string_view json_data{ json_data_ptr, json_str.size() };
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.cpp
index 7cf42350a..6f08487e4 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.cpp
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.cpp
@@ -12,10 +12,12 @@
 
 #include <memory>
 #include <vector>
+#include <cstddef>
 #include <cstdlib>
 #include <cstring>
 #include <string>
 #include <stdexcept>
+#include <type_traits>
 #include <unistd.h>
 #include <mutex>
 #include <fmt/format.h>
@@ -215,6 +217,59 @@ class DecodeBuffer
     bool is_device_ = false;
 };
 
+struct DecodePixelSpec
+{
+    uint32_t num_channels = 3;
+    uint32_t bytes_per_sample = 1;
+    nvimgcodecSampleDataType_t sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT8;
+    nvimgcodecColorSpec_t color_spec = NVIMGCODEC_COLORSPEC_SRGB;
+};
+
+DecodePixelSpec resolve_decode_pixel_spec(const IfdInfo& ifd_info)
+{
+    DecodePixelSpec spec;
+    spec.num_channels = ifd_info.num_channels > 0 ? ifd_info.num_channels : 3;
+
+    const uint32_t bits = ifd_info.bits_per_sample > 0 ? ifd_info.bits_per_sample : 8;
+    if (bits > 8)
+    {
+        spec.bytes_per_sample = 2;
+        spec.sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT16;
+    }
+    else
+    {
+        spec.bytes_per_sample = 1;
+        spec.sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT8;
+    }
+
+    int64_t photometric = -1;
+    auto tag_it = ifd_info.tiff_tags.find("PHOTOMETRIC");
+    if (tag_it != ifd_info.tiff_tags.end())
+    {
+        std::visit(
+            [&](const auto& v)
+            {
+                using T = std::decay_t<decltype(v)>;
+                if constexpr (std::is_integral_v<T>)
+                {
+                    photometric = static_cast<int64_t>(v);
+                }
+            },
+            tag_it->second);
+    }
+
+    if (spec.num_channels == 1 || photometric == 0 || photometric == 1)
+    {
+        spec.color_spec = NVIMGCODEC_COLORSPEC_GRAY;
+    }
+    else
+    {
+        spec.color_spec = NVIMGCODEC_COLORSPEC_SRGB;
+    }
+
+    return spec;
+}
+
 // ============================================================================
 // IFD-Level Region Decoding (Primary Decode Function)
 // ============================================================================
@@ -331,24 +386,24 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info,
         output_image_info.struct_size = sizeof(nvimgcodecImageInfo_t);
         output_image_info.struct_next = nullptr;
 
-        // Use UNCHANGED to preserve original format (RGB or grayscale)
+        const DecodePixelSpec decode_spec = resolve_decode_pixel_spec(ifd_info);
+        // Use UNCHANGED to preserve original planar packing and channel ordering.
         output_image_info.sample_format = NVIMGCODEC_SAMPLEFORMAT_I_UNCHANGED;
-        output_image_info.color_spec = NVIMGCODEC_COLORSPEC_SRGB;
+        output_image_info.color_spec = decode_spec.color_spec;
         output_image_info.chroma_subsampling = NVIMGCODEC_SAMPLING_NONE;
         output_image_info.num_planes = 1;
         output_image_info.buffer_kind = buffer_kind;
 
-        uint32_t num_channels = ifd_info.num_channels > 0 ? ifd_info.num_channels : 3;
-        size_t row_stride = width * num_channels;
+        uint32_t num_channels = decode_spec.num_channels;
+        size_t row_stride = static_cast<size_t>(width) * num_channels * decode_spec.bytes_per_sample;
         size_t buffer_size = row_stride * height;
 
         output_image_info.plane_info[0].height = height;
         output_image_info.plane_info[0].width = width;
         output_image_info.plane_info[0].num_channels = num_channels;
         output_image_info.plane_info[0].row_stride = row_stride;
-        output_image_info.plane_info[0].sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT8;
-        // nvImageCodec >=0.7.0 removed buffer_size from nvimgcodecImageInfo_t;
-        // the required size is deterministic: row_stride * height per plane.
+        output_image_info.plane_info[0].sample_type = decode_spec.sample_type;
+        // Buffer size is inferred from plane_info rather than set explicitly.
         output_image_info.cuda_stream = cuda_stream;
 
         // Step 4: Provide output buffer
@@ -623,9 +678,9 @@ std::vector<BatchDecodeResult> decode_batch_regions_nvimgcodec(
             }
 
             const auto& region = regions[i];
-            // Use num_channels from IFD info, fallback to 3 (RGB)
-            uint32_t num_channels = (ifd_info.num_channels > 0) ? ifd_info.num_channels : 3;
-            size_t row_stride = region.width * num_channels;
+            const DecodePixelSpec decode_spec = resolve_decode_pixel_spec(ifd_info);
+            uint32_t num_channels = decode_spec.num_channels;
+            size_t row_stride = static_cast<size_t>(region.width) * num_channels * decode_spec.bytes_per_sample;
             size_t buffer_size = row_stride * region.height;
 
             if (!decode_buffers[i].allocate(buffer_size, use_device_memory))
@@ -641,9 +696,9 @@ std::vector<BatchDecodeResult> decode_batch_regions_nvimgcodec(
             output_image_info.struct_type = NVIMGCODEC_STRUCTURE_TYPE_IMAGE_INFO;
             output_image_info.struct_size = sizeof(nvimgcodecImageInfo_t);
             output_image_info.struct_next = nullptr;
-            // Use UNCHANGED to preserve original format (RGB or grayscale)
+            // Use UNCHANGED to preserve original planar packing and channel ordering.
             output_image_info.sample_format = NVIMGCODEC_SAMPLEFORMAT_I_UNCHANGED;
-            output_image_info.color_spec = NVIMGCODEC_COLORSPEC_SRGB;
+            output_image_info.color_spec = decode_spec.color_spec;
             output_image_info.chroma_subsampling = NVIMGCODEC_SAMPLING_NONE;
             output_image_info.num_planes = 1;
             output_image_info.buffer_kind = buffer_kind;
@@ -652,7 +707,7 @@ std::vector<BatchDecodeResult> decode_batch_regions_nvimgcodec(
             output_image_info.plane_info[0].width = region.width;
             output_image_info.plane_info[0].num_channels = num_channels;
             output_image_info.plane_info[0].row_stride = row_stride;
-            output_image_info.plane_info[0].sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT8;
+            output_image_info.plane_info[0].sample_type = decode_spec.sample_type;
             output_image_info.cuda_stream = cuda_stream;
 
             nvimgcodecImage_t image_raw = nullptr;
@@ -926,8 +981,9 @@ BatchDecodeState schedule_batch_decode(
             }
 
             const auto& region = regions[i];
-            uint32_t num_channels = (ifd_info.num_channels > 0) ? ifd_info.num_channels : 3;
-            size_t row_stride = region.width * num_channels;
+            const DecodePixelSpec decode_spec = resolve_decode_pixel_spec(ifd_info);
+            uint32_t num_channels = decode_spec.num_channels;
+            size_t row_stride = static_cast<size_t>(region.width) * num_channels * decode_spec.bytes_per_sample;
             size_t buffer_size = row_stride * region.height;
 
             void* buffer_ptr = nullptr;
@@ -960,7 +1016,7 @@ BatchDecodeState schedule_batch_decode(
             output_image_info.struct_size = sizeof(nvimgcodecImageInfo_t);
             output_image_info.struct_next = nullptr;
             output_image_info.sample_format = NVIMGCODEC_SAMPLEFORMAT_I_UNCHANGED;
-            output_image_info.color_spec = NVIMGCODEC_COLORSPEC_SRGB;
+            output_image_info.color_spec = decode_spec.color_spec;
             output_image_info.chroma_subsampling = NVIMGCODEC_SAMPLING_NONE;
             output_image_info.num_planes = 1;
             output_image_info.buffer_kind = buffer_kind;
@@ -969,7 +1025,7 @@ BatchDecodeState schedule_batch_decode(
             output_image_info.plane_info[0].width = region.width;
             output_image_info.plane_info[0].num_channels = num_channels;
             output_image_info.plane_info[0].row_stride = row_stride;
-            output_image_info.plane_info[0].sample_type = NVIMGCODEC_SAMPLE_DATA_TYPE_UINT8;
+            output_image_info.plane_info[0].sample_type = decode_spec.sample_type;
             output_image_info.cuda_stream = impl->cuda_stream;
 
             nvimgcodecImage_t image_raw = nullptr;
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.h
index 13ab9ffdf..8080389ee 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_decoder.h
@@ -123,7 +123,7 @@ struct BatchDecodeState
 /**
  * Decode multiple regions of interest (ROIs) from a single IFD in a TIFF file
  *
- * Uses nvImageCodec batch decoding API (v0.7.0+):
+ * Uses the nvImageCodec batch decoding API:
  * 1. Read image into CodeStream (main_code_stream)
  * 2. Call get_sub_code_stream() for each ROI with different regions
  * 3. Decode all ROIs in a single decoder.decode() call
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.cpp
index d9352b0dc..19be38603 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.cpp
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.cpp
@@ -12,8 +12,7 @@
 //
 // - Vendor-specific metadata blobs (MED_APERIO, MED_PHILIPS, ...)
 // - IFD (Image File Directory) enumeration for pyramidal TIFFs
-// - nvImageCodec 0.7.0+: direct TIFF tag queries via NVIMGCODEC_METADATA_KIND_TIFF_TAG
-// - nvImageCodec 0.8.0+: TIFF offset parsing, pagination, SubIFD support, TIFF_TAG_LIST
+// - Direct TIFF tag queries via NVIMGCODEC_METADATA_KIND_TIFF_TAG
 //   (COMPRESSION, SUBFILETYPE, IMAGEDESCRIPTION, JPEGTABLES, ...)
 //
 // ============================================================================
@@ -21,7 +20,7 @@
 #include "nvimgcodec_tiff_parser.h"
 
 #include <algorithm>  // for std::transform
-#include <cstdio>     // for stderr
+#include <cstddef>    // for size_t
 #include <cstdlib>    // for std::atexit, std::getenv, std::strtol
 #include <cstring>    // for strlen
 #include <thread>     // for std::thread::hardware_concurrency
@@ -54,7 +53,7 @@ namespace cuslide2::nvimgcodec
 // nvImageCodec extension path auto-detection
 // ============================================================================
 //
-// Workaround for a known nvImageCodec pre-0.8.0 bug: when the library is installed
+// Workaround for older nvImageCodec extension-path behavior: when the library is installed
 // via conda, its internal extension search path doesn't find the codec plugins
 // (libtiff_ext.so, libnvjpeg_ext.so, etc.) in $PREFIX/lib/extensions/.
 //
@@ -66,8 +65,8 @@ namespace cuslide2::nvimgcodec
 //      NvimgcodecLoadSymbol() + dladdr() (not the stub address)
 //   3. Probes <lib_dir>/extensions/ and <lib_dir>/../extensions/
 //
-// NOTE: nvImageCodec 0.8.0 fixes the extension search path for conda installs.
-// This workaround is kept as a safety net for non-standard install layouts.
+// Keep this path probe in place for compatibility across mixed environments,
+// including older deployments where extension auto-discovery is incomplete.
 //
 
 static std::string detect_nvimgcodec_extensions_path()
@@ -220,14 +219,6 @@ static int compute_max_decoder_threads()
 }
 
 
-// nvimgcodec API compatibility
-//
-// TIFF-tag retrieval via nvimgcodecDecoderGetMetadata.
-// The required enum values (NVIMGCODEC_METADATA_KIND_TIFF_TAG,
-// NVIMGCODEC_METADATA_VALUE_TYPE_ASCII, etc.) are defined in nvimgcodec.h
-// which is always present in our build tree.
-#define CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA 1
-
 // Helper: convert typed TIFF tag value to a string representation.
 static std::string tiff_tag_value_to_string(const TiffTagValue& value)
 {
@@ -273,6 +264,41 @@ static std::string tiff_tag_value_to_string(const TiffTagValue& value)
     }, value);
 }
 
+// RATIONAL/SRATIONAL tags arrive as numerator/denominator pairs. The pair count is
+// derived from the buffer size rather than value_count, which may report either the
+// number of rationals or the number of underlying LONGs.
+template <typename T>
+static bool extract_rational_value(const std::vector<uint8_t>& buffer, TiffTagValue& out_value)
+{
+    constexpr size_t pair_size = sizeof(T) * 2;
+    const size_t pair_count = buffer.size() / pair_size;
+    if (pair_count == 0)
+    {
+        return false;
+    }
+
+    const T* vals = reinterpret_cast<const T*>(buffer.data());
+    auto as_double = [](T numerator, T denominator) -> double {
+        return (denominator == 0) ? 0.0 : static_cast<double>(numerator) / static_cast<double>(denominator);
+    };
+
+    if (pair_count == 1)
+    {
+        out_value = as_double(vals[0], vals[1]);
+    }
+    else
+    {
+        std::vector<double> converted;
+        converted.reserve(pair_count);
+        for (size_t i = 0; i < pair_count; ++i)
+        {
+            converted.emplace_back(as_double(vals[2 * i], vals[2 * i + 1]));
+        }
+        out_value = std::move(converted);
+    }
+    return true;
+}
+
 // Unified extraction function: handles both single values and arrays.
 template <typename T>
 static bool extract_tag_value(const std::vector<uint8_t>& buffer, int value_count, TiffTagValue& out_value)
@@ -360,33 +386,46 @@ NvImageCodecTiffParserManager::NvImageCodecTiffParserManager()
         exec_params.backends = nullptr;
 
         status = nvimgcodecDecoderCreate(instance_, &decoder_, &exec_params, nullptr);
-
         if (status != NVIMGCODEC_STATUS_SUCCESS)
         {
-            nvimgcodecInstanceDestroy(instance_);
-            instance_ = nullptr;
-            status_message_ = fmt::format("Failed to create decoder for metadata extraction (status: {})",
-                                         static_cast<int>(status));
-            #ifdef DEBUG
-            fmt::print("⚠️  {}\n", status_message_);
-            #endif // DEBUG
-            return;
+            // If CUDA initialization is unavailable (for example, no GPU device
+            // in the current runtime), retry with CPU-only decoder so metadata
+            // extraction and CPU decode paths can still run.
+            exec_params.device_id = NVIMGCODEC_DEVICE_CPU_ONLY;
+            status = nvimgcodecDecoderCreate(instance_, &decoder_, &exec_params, nullptr);
+            if (status != NVIMGCODEC_STATUS_SUCCESS)
+            {
+                nvimgcodecInstanceDestroy(instance_);
+                instance_ = nullptr;
+                status_message_ = fmt::format("Failed to create decoder for metadata extraction (status: {})",
+                                             static_cast<int>(status));
+                #ifdef DEBUG
+                fmt::print("⚠️  {}\n", status_message_);
+                #endif // DEBUG
+                return;
+            }
+            status_message_ = "nvImageCodec TIFF parser initialized in CPU-only mode";
         }
 
         initialized_ = true;
-        status_message_ = "nvImageCodec TIFF parser initialized successfully (with metadata extraction support)";
+        if (status_message_.empty())
+        {
+            status_message_ = "nvImageCodec TIFF parser initialized successfully (with metadata extraction support)";
+        }
         #ifdef DEBUG
         fmt::print("✅ {}\n", status_message_);
         #endif // DEBUG
 
-        // Register atexit() AFTER nvimgcodecDecoderCreate (which initializes
-        // CUDA internally).  atexit handlers run in LIFO order, so this
-        // handler fires BEFORE the CUDA runtime's own atexit cleanup,
-        // guaranteeing that the CUDA context is still alive when we call
-        // nvimgcodecDecoderDestroy / nvimgcodecInstanceDestroy.
-        std::atexit([]() {
-            NvImageCodecTiffParserManager::instance().shutdown();
-        });
+        // The decoder and instance are intentionally not torn down at process exit.
+        // Registering an atexit handler after nvimgcodecDecoderCreate does not make
+        // the teardown safe: the CUDA driver releases its own state on a schedule
+        // that cannot be ordered against libc exit handlers, so
+        // nvimgcodecDecoderDestroy intermittently frees pointers the driver has
+        // already reclaimed and aborts in free() inside libcuda. Both orderings
+        // were observed to abort. These are process-lifetime singletons, so letting
+        // the OS and the driver reclaim them at exit is the reliable choice.
+        // shutdown() remains available for callers that need to release the decoder
+        // deterministically while CUDA is still up.
     }
     catch (const std::exception& e)
     {
@@ -430,9 +469,10 @@ void NvImageCodecTiffParserManager::shutdown()
 
 NvImageCodecTiffParserManager::~NvImageCodecTiffParserManager()
 {
-    // shutdown() is normally called via atexit() while the CUDA context is
-    // still alive.  The destructor calls it again as a safety net — it is
-    // idempotent, so the second call is a harmless no-op.
+    // Not reached at process exit: instance() intentionally leaks the singleton so
+    // that teardown is driven solely by the atexit() handler, which is sequenced
+    // before the CUDA driver's own cleanup. Kept for completeness if an instance is
+    // ever destroyed explicitly while CUDA is still up. shutdown() is idempotent.
     shutdown();
 }
 
@@ -547,6 +587,11 @@ bool TiffFileParser::parse_tiff_structure()
         #endif // DEBUG
     }
 
+    int substream_failures = 0;
+    int image_info_failures = 0;
+    int first_substream_status = 0;
+    int first_image_info_status = 0;
+
     // Get information for each IFD
     for (uint32_t i = 0; i < num_ifds; ++i)
     {
@@ -559,6 +604,11 @@ bool TiffFileParser::parse_tiff_structure()
         view.struct_size = sizeof(nvimgcodecCodeStreamView_t);
         view.struct_next = nullptr;
         view.image_idx = i;  // Note: nvImageCodec uses 'image_idx' not 'image_index'
+        view.bitstream_offset = 0;
+        view.region.struct_type = NVIMGCODEC_STRUCTURE_TYPE_REGION;
+        view.region.struct_size = sizeof(nvimgcodecRegion_t);
+        view.region.struct_next = nullptr;
+        view.region.ndim = 0; // No ROI for parser-time IFD discovery.
 
         // Get sub-code stream for this IFD
         status = nvimgcodecCodeStreamGetSubCodeStream(main_code_stream_,
@@ -567,6 +617,11 @@ bool TiffFileParser::parse_tiff_structure()
 
         if (status != NVIMGCODEC_STATUS_SUCCESS)
         {
+            ++substream_failures;
+            if (first_substream_status == 0)
+            {
+                first_substream_status = static_cast<int>(status);
+            }
             #ifdef DEBUG
             fmt::print("❌ Failed to get sub-code stream for IFD {} (status: {})\n",
                       i, static_cast<int>(status));
@@ -589,6 +644,11 @@ bool TiffFileParser::parse_tiff_structure()
 
         if (status != NVIMGCODEC_STATUS_SUCCESS)
         {
+            ++image_info_failures;
+            if (first_image_info_status == 0)
+            {
+                first_image_info_status = static_cast<int>(status);
+            }
             #ifdef DEBUG
             fmt::print("❌ Failed to get image info for IFD {} (status: {})\n",
                       i, static_cast<int>(status));
@@ -631,9 +691,8 @@ bool TiffFileParser::parse_tiff_structure()
         // Extract TIFF metadata using available methods
         extract_tiff_tags(ifd_info);
 
-        // codec_name returns "tiff" (container format), not compression type.
-        // Individual TIFF tags are queried via NVIMGCODEC_METADATA_KIND_TIFF_TAG,
-        // and TIFF_TAG_LIST enumeration is available in v0.8.0+.
+        // codec_name reports "tiff" (the container format) rather than the
+        // compression type, so compression is inferred from the TIFF tags below.
 
         if (ifd_info.codec == "tiff")
         {
@@ -750,6 +809,19 @@ bool TiffFileParser::parse_tiff_structure()
         #endif // DEBUG
     }
 
+    if (ifd_infos_.empty())
+    {
+        parse_error_ = fmt::format(
+            "Failed to parse TIFF IFDs via nvImageCodec: num_images={}, "
+            "substream_failures={} (first_status={}), image_info_failures={} (first_status={})",
+            num_ifds,
+            substream_failures,
+            first_substream_status,
+            image_info_failures,
+            first_image_info_status);
+        return false;
+    }
+
     return true;
 }
 
@@ -1005,11 +1077,10 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info)
     }
 
     // ========================================================================
-    // Direct TIFF Tag Retrieval by ID (nvImageCodec 0.7.0+)
+    // Direct TIFF Tag Retrieval by ID
     // ========================================================================
     // Query a fixed set of common TIFF tags individually. Not all tags exist on all IFDs.
 
-#if CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA
     std::vector<std::pair<uint16_t, std::string>> tiff_tags_to_query = {
         {254, "SUBFILETYPE"},
         {256, "IMAGEWIDTH"},
@@ -1021,6 +1092,9 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info)
         {271, "MAKE"},
         {272, "MODEL"},
         {277, "SAMPLESPERPIXEL"},
+        {282, "XRESOLUTION"},
+        {283, "YRESOLUTION"},
+        {296, "RESOLUTIONUNIT"},
         {305, "SOFTWARE"},
         {306, "DATETIME"},
         {322, "TILEWIDTH"},
@@ -1106,6 +1180,12 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info)
             case NVIMGCODEC_METADATA_VALUE_TYPE_SLONG:
                 extract_tag_value<int32_t>(buffer, metadata.value_count, tag_value);
                 break;
+            case NVIMGCODEC_METADATA_VALUE_TYPE_RATIONAL:
+                extract_rational_value<uint32_t>(buffer, tag_value);
+                break;
+            case NVIMGCODEC_METADATA_VALUE_TYPE_SRATIONAL:
+                extract_rational_value<int32_t>(buffer, tag_value);
+                break;
             case NVIMGCODEC_METADATA_VALUE_TYPE_LONG8:
             case NVIMGCODEC_METADATA_VALUE_TYPE_IFD8:
                 extract_tag_value<uint64_t>(buffer, metadata.value_count, tag_value);
@@ -1139,6 +1219,20 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info)
         }
     }
 
+    // Tag retrieval goes through the nvImageCodec decoder, so it yields nothing when no
+    // decoder can be created (for example, no usable CUDA device). Without tags there is
+    // no ImageDescription, which silently disables OME and vendor format detection.
+    // Report it once per file so the degraded metadata is traceable.
+    if (extracted_count == 0 && !tag_extraction_warned_)
+    {
+        tag_extraction_warned_ = true;
+        fmt::print(stderr,
+                   "[cuslide2] No TIFF tags returned by nvImageCodec for '{}'. "
+                   "OME/vendor metadata and pixel spacing will be unavailable. "
+                   "Tag retrieval requires a working decoder (CUDA device).\n",
+                   file_path_);
+    }
+
     // Populate image_description if it wasn't already filled via vendor metadata.
     auto desc_it = ifd_info.tiff_tags.find("IMAGEDESCRIPTION");
     if (desc_it != ifd_info.tiff_tags.end() && ifd_info.image_description.empty())
@@ -1151,9 +1245,8 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info)
     {
         return;  // Have compression info, no need for heuristics
     }
-#endif // CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA
 
-    // Fallback: file extension heuristics when COMPRESSION tag is not present in file
+    // Fallback: file extension heuristics when the file carries no COMPRESSION tag
 #ifdef DEBUG
     fmt::print("  ℹ️  COMPRESSION tag not available, using file extension heuristics\n");
 #endif // DEBUG
@@ -1200,13 +1293,11 @@ std::vector<int> TiffFileParser::query_metadata_kinds(uint32_t ifd_index) const
         kinds.push_back(kind);
     }
 
-    // Also include TIFF_TAG kind if any tags were extracted (only supported on newer nvimgcodec).
-#if CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA
+    // Also include TIFF_TAG kind if any tags were extracted.
     if (!ifd_infos_[ifd_index].tiff_tags.empty())
     {
         kinds.insert(kinds.begin(), NVIMGCODEC_METADATA_KIND_TIFF_TAG);
     }
-#endif
 
     return kinds;
 }
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.h
index 5bd7ac52d..64925b869 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.h
@@ -89,7 +89,7 @@ struct IfdInfo
     };
     std::map<int, MetadataBlob> metadata_blobs;
 
-    // nvImageCodec 0.7.0+: Individual TIFF tag storage with typed values (0.8.0+ adds TIFF_TAG_LIST)
+    // Individual TIFF tag storage with typed values.
     // tag_name -> TiffTagValue (variant with typed storage)
     std::unordered_map<std::string, TiffTagValue> tiff_tags;
 
@@ -198,7 +198,7 @@ class TiffFileParser
     /**
      * @brief Get a specific TIFF tag value as string
      *
-     * Returns TIFF tags queried via nvImageCodec metadata API (v0.7.0+)
+     * Returns TIFF tags queried via the nvImageCodec metadata API,
      * or inferred from file extension and vendor metadata as fallback.
      *
      * @param ifd_index IFD index
@@ -310,6 +310,9 @@ class TiffFileParser
 
     // Configuration: Maximum size for binary TIFF tag data (0 = unlimited)
     size_t max_binary_tag_size_ = 0;
+
+    // Limits the "no TIFF tags returned" diagnostic to one message per file.
+    bool tag_extraction_warned_ = false;
 };
 
 /**
@@ -328,8 +331,16 @@ class NvImageCodecTiffParserManager
      */
     static NvImageCodecTiffParserManager& instance()
     {
-        static NvImageCodecTiffParserManager manager;
-        return manager;
+        // Deliberately never destroyed. A static local would register its destructor
+        // via __cxa_atexit *after* the constructor registers shutdown() with
+        // std::atexit, and exit handlers run in reverse registration order, so the
+        // destructor would always run first and tear down CUDA-backed decoder state
+        // at an unsequenced point relative to the CUDA driver's own teardown. That
+        // races and aborts in free() inside libcuda. Leaking the manager leaves
+        // shutdown() to the atexit handler, which is registered after CUDA is
+        // initialized and therefore runs while the driver is still alive.
+        static NvImageCodecTiffParserManager* manager = new NvImageCodecTiffParserManager();
+        return *manager;
     }
 
     /**
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp
index 9bb888c78..51664db7b 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp
@@ -14,6 +14,7 @@
 #include <memory>
 #include <random>
 #include <thread>
+#include <cstring>
 #include <fstream>
 #include <langinfo.h>
 
@@ -112,7 +113,12 @@ IFD::IFD(TIFF* tiff, uint16_t index, ifd_offset_t offset) : tiff_(tiff), ifd_ind
 // ============================================================================
 
 #ifdef CUCIM_HAS_NVIMGCODEC
-IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_info)
+IFD::IFD(TIFF* tiff,
+         uint16_t index,
+         const ::cuslide2::nvimgcodec::IfdInfo& ifd_info,
+         size_t parser_idx,
+         uint32_t local_ifd_idx,
+         uint64_t file_hash)
     : tiff_(tiff), ifd_index_(index), ifd_offset_(index)
 {
     PROF_SCOPED_RANGE(PROF_EVENT(ifd_ifd));  // Use standard ifd_ifd profiler event
@@ -126,6 +132,8 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_
     height_ = ifd_info.height;
     samples_per_pixel_ = ifd_info.num_channels;
     bits_per_sample_ = ifd_info.bits_per_sample;
+    parser_index_ = parser_idx;
+    local_ifd_index_ = local_ifd_idx;
 
     #ifdef DEBUG
     fmt::print("   Dimensions: {}x{}, {} channels, {} bits/sample\n",
@@ -142,14 +150,15 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_
     // Get ImageDescription from nvImageCodec
     image_description_ = ifd_info.image_description;
 
-    // Extract TIFF tags from TiffFileParser
-    if (tiff->nvimgcodec_parser_) {
+    // Extract TIFF tags from the parser that owns this IFD.
+    const auto& parser = (parser_idx == 0) ? *tiff->nvimgcodec_parser_ : *tiff->companion_parsers_[parser_idx - 1];
+    {
         // Software and Model tags
-        software_ = tiff->nvimgcodec_parser_->get_tiff_tag(index, "SOFTWARE");
-        model_ = tiff->nvimgcodec_parser_->get_tiff_tag(index, "MODEL");
+        software_ = parser.get_tiff_tag(local_ifd_idx, "SOFTWARE");
+        model_ = parser.get_tiff_tag(local_ifd_idx, "MODEL");
 
         // SUBFILETYPE for IFD classification
-        int subfile_type = tiff->nvimgcodec_parser_->get_subfile_type(index);
+        int subfile_type = parser.get_subfile_type(local_ifd_idx);
         if (subfile_type >= 0) {
             subfile_type_ = static_cast<uint64_t>(subfile_type);
             #ifdef DEBUG
@@ -158,7 +167,7 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_
         }
 
         // Check for JPEGTables (abbreviated JPEG indicator)
-        std::string jpeg_tables = tiff->nvimgcodec_parser_->get_tiff_tag(index, "JPEGTABLES");
+        std::string jpeg_tables = parser.get_tiff_tag(local_ifd_idx, "JPEGTABLES");
         if (!jpeg_tables.empty()) {
             #ifdef DEBUG
             fmt::print("   ✅ JPEGTables detected (abbreviated JPEG)\n");
@@ -166,8 +175,9 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_
         }
 
         // Tile dimensions (if available from TIFF tags)
-        std::string tile_w_str = tiff->nvimgcodec_parser_->get_tiff_tag(index, "TILEWIDTH");
-        std::string tile_h_str = tiff->nvimgcodec_parser_->get_tiff_tag(index, "TILELENGTH");
+        std::string tile_w_str = parser.get_tiff_tag(local_ifd_idx, "TILEWIDTH");
+        std::string tile_h_str = parser.get_tiff_tag(local_ifd_idx, "TILELENGTH");
+        std::string photo_str = parser.get_tiff_tag(local_ifd_idx, "PHOTOMETRIC");
 
         if (!tile_w_str.empty() && !tile_h_str.empty()) {
             try {
@@ -191,20 +201,81 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_
             fmt::print("   Not tiled (strip-based or whole image)\n");
             #endif
         }
+
+        if (!photo_str.empty())
+        {
+            try
+            {
+                photometric_ = static_cast<uint16_t>(std::stoul(photo_str));
+            }
+            catch (...)
+            {
+                // Keep fallback below.
+            }
+        }
+
+        // Resolution tags drive the pixel spacing reported to callers. XRESOLUTION and
+        // YRESOLUTION are RATIONAL tags already reduced to a decimal by the parser.
+        // Missing or unparsable tags leave the declared defaults in place.
+        std::string res_unit_str = parser.get_tiff_tag(local_ifd_idx, "RESOLUTIONUNIT");
+        std::string x_res_str = parser.get_tiff_tag(local_ifd_idx, "XRESOLUTION");
+        std::string y_res_str = parser.get_tiff_tag(local_ifd_idx, "YRESOLUTION");
+
+        if (!res_unit_str.empty())
+        {
+            try
+            {
+                resolution_unit_ = static_cast<uint16_t>(std::stoul(res_unit_str));
+            }
+            catch (...)
+            {
+            }
+        }
+        if (!x_res_str.empty())
+        {
+            try
+            {
+                const float parsed = std::stof(x_res_str);
+                if (parsed > 0.0f)
+                {
+                    x_resolution_ = parsed;
+                }
+            }
+            catch (...)
+            {
+            }
+        }
+        if (!y_res_str.empty())
+        {
+            try
+            {
+                const float parsed = std::stof(y_res_str);
+                if (parsed > 0.0f)
+                {
+                    y_resolution_ = parsed;
+                }
+            }
+            catch (...)
+            {
+            }
+        }
+        #ifdef DEBUG
+        fmt::print("   Resolution: {}x{} (unit={})\n", x_resolution_, y_resolution_, resolution_unit_);
+        #endif
     }
 
     // Set format defaults
     planar_config_ = cuslide::tiff::PLANARCONFIG_CONTIG;  // nvImageCodec outputs interleaved
-    photometric_ = cuslide::tiff::PHOTOMETRIC_RGB;
+    if (photometric_ == 0)
+    {
+        photometric_ = (samples_per_pixel_ == 1) ? cuslide::tiff::PHOTOMETRIC_MINISBLACK
+                                                  : cuslide::tiff::PHOTOMETRIC_RGB;
+    }
     predictor_ = 1;  // No predictor
 
-    // Resolution info (defaults - may not be available from nvImageCodec)
-    resolution_unit_ = 1;  // No absolute unit
-    x_resolution_ = 1.0f;
-    y_resolution_ = 1.0f;
-
     // Calculate hash for caching (include file hash for cross-file uniqueness)
-    hash_value_ = tiff->file_handle_shared_.get()->hash_value ^ cucim::codec::splitmix64(index);
+    hash_value_ = (file_hash != 0) ? (file_hash ^ cucim::codec::splitmix64(local_ifd_idx))
+                                   : (tiff->file_handle_shared_.get()->hash_value ^ cucim::codec::splitmix64(index));
 
 #ifdef CUCIM_HAS_NVIMGCODEC
     // Store reference to nvImageCodec sub-stream
@@ -254,7 +325,7 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
     #endif
 
 #ifdef CUCIM_HAS_NVIMGCODEC
-    if (!nvimgcodec_sub_stream_ || !tiff->nvimgcodec_parser_)
+    if (!nvimgcodec_sub_stream_)
     {
         throw std::runtime_error(fmt::format(
             "IFD[{}]: nvImageCodec parser not available", ifd_index_));
@@ -272,10 +343,11 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
     const uint64_t location_len = request->location_len;
 
     uint32_t batch_size = request->batch_size;
-    int32_t n_ch = samples_per_pixel_;
+    int32_t n_ch = static_cast<int32_t>(std::max<uint32_t>(1, samples_per_pixel_));
     int ndim = 3;
+    const size_t bytes_per_sample = std::max<size_t>(1, bits_per_sample_ / 8);
 
-    size_t one_raster_size = static_cast<size_t>(w) * static_cast<size_t>(h) * samples_per_pixel_;
+    size_t one_raster_size = static_cast<size_t>(w) * static_cast<size_t>(h) * n_ch * bytes_per_sample;
     size_t raster_size = one_raster_size;
 
     void* raster = nullptr;
@@ -419,7 +491,7 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
         // loader->next_data() during Python iteration.
         out_image_data->container.data = raster;  // may be nullptr for batch iteration
         out_image_data->container.device = DLDevice{ static_cast<DLDeviceType>(raster_type), out_device.index() };
-        out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 };
+        out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast<uint8_t>(bits_per_sample_), 1 };
         out_image_data->container.ndim = ndim;
         out_image_data->container.shape = static_cast<int64_t*>(cucim_malloc(ndim * sizeof(int64_t)));
         if (ndim == 4)
@@ -456,9 +528,10 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
         output_buffer = static_cast<uint8_t*>(out_buf->data);
     }
 
-    // Get IFD info from TiffFileParser
-    const auto& ifd_info = tiff->nvimgcodec_parser_->get_ifd(static_cast<uint32_t>(ifd_index_));
-    auto main_code_stream = tiff->nvimgcodec_parser_->get_main_code_stream();
+    // Get IFD info from the parser that owns this global IFD.
+    const auto& parser = tiff->parser_for_global_ifd(ifd_index_);
+    const auto& ifd_info = parser.get_ifd(local_ifd_index_);
+    auto main_code_stream = parser.get_main_code_stream();
 
     // ====================================================================
     // TILE-LEVEL CACHING PATH
@@ -469,13 +542,67 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
     int64_t ex = sx + w - 1;
     int64_t ey = sy + h - 1;
 
+    // A request can fall entirely outside the image. nvImageCodec fills out-of-bounds
+    // pixels for a region that partially overlaps, but rejects one with no
+    // intersection at all, which would surface as a decode failure. read_region()
+    // is specified to return an all-background raster of the requested size here, so
+    // produce it directly instead of consulting the decoder.
+    if (ex < 0 || ey < 0 || sx >= static_cast<int64_t>(width_) || sy >= static_cast<int64_t>(height_))
+    {
+        if (!output_buffer)
+        {
+            if (out_device.type() == cucim::io::DeviceType::kCUDA)
+            {
+                if (cudaMalloc(reinterpret_cast<void**>(&output_buffer), one_raster_size) != cudaSuccess)
+                {
+                    throw std::runtime_error("Failed to allocate GPU buffer for out-of-bounds region");
+                }
+            }
+            else
+            {
+                output_buffer = static_cast<uint8_t*>(cucim_malloc(one_raster_size));
+                if (!output_buffer)
+                {
+                    throw std::runtime_error("Failed to allocate host buffer for out-of-bounds region");
+                }
+            }
+        }
+
+        if (out_device.type() == cucim::io::DeviceType::kCUDA)
+        {
+            if (cudaMemset(output_buffer, 0, one_raster_size) != cudaSuccess)
+            {
+                throw std::runtime_error("Failed to clear GPU buffer for out-of-bounds region");
+            }
+        }
+        else
+        {
+            std::memset(output_buffer, 0, one_raster_size);
+        }
+
+        out_image_data->container.data = output_buffer;
+        out_image_data->container.device =
+            DLDevice{ static_cast<DLDeviceType>(out_device.type()), out_device.index() };
+        out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast<uint8_t>(bits_per_sample_), 1 };
+        out_image_data->container.ndim = 3;
+        out_image_data->container.shape = static_cast<int64_t*>(cucim_malloc(3 * sizeof(int64_t)));
+        out_image_data->container.shape[0] = h;
+        out_image_data->container.shape[1] = w;
+        out_image_data->container.shape[2] = n_ch;
+        out_image_data->container.strides = nullptr;
+        out_image_data->container.byte_offset = 0;
+
+        return true;
+    }
+
     // Tile caching is applicable when:
     //  1. The image is tiled (tile_width_ > 0 && tile_height_ > 0)
     //  2. The ROI is fully within the image bounds (no boundary handling)
     //  3. An image cache is configured (not kNoCache)
-    //  4. The image is 8-bit RGB (the only format the assembly path supports)
+    //  4. The image has a scalar sample format this assembly path supports
     bool use_tile_caching = (tile_width_ > 0 && tile_height_ > 0 &&
-                             bits_per_sample_ == 8 && samples_per_pixel_ == 3 &&
+                             (bits_per_sample_ == 8 || bits_per_sample_ == 16) &&
+                             samples_per_pixel_ >= 1 &&
                              sx >= 0 && sy >= 0 &&
                              ex < static_cast<int64_t>(width_) &&
                              ey < static_cast<int64_t>(height_));
@@ -497,7 +624,8 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
 
         const uint32_t tw = tile_width_;
         const uint32_t th = tile_height_;
-        const uint32_t samples = samples_per_pixel_;
+        const uint32_t samples = std::max<uint32_t>(1, samples_per_pixel_);
+        const uint32_t bytes_per_px = static_cast<uint32_t>(std::max<size_t>(1, bits_per_sample_ / 8));
         const uint8_t background_value = tiff->background_value_;
 
         // Tile grid offsets that the ROI overlaps
@@ -542,7 +670,7 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
             host_raster = host_raster_owner.get();
         }
 
-        const uint32_t dest_row_stride = static_cast<uint32_t>(w) * samples;
+        const uint32_t dest_row_stride = static_cast<uint32_t>(w) * samples * bytes_per_px;
         uint8_t* dest_row_ptr = host_raster;
 
         #ifdef DEBUG
@@ -567,15 +695,15 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
                 const uint32_t tile_pixel_sx = (tile_col == offset_sx) ? pixel_offset_sx : 0;
                 const uint32_t tile_pixel_ex = (tile_col == offset_ex) ? pixel_offset_ex : (tw - 1);
                 const uint32_t copy_cols = tile_pixel_ex - tile_pixel_sx + 1;
-                const uint32_t copy_bytes_per_row = copy_cols * samples;
+                const uint32_t copy_bytes_per_row = copy_cols * samples * bytes_per_px;
 
                 // Actual decoded tile dimensions (clipped to image bounds for edge tiles)
                 const uint32_t tile_origin_x = tile_col * tw;
                 const uint32_t tile_origin_y = tile_row * th;
                 const uint32_t actual_tw = std::min(tw, width_ - tile_origin_x);
                 const uint32_t actual_th = std::min(th, height_ - tile_origin_y);
-                const size_t tile_buf_size = static_cast<size_t>(actual_tw) * actual_th * samples;
-                const uint32_t tile_row_stride = actual_tw * samples;
+                const size_t tile_buf_size = static_cast<size_t>(actual_tw) * actual_th * samples * bytes_per_px;
+                const uint32_t tile_row_stride = actual_tw * samples * bytes_per_px;
 
                 // Hash for per-tile locking
                 const uint64_t index_hash = ifd_hash ^
@@ -677,7 +805,7 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
                 // If insert() succeeded, the cache also holds a shared reference.
                 // If insert() failed, tile_value is the sole owner — freed after this iteration.
                 const uint32_t src_byte_offset =
-                    (tile_pixel_sy * actual_tw + tile_pixel_sx) * samples;
+                    (tile_pixel_sy * actual_tw + tile_pixel_sx) * samples * bytes_per_px;
 
                 for (uint32_t r = 0; r < copy_rows; ++r)
                 {
@@ -733,12 +861,12 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
         // Set up output metadata
         out_image_data->container.data = output_buffer;
         out_image_data->container.device = DLDevice{ static_cast<DLDeviceType>(raster_type), out_device.index() };
-        out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 };
+        out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast<uint8_t>(bits_per_sample_), 1 };
         out_image_data->container.ndim = 3;
         out_image_data->container.shape = static_cast<int64_t*>(cucim_malloc(3 * sizeof(int64_t)));
         out_image_data->container.shape[0] = h;
         out_image_data->container.shape[1] = w;
-        out_image_data->container.shape[2] = samples_per_pixel_;
+        out_image_data->container.shape[2] = std::max<uint32_t>(1, samples_per_pixel_);
         out_image_data->container.strides = nullptr;
         out_image_data->container.byte_offset = 0;
 
@@ -776,12 +904,12 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff,
         // Set up output metadata
         out_image_data->container.data = output_buffer;
         out_image_data->container.device = DLDevice{ static_cast<DLDeviceType>(out_device.type()), out_device.index() };
-        out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 };
+        out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast<uint8_t>(bits_per_sample_), 1 };
         out_image_data->container.ndim = 3;
         out_image_data->container.shape = static_cast<int64_t*>(cucim_malloc(3 * sizeof(int64_t)));
         out_image_data->container.shape[0] = h;
         out_image_data->container.shape[1] = w;
-        out_image_data->container.shape[2] = samples_per_pixel_;
+        out_image_data->container.shape[2] = std::max<uint32_t>(1, samples_per_pixel_);
         out_image_data->container.strides = nullptr;
         out_image_data->container.byte_offset = 0;
 
@@ -1002,9 +1130,13 @@ bool IFD::is_compression_supported() const
 
 bool IFD::is_read_optimizable() const
 {
-    return is_compression_supported() && bits_per_sample_ == 8 && samples_per_pixel_ == 3 &&
+    return is_compression_supported() && (bits_per_sample_ == 8 || bits_per_sample_ == 16) &&
+           samples_per_pixel_ >= 1 &&
            (tile_width_ != 0 && tile_height_ != 0) && planar_config_ == cuslide::tiff::PLANARCONFIG_CONTIG &&
-           (photometric_ == cuslide::tiff::PHOTOMETRIC_RGB || photometric_ == cuslide::tiff::PHOTOMETRIC_YCBCR) &&
+           (photometric_ == cuslide::tiff::PHOTOMETRIC_RGB ||
+            photometric_ == cuslide::tiff::PHOTOMETRIC_YCBCR ||
+            photometric_ == cuslide::tiff::PHOTOMETRIC_MINISBLACK ||
+            photometric_ == cuslide::tiff::PHOTOMETRIC_MINISWHITE) &&
            !tiff_->is_in_read_config(TIFF::kUseLibTiff);
 }
 
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h
index ec2e24b72..2410eb9fe 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -34,7 +34,12 @@ class EXPORT_VISIBLE IFD : public std::enable_shared_from_this<IFD>
 public:
     IFD(TIFF* tiff, uint16_t index, ifd_offset_t offset);
 #ifdef CUCIM_HAS_NVIMGCODEC
-    IFD(TIFF* tiff, uint16_t index, const cuslide2::nvimgcodec::IfdInfo& ifd_info);
+    IFD(TIFF* tiff,
+        uint16_t index,
+        const cuslide2::nvimgcodec::IfdInfo& ifd_info,
+        size_t parser_idx = 0,
+        uint32_t local_ifd_idx = 0,
+        uint64_t file_hash = 0);
 #endif
     ~IFD();
 
@@ -121,6 +126,8 @@ class EXPORT_VISIBLE IFD : public std::enable_shared_from_this<IFD>
     // nvImageCodec-specific members
     nvimgcodecCodeStream_t nvimgcodec_sub_stream_ = nullptr;
     std::string codec_name_;  // codec name from nvImageCodec (jpeg, jpeg2k, deflate, etc.)
+    size_t parser_index_ = 0;
+    uint32_t local_ifd_index_ = 0;
 #endif
 
     /**
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp
new file mode 100644
index 000000000..2e3e52bf4
--- /dev/null
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp
@@ -0,0 +1,170 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "ome_xml.h"
+
+#include <algorithm>
+#include <string_view>
+
+#include <pugixml.hpp>
+
+namespace cuslide::tiff::ome
+{
+
+namespace
+{
+std::string_view local_name(const char* name)
+{
+    if (!name)
+    {
+        return {};
+    }
+    std::string_view value(name);
+    auto pos = value.find(':');
+    return (pos == std::string_view::npos) ? value : value.substr(pos + 1);
+}
+
+pugi::xml_node find_child_by_local_name(const pugi::xml_node& node, std::string_view name)
+{
+    for (auto child : node.children())
+    {
+        if (local_name(child.name()) == name)
+        {
+            return child;
+        }
+    }
+    return {};
+}
+} // namespace
+
+bool parse(const std::string& xml_text, Model* out_model, std::string* out_error)
+{
+    if (!out_model)
+    {
+        if (out_error)
+        {
+            *out_error = "Output model pointer is null";
+        }
+        return false;
+    }
+
+    *out_model = Model{};
+
+    pugi::xml_document doc;
+    pugi::xml_parse_result parsed = doc.load_string(xml_text.c_str());
+    if (!parsed)
+    {
+        if (out_error)
+        {
+            *out_error = parsed.description();
+        }
+        return false;
+    }
+
+    pugi::xml_node ome_node = doc.document_element();
+    if (local_name(ome_node.name()) != "OME")
+    {
+        if (out_error)
+        {
+            *out_error = "XML root is not OME";
+        }
+        return false;
+    }
+
+    pugi::xml_node image_node = find_child_by_local_name(ome_node, "Image");
+    if (!image_node)
+    {
+        if (out_error)
+        {
+            *out_error = "OME Image node not found";
+        }
+        return false;
+    }
+
+    pugi::xml_node pixels_node = find_child_by_local_name(image_node, "Pixels");
+    if (!pixels_node)
+    {
+        if (out_error)
+        {
+            *out_error = "OME Pixels node not found";
+        }
+        return false;
+    }
+
+    Model model;
+    model.image_id = image_node.attribute("ID").as_string("");
+    model.image_name = image_node.attribute("Name").as_string("");
+
+    Pixels pixels;
+    pixels.size_x = pixels_node.attribute("SizeX").as_uint(0);
+    pixels.size_y = pixels_node.attribute("SizeY").as_uint(0);
+    pixels.size_c = std::max<uint32_t>(1, pixels_node.attribute("SizeC").as_uint(1));
+    pixels.size_z = std::max<uint32_t>(1, pixels_node.attribute("SizeZ").as_uint(1));
+    pixels.size_t = std::max<uint32_t>(1, pixels_node.attribute("SizeT").as_uint(1));
+    pixels.type = pixels_node.attribute("Type").as_string("");
+    pixels.dimension_order = pixels_node.attribute("DimensionOrder").as_string("XYCZT");
+
+    if (auto attr = pixels_node.attribute("PhysicalSizeX"))
+    {
+        pixels.physical_size_x = attr.as_double();
+        pixels.has_physical_size_x = true;
+    }
+    if (auto attr = pixels_node.attribute("PhysicalSizeY"))
+    {
+        pixels.physical_size_y = attr.as_double();
+        pixels.has_physical_size_y = true;
+    }
+    pixels.physical_size_x_unit = pixels_node.attribute("PhysicalSizeXUnit").as_string("");
+    pixels.physical_size_y_unit = pixels_node.attribute("PhysicalSizeYUnit").as_string("");
+
+    uint32_t channel_idx = 0;
+    for (auto child : pixels_node.children())
+    {
+        const auto child_name = local_name(child.name());
+        if (child_name == "Channel")
+        {
+            Channel channel;
+            channel.index = channel_idx++;
+            channel.id = child.attribute("ID").as_string("");
+            channel.name = child.attribute("Name").as_string("");
+            channel.samples_per_pixel = std::max<uint32_t>(1, child.attribute("SamplesPerPixel").as_uint(1));
+            pixels.channels.emplace_back(std::move(channel));
+        }
+        else if (child_name == "TiffData")
+        {
+            TiffDataEntry entry;
+            entry.ifd = child.attribute("IFD").as_uint(0);
+            entry.first_c = child.attribute("FirstC").as_uint(0);
+            entry.first_z = child.attribute("FirstZ").as_uint(0);
+            entry.first_t = child.attribute("FirstT").as_uint(0);
+            entry.plane_count = std::max<uint32_t>(1, child.attribute("PlaneCount").as_uint(1));
+
+            pugi::xml_node uuid_node = find_child_by_local_name(child, "UUID");
+            if (uuid_node)
+            {
+                entry.uuid = uuid_node.child_value();
+                entry.file_name = uuid_node.attribute("FileName").as_string("");
+            }
+
+            pixels.tiff_data.emplace_back(std::move(entry));
+        }
+    }
+
+    if (pixels.size_x == 0 || pixels.size_y == 0)
+    {
+        if (out_error)
+        {
+            *out_error = "OME Pixels is missing SizeX/SizeY";
+        }
+        return false;
+    }
+
+    model.pixels = std::move(pixels);
+    model.valid = true;
+    *out_model = std::move(model);
+    return true;
+}
+
+} // namespace cuslide::tiff::ome
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h
new file mode 100644
index 000000000..00de44bde
--- /dev/null
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h
@@ -0,0 +1,63 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <cstdint>
+#include <string>
+#include <vector>
+
+namespace cuslide::tiff::ome
+{
+
+struct Channel
+{
+    uint32_t index = 0;
+    std::string id;
+    std::string name;
+    uint32_t samples_per_pixel = 1;
+};
+
+struct TiffDataEntry
+{
+    uint32_t ifd = 0;
+    uint32_t first_c = 0;
+    uint32_t first_z = 0;
+    uint32_t first_t = 0;
+    uint32_t plane_count = 1;
+    std::string uuid;
+    std::string file_name;
+};
+
+struct Pixels
+{
+    uint32_t size_x = 0;
+    uint32_t size_y = 0;
+    uint32_t size_c = 1;
+    uint32_t size_z = 1;
+    uint32_t size_t = 1;
+    std::string type;
+    std::string dimension_order = "XYCZT";
+    double physical_size_x = 0.0;
+    double physical_size_y = 0.0;
+    std::string physical_size_x_unit;
+    std::string physical_size_y_unit;
+    bool has_physical_size_x = false;
+    bool has_physical_size_y = false;
+    std::vector<Channel> channels;
+    std::vector<TiffDataEntry> tiff_data;
+};
+
+struct Model
+{
+    bool valid = false;
+    Pixels pixels;
+    std::string image_id;
+    std::string image_name;
+};
+
+bool parse(const std::string& xml_text, Model* out_model, std::string* out_error = nullptr);
+
+} // namespace cuslide::tiff::ome
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp
index 01bf11c20..6797c6b96 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp
@@ -8,13 +8,15 @@
 #include <fcntl.h>
 
 #include <algorithm>
+#include <filesystem>
+#include <functional>
 #include <string>
 #include <string_view>
-#include <type_traits>
 
 #include <fmt/format.h>
 #include <nlohmann/json.hpp>
 #include <pugixml.hpp>
+#include "ome_xml.h"
 #include "tiff_constants.h"
 
 #include <cucim/codec/base64.h>
@@ -26,20 +28,6 @@
 
 static constexpr int DEFAULT_IFD_SIZE = 32;
 
-#ifdef CUCIM_HAS_NVIMGCODEC
-constexpr int kMetadataKindMedAperio = NVIMGCODEC_METADATA_KIND_MED_APERIO;
-constexpr int kMetadataKindMedPhilips = NVIMGCODEC_METADATA_KIND_MED_PHILIPS;
-constexpr int kMetadataKindMedVentana = NVIMGCODEC_METADATA_KIND_MED_VENTANA;
-constexpr int kMetadataKindMedLeica = NVIMGCODEC_METADATA_KIND_MED_LEICA;
-constexpr int kMetadataKindMedTrestle = NVIMGCODEC_METADATA_KIND_MED_TRESTLE;
-#else
-constexpr int kMetadataKindMedAperio = 5;
-constexpr int kMetadataKindMedPhilips = 6;
-constexpr int kMetadataKindMedVentana = 7;
-constexpr int kMetadataKindMedLeica = 8;
-constexpr int kMetadataKindMedTrestle = 9;
-#endif
-
 using json = nlohmann::json;
 
 namespace cuslide::tiff
@@ -82,44 +70,30 @@ static void parse_string_array(const char* values, json& arr, PhilipsMetadataTyp
         {
             if (text[next_pos - 1] != '\\')
             {
-                try
-                {
-                    std::string val(text.substr(pos + 1, next_pos - pos - 1));
-                    switch (type)
-                    {
-                    case PhilipsMetadataType::IString:
-                        arr.emplace_back(std::move(val));
-                        break;
-                    case PhilipsMetadataType::IDouble:
-                        arr.emplace_back(std::stod(val));
-                        break;
-                    case PhilipsMetadataType::IUInt16:
-                    case PhilipsMetadataType::IUInt32:
-                    case PhilipsMetadataType::IUInt64:
-                        arr.emplace_back(std::stoul(val));
-                        break;
-                    }
-                }
-                catch (const std::exception&)
+                switch (type)
                 {
+                case PhilipsMetadataType::IString:
+                    arr.emplace_back(std::string(text.substr(pos + 1, next_pos - pos - 1)));
+                    break;
+                case PhilipsMetadataType::IDouble:
+                    arr.emplace_back(std::stod(std::string(text.substr(pos + 1, next_pos - pos - 1))));
+                    break;
+                case PhilipsMetadataType::IUInt16:
+                case PhilipsMetadataType::IUInt32:
+                case PhilipsMetadataType::IUInt64:
+                    arr.emplace_back(std::stoul(std::string(text.substr(pos + 1, next_pos - pos - 1))));
+                    break;
                 }
                 pos = next_pos + 1;
             }
         }
     }
 }
-static constexpr int kMaxXmlParseDepth = 32;
-
 static void parse_philips_tiff_metadata(const pugi::xml_node& node,
                                         json& metadata,
                                         const char* name,
-                                        PhilipsMetadataStage stage,
-                                        int depth = 0)
+                                        PhilipsMetadataStage stage)
 {
-    if (depth > kMaxXmlParseDepth)
-    {
-        return;
-    }
     switch (stage)
     {
     case PhilipsMetadataStage::ROOT:
@@ -130,8 +104,7 @@ static void parse_philips_tiff_metadata(const pugi::xml_node& node,
             const pugi::xml_attribute& attr_attribute = attr.attribute("Name");
             if (attr_attribute)
             {
-                parse_philips_tiff_metadata(
-                    attr, metadata, attr_attribute.value(), PhilipsMetadataStage::ELEMENT, depth + 1);
+                parse_philips_tiff_metadata(attr, metadata, attr_attribute.value(), PhilipsMetadataStage::ELEMENT);
             }
         }
         break;
@@ -196,8 +169,7 @@ static void parse_philips_tiff_metadata(const pugi::xml_node& node,
                         {
                             auto& item_iter = item_array_iter.first->emplace_back(json{});
                             parse_philips_tiff_metadata(
-                                data_node, item_iter, nullptr, PhilipsMetadataStage::PIXEL_DATA_PRESENTATION,
-                                depth + 1);
+                                data_node, item_iter, nullptr, PhilipsMetadataStage::PIXEL_DATA_PRESENTATION);
                         }
                     }
                     break;
@@ -249,20 +221,13 @@ static std::string strip_string(const std::string& str)
     }
 }
 
-static constexpr size_t kMaxImageDescriptionBytes = 1 * 1024 * 1024; // 1 MB
-
 static void parse_aperio_svs_metadata(std::shared_ptr<IFD>& first_ifd, json& metadata)
 {
     (void)metadata;
     std::string& desc = first_ifd->image_description();
 
-    if (desc.size() > kMaxImageDescriptionBytes)
-    {
-        fmt::print(stderr, "[Warning] Aperio SVS ImageDescription exceeds 1 MB limit ({}); skipping metadata parse\n",
-                   desc.size());
-        return;
-    }
-
+    // Assumes that metadata's image description starts with 'Aperio '.
+    // It is handled by 'resolve_vendor_format()'
     std::vector<std::string> items = split_string(desc, "|");
     if (items.size() < 1)
     {
@@ -379,6 +344,8 @@ TIFF::TIFF(const cucim::filesystem::Path& file_path) : file_path_(file_path)
 
     // Initialize metadata container
     metadata_ = new json{};
+    parser_file_hashes_.clear();
+    parser_file_hashes_.push_back(file_handle_shared_ ? file_handle_shared_->hash_value : 0);
 }
 
 TIFF::TIFF(const cucim::filesystem::Path& file_path, uint64_t read_config) : TIFF(file_path)
@@ -453,6 +420,18 @@ void TIFF::construct_ifds()
 
     ifd_offsets_.clear();
     ifds_.clear();
+    parser_local_to_global_ifd_.clear();
+    global_to_parser_local_ifd_.clear();
+    companion_parsers_.clear();
+    companion_parser_paths_.clear();
+    if (parser_file_hashes_.empty())
+    {
+        parser_file_hashes_.push_back(file_handle_shared_ ? file_handle_shared_->hash_value : 0);
+    }
+    else
+    {
+        parser_file_hashes_.resize(1);
+    }
 
     uint32_t ifd_count = nvimgcodec_parser_->get_ifd_count();
     #ifdef DEBUG
@@ -472,8 +451,16 @@ void TIFF::construct_ifds()
             ifd_offsets_.push_back(ifd_index);
 
             // Create IFD from nvImageCodec metadata using NEW constructor
-            auto ifd = std::make_shared<cuslide::tiff::IFD>(this, ifd_index, ifd_info);
+            auto ifd = std::make_shared<cuslide::tiff::IFD>(
+                this,
+                ifd_index,
+                ifd_info,
+                0,
+                ifd_index,
+                parser_file_hashes_[0]);
             ifds_.emplace_back(std::move(ifd));
+            parser_local_to_global_ifd_[{ 0, ifd_index }] = ifds_.size() - 1;
+            global_to_parser_local_ifd_.emplace_back(0, ifd_index);
 
             #ifdef DEBUG
             fmt::print("  ✅ IFD[{}]: {}x{}, {} channels, codec: {}\n",
@@ -527,6 +514,104 @@ void TIFF::construct_ifds()
               level_to_ifd_idx_.size(), associated_images_.size());
     #endif // DEBUG
 }
+
+uint64_t TIFF::_hash_path(const std::string& path)
+{
+    return static_cast<uint64_t>(std::hash<std::string>{}(path));
+}
+
+size_t TIFF::_find_or_add_parser_context(const std::string& abs_path)
+{
+    if (abs_path == std::string(file_path_.c_str()))
+    {
+        return 0;
+    }
+    for (size_t i = 0; i < companion_parser_paths_.size(); ++i)
+    {
+        if (companion_parser_paths_[i] == abs_path)
+        {
+            return i + 1;
+        }
+    }
+
+    auto parser = std::make_unique<cuslide2::nvimgcodec::TiffFileParser>(abs_path);
+    if (!parser || !parser->is_valid())
+    {
+        throw std::runtime_error(fmt::format("Failed to open OME companion TIFF '{}'", abs_path));
+    }
+    companion_parser_paths_.emplace_back(abs_path);
+    companion_parsers_.emplace_back(std::move(parser));
+    parser_file_hashes_.emplace_back(_hash_path(abs_path));
+    return companion_parsers_.size();
+}
+
+const cuslide2::nvimgcodec::TiffFileParser& TIFF::parser_for_global_ifd(size_t global_ifd_idx) const
+{
+    if (global_ifd_idx >= global_to_parser_local_ifd_.size())
+    {
+        throw std::out_of_range(fmt::format("Global IFD {} is out of range", global_ifd_idx));
+    }
+    const size_t parser_idx = global_to_parser_local_ifd_[global_ifd_idx].first;
+    if (parser_idx == 0)
+    {
+        return *nvimgcodec_parser_;
+    }
+    const size_t companion_idx = parser_idx - 1;
+    if (companion_idx >= companion_parsers_.size() || !companion_parsers_[companion_idx])
+    {
+        throw std::out_of_range(fmt::format("Companion parser {} is unavailable", parser_idx));
+    }
+    return *companion_parsers_[companion_idx];
+}
+
+uint32_t TIFF::local_ifd_for_global_ifd(size_t global_ifd_idx) const
+{
+    if (global_ifd_idx >= global_to_parser_local_ifd_.size())
+    {
+        throw std::out_of_range(fmt::format("Global IFD {} is out of range", global_ifd_idx));
+    }
+    return global_to_parser_local_ifd_[global_ifd_idx].second;
+}
+
+uint64_t TIFF::file_hash_for_global_ifd(size_t global_ifd_idx) const
+{
+    if (global_ifd_idx >= global_to_parser_local_ifd_.size())
+    {
+        throw std::out_of_range(fmt::format("Global IFD {} is out of range", global_ifd_idx));
+    }
+    const size_t parser_idx = global_to_parser_local_ifd_[global_ifd_idx].first;
+    if (parser_idx >= parser_file_hashes_.size())
+    {
+        return file_handle_shared_ ? file_handle_shared_->hash_value : 0;
+    }
+    return parser_file_hashes_[parser_idx];
+}
+
+size_t TIFF::_ensure_global_ifd(size_t parser_idx, uint32_t local_ifd_idx)
+{
+    auto map_it = parser_local_to_global_ifd_.find({ parser_idx, local_ifd_idx });
+    if (map_it != parser_local_to_global_ifd_.end())
+    {
+        return map_it->second;
+    }
+
+    const auto& parser = (parser_idx == 0) ? *nvimgcodec_parser_ : *companion_parsers_[parser_idx - 1];
+    const auto& ifd_info = parser.get_ifd(local_ifd_idx);
+    const size_t global_ifd_idx = ifds_.size();
+    auto ifd = std::make_shared<cuslide::tiff::IFD>(
+        this,
+        static_cast<uint16_t>(global_ifd_idx),
+        ifd_info,
+        parser_idx,
+        local_ifd_idx,
+        parser_file_hashes_[parser_idx]);
+    ifd_offsets_.push_back(global_ifd_idx);
+    ifds_.emplace_back(std::move(ifd));
+    parser_local_to_global_ifd_[{ parser_idx, local_ifd_idx }] = global_ifd_idx;
+    global_to_parser_local_ifd_.emplace_back(parser_idx, local_ifd_idx);
+    return global_ifd_idx;
+}
+
 void TIFF::resolve_vendor_format()
 {
     PROF_SCOPED_RANGE(PROF_EVENT(tiff_resolve_vendor_format));
@@ -548,19 +633,20 @@ void TIFF::resolve_vendor_format()
     {
         bool is_aperio = false;
 
+        // Method 1: Check ImageDescription starts with "Aperio "
         auto& image_desc = first_ifd->image_description();
         std::string_view prefix("Aperio ");
-        if (image_desc.size() >= prefix.size() &&
-            std::mismatch(prefix.begin(), prefix.end(), image_desc.begin()).first == prefix.end())
+        auto res = std::mismatch(prefix.begin(), prefix.end(), image_desc.begin());
+        if (res.first == prefix.end())
         {
             is_aperio = true;
         }
 
-        // Method 2: Check metadata_blobs for Aperio (MED_APERIO)
+        // Method 2: Check metadata_blobs for Aperio
         if (!is_aperio && nvimgcodec_parser_)
         {
             const auto& metadata_blobs = nvimgcodec_parser_->get_metadata_blobs(0);
-            if (metadata_blobs.find(kMetadataKindMedAperio) != metadata_blobs.end())
+            if (metadata_blobs.find(5) != metadata_blobs.end())  // MED_APERIO = 5
             {
                 is_aperio = true;
                 #ifdef DEBUG
@@ -575,19 +661,21 @@ void TIFF::resolve_vendor_format()
         }
     }
 
-    // Detect Philips TIFF
-    // Check for Philips TIFF using SOFTWARE tag, ImageDescription XML, or nvImageCodec metadata kind
+    // Detect Philips TIFF. Not every file carries a SOFTWARE tag, so fall back to
+    // the Philips XML in ImageDescription and to the vendor metadata kind.
     {
         bool is_philips = false;
 
+        // Method 1: Check SOFTWARE tag
         std::string_view prefix("Philips");
-        if (software.size() >= prefix.size() &&
-            std::mismatch(prefix.begin(), prefix.end(), software.begin()).first == prefix.end())
+        auto res = std::mismatch(prefix.begin(), prefix.end(), software.begin());
+        if (res.first == prefix.end())
         {
             is_philips = true;
         }
 
-        // Method 2: Check for Philips XML structure in ImageDescription (fallback)
+        // Method 2: Check for Philips XML structure in ImageDescription
+        // (used when the file has no SOFTWARE tag)
         if (!is_philips)
         {
             auto& image_desc = first_ifd->image_description();
@@ -598,11 +686,11 @@ void TIFF::resolve_vendor_format()
             }
         }
 
-        // Method 3: Check metadata_blobs for Philips (MED_PHILIPS)
+        // Method 3: Check metadata_blobs for Philips
         if (!is_philips && nvimgcodec_parser_)
         {
             const auto& metadata_blobs = nvimgcodec_parser_->get_metadata_blobs(0);
-            if (metadata_blobs.find(kMetadataKindMedPhilips) != metadata_blobs.end())
+            if (metadata_blobs.find(6) != metadata_blobs.end())  // MED_PHILIPS = 6
             {
                 is_philips = true;
                 #ifdef DEBUG
@@ -642,28 +730,28 @@ void TIFF::resolve_vendor_format()
             }
         }
 
-        const std::remove_reference_t<decltype(nvimgcodec_parser_->get_metadata_blobs(0))>* blobs =
-            nvimgcodec_parser_ ? &nvimgcodec_parser_->get_metadata_blobs(0) : nullptr;
-
         // Hamamatsu NDPI: extension .ndpi or MODEL contains "Hamamatsu"
         if (file_ext == ".ndpi" || model.find("Hamamatsu") != std::string::npos)
         {
             tiff_type_ = TiffType::Hamamatsu;
         }
-        // Ventana BIF: extension .bif or nvImageCodec MED_VENTANA blob
+        // Ventana BIF: extension .bif or nvImageCodec MED_VENTANA blob (kind=7)
         else if (file_ext == ".bif" ||
-                 (blobs && blobs->find(kMetadataKindMedVentana) != blobs->end()))
+                 (nvimgcodec_parser_ &&
+                  nvimgcodec_parser_->get_metadata_blobs(0).find(7) != nvimgcodec_parser_->get_metadata_blobs(0).end()))
         {
             tiff_type_ = TiffType::Ventana;
         }
-        // Leica SCN: extension .scn or nvImageCodec MED_LEICA blob
+        // Leica SCN: extension .scn or nvImageCodec MED_LEICA blob (kind=8, when not NDPI)
         else if (file_ext == ".scn" ||
-                 (blobs && blobs->find(kMetadataKindMedLeica) != blobs->end()))
+                 (nvimgcodec_parser_ &&
+                  nvimgcodec_parser_->get_metadata_blobs(0).find(8) != nvimgcodec_parser_->get_metadata_blobs(0).end()))
         {
             tiff_type_ = TiffType::Leica;
         }
-        // Trestle TIFF: nvImageCodec MED_TRESTLE blob
-        else if (blobs && blobs->find(kMetadataKindMedTrestle) != blobs->end())
+        // Trestle TIFF: nvImageCodec MED_TRESTLE blob (kind=9)
+        else if (nvimgcodec_parser_ &&
+                 nvimgcodec_parser_->get_metadata_blobs(0).find(9) != nvimgcodec_parser_->get_metadata_blobs(0).end())
         {
             tiff_type_ = TiffType::Trestle;
         }
@@ -672,13 +760,14 @@ void TIFF::resolve_vendor_format()
                  image_desc.find("ome.xsd") != std::string::npos)
         {
             tiff_type_ = TiffType::OmeTiff;
+            _populate_ome_tiff_metadata(ifd_count, json_metadata, first_ifd);
         }
         // Vectra QPTIFF: extension .qptiff or ImageDescription contains PerkinElmer
         else if (file_ext == ".qptiff" ||
                  image_desc.find("PerkinElmer") != std::string::npos ||
                  image_desc.find("Vectra") != std::string::npos)
         {
-            tiff_type_ = TiffType::QpTiff;
+            tiff_type_ = TiffType::Qptiff;
         }
 
         #ifdef DEBUG
@@ -715,6 +804,238 @@ void TIFF::resolve_vendor_format()
     }
 }
 
+void TIFF::_populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr<IFD>& first_ifd)
+{
+    (void)ifd_count;
+    json* json_metadata = reinterpret_cast<json*>(metadata);
+
+    ome::Model model;
+    std::string parse_error;
+    if (!ome::parse(first_ifd->image_description(), &model, &parse_error) || !model.valid)
+    {
+        // The file advertised OME XML but it could not be used, so callers silently lose
+        // channel/plane information. Report why rather than discarding the reason.
+        fmt::print(stderr, "[cuslide2] OME XML present but not usable ({} bytes): {}\n",
+                   first_ifd->image_description().size(),
+                   parse_error.empty() ? "incomplete OME model" : parse_error);
+        return;
+    }
+
+    ome_size_c_ = std::max<int64_t>(1, model.pixels.size_c);
+    ome_size_z_ = std::max<int64_t>(1, model.pixels.size_z);
+    ome_size_t_ = std::max<int64_t>(1, model.pixels.size_t);
+
+    ome_channel_names_.clear();
+    ome_channel_name_to_index_.clear();
+    if (!model.pixels.channels.empty())
+    {
+        ome_channel_names_.reserve(model.pixels.channels.size());
+        for (const auto& channel : model.pixels.channels)
+        {
+            std::string channel_name = channel.name.empty() ? fmt::format("C{}", channel.index) : channel.name;
+            if (!ome_channel_name_to_index_.count(channel_name))
+            {
+                ome_channel_name_to_index_.emplace(channel_name, static_cast<int64_t>(ome_channel_names_.size()));
+            }
+            ome_channel_names_.emplace_back(std::move(channel_name));
+        }
+    }
+    else
+    {
+        for (int64_t c = 0; c < ome_size_c_; ++c)
+        {
+            ome_channel_names_.emplace_back(fmt::format("C{}", c));
+            ome_channel_name_to_index_.emplace(ome_channel_names_.back(), c);
+        }
+    }
+
+    std::vector<char> logical_axes;
+    logical_axes.reserve(3);
+    for (char axis : model.pixels.dimension_order)
+    {
+        if (axis == 'C' || axis == 'Z' || axis == 'T')
+        {
+            logical_axes.emplace_back(axis);
+        }
+    }
+
+    auto increment_plane = [&](int64_t& c, int64_t& z, int64_t& t) {
+        for (auto it = logical_axes.rbegin(); it != logical_axes.rend(); ++it)
+        {
+            switch (*it)
+            {
+            case 'C':
+                c = (c + 1) % ome_size_c_;
+                if (c != 0)
+                    return;
+                break;
+            case 'Z':
+                z = (z + 1) % ome_size_z_;
+                if (z != 0)
+                    return;
+                break;
+            case 'T':
+                t = (t + 1) % ome_size_t_;
+                if (t != 0)
+                    return;
+                break;
+            default:
+                break;
+            }
+        }
+    };
+
+    std::vector<std::tuple<int64_t, int64_t, int64_t, size_t>> pending_planes;
+    pending_planes.reserve(model.pixels.tiff_data.size());
+
+    const std::filesystem::path base_path = std::filesystem::path(file_path_.c_str()).parent_path();
+    for (const auto& tiff_data : model.pixels.tiff_data)
+    {
+        size_t parser_idx = 0;
+        if (!tiff_data.file_name.empty())
+        {
+            std::filesystem::path candidate = base_path / tiff_data.file_name;
+            candidate = candidate.lexically_normal();
+            parser_idx = _find_or_add_parser_context(candidate.string());
+        }
+
+        const auto& parser = (parser_idx == 0) ? *nvimgcodec_parser_ : *companion_parsers_[parser_idx - 1];
+        int64_t c = tiff_data.first_c;
+        int64_t z = tiff_data.first_z;
+        int64_t t = tiff_data.first_t;
+
+        for (uint32_t p = 0; p < tiff_data.plane_count; ++p)
+        {
+            uint32_t local_ifd_idx = tiff_data.ifd + p;
+            if (local_ifd_idx >= parser.get_ifd_count())
+            {
+                break;
+            }
+            size_t global_ifd_idx = _ensure_global_ifd(parser_idx, local_ifd_idx);
+            pending_planes.emplace_back(c, z, t, global_ifd_idx);
+            increment_plane(c, z, t);
+        }
+    }
+
+    ome_plane_to_ifd_.clear();
+
+    // Distinct IFD dimensions, largest first: index into this list is the resolution level.
+    std::vector<std::pair<uint32_t, uint32_t>> level_dims;
+    if (!pending_planes.empty())
+    {
+        level_dims.reserve(ifds_.size());
+        for (const auto& ifd : ifds_)
+        {
+            std::pair<uint32_t, uint32_t> dim = { ifd->width(), ifd->height() };
+            if (std::find(level_dims.begin(), level_dims.end(), dim) == level_dims.end())
+            {
+                level_dims.emplace_back(dim);
+            }
+        }
+        std::sort(level_dims.begin(), level_dims.end(), [](const auto& a, const auto& b) {
+            if (a.first != b.first)
+            {
+                return a.first > b.first;
+            }
+            return a.second > b.second;
+        });
+
+        auto level_for_ifd = [&level_dims, this](size_t ifd_idx) -> uint16_t {
+            const auto& ifd = ifds_[ifd_idx];
+            std::pair<uint32_t, uint32_t> dim = { ifd->width(), ifd->height() };
+            auto it = std::find(level_dims.begin(), level_dims.end(), dim);
+            return (it == level_dims.end()) ? 0 : static_cast<uint16_t>(std::distance(level_dims.begin(), it));
+        };
+
+        for (const auto& [c, z, t, ifd_idx] : pending_planes)
+        {
+            ome_plane_to_ifd_[std::make_tuple(c, z, t, level_for_ifd(ifd_idx))] = ifd_idx;
+        }
+    }
+    has_ome_plane_index_ = !ome_plane_to_ifd_.empty();
+
+    if (has_ome_plane_index_)
+    {
+        std::map<uint16_t, size_t> level_representative_ifd;
+        for (const auto& [key, ifd_idx] : ome_plane_to_ifd_)
+        {
+            if (std::get<0>(key) == 0 && std::get<1>(key) == 0 && std::get<2>(key) == 0)
+            {
+                level_representative_ifd[std::get<3>(key)] = ifd_idx;
+            }
+        }
+        if (level_representative_ifd.empty())
+        {
+            for (const auto& [key, ifd_idx] : ome_plane_to_ifd_)
+            {
+                uint16_t level = std::get<3>(key);
+                if (!level_representative_ifd.count(level))
+                {
+                    level_representative_ifd[level] = ifd_idx;
+                }
+            }
+        }
+        // OME TiffData entries enumerate the full-resolution planes only; pyramid
+        // sub-resolutions are additional IFDs the XML never references. Mapping planes
+        // alone therefore yields level 0 and drops every sub-resolution, making them
+        // unreachable through read_region(). Backfill the levels the plane index missed.
+        for (size_t level = 0; level < level_dims.size(); ++level)
+        {
+            const auto level_key = static_cast<uint16_t>(level);
+            if (level_representative_ifd.count(level_key) != 0)
+            {
+                continue;
+            }
+            for (size_t ifd_idx = 0; ifd_idx < ifds_.size(); ++ifd_idx)
+            {
+                if (ifds_[ifd_idx]->width() == level_dims[level].first &&
+                    ifds_[ifd_idx]->height() == level_dims[level].second)
+                {
+                    level_representative_ifd[level_key] = ifd_idx;
+                    break;
+                }
+            }
+        }
+
+        if (!level_representative_ifd.empty())
+        {
+            // std::map iterates in ascending level order, so levels stay largest-first.
+            level_to_ifd_idx_.clear();
+            for (const auto& [level, ifd_idx] : level_representative_ifd)
+            {
+                (void)level;
+                level_to_ifd_idx_.push_back(ifd_idx);
+            }
+        }
+    }
+
+    if (json_metadata)
+    {
+        json ome_metadata;
+        ome_metadata.emplace("size_c", ome_size_c_);
+        ome_metadata.emplace("size_z", ome_size_z_);
+        ome_metadata.emplace("size_t", ome_size_t_);
+        ome_metadata.emplace("size_x", model.pixels.size_x);
+        ome_metadata.emplace("size_y", model.pixels.size_y);
+        ome_metadata.emplace("type", model.pixels.type);
+        ome_metadata.emplace("dimension_order", model.pixels.dimension_order);
+        ome_metadata.emplace("channel_names", ome_channel_names_);
+        if (model.pixels.has_physical_size_x)
+        {
+            ome_metadata.emplace("physical_size_x", model.pixels.physical_size_x);
+            ome_metadata.emplace("physical_size_x_unit", model.pixels.physical_size_x_unit);
+        }
+        if (model.pixels.has_physical_size_y)
+        {
+            ome_metadata.emplace("physical_size_y", model.pixels.physical_size_y);
+            ome_metadata.emplace("physical_size_y_unit", model.pixels.physical_size_y_unit);
+        }
+        ome_metadata.emplace("plane_count", pending_planes.size());
+        ome_metadata.emplace("companion_file_count", companion_parsers_.size());
+        (*json_metadata)["ome"] = std::move(ome_metadata);
+    }
+}
+
 void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr<IFD>& first_ifd)
 {
     json* json_metadata = reinterpret_cast<json*>(metadata);
@@ -785,12 +1106,6 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s
             pixel_spacings.emplace_back(std::pair{ spacing_x, spacing_y });
         }
 
-        if (pixel_spacings.empty())
-        {
-            fmt::print(stderr, "[Warning] Philips TIFF: no DICOM_PIXEL_SPACING entries found\n");
-            return;
-        }
-
         double spacing_x_l0 = pixel_spacings[0].first;
         double spacing_y_l0 = pixel_spacings[0].second;
 
@@ -806,12 +1121,8 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s
             // NOTE: In Philips TIFF, pyramid levels can be strip-based (tile_width==0)
             // So we can't use tile_width to identify associated images
             auto& image_desc = ifd->image_description();
-            bool is_macro = (image_desc.size() >= macro_prefix.size() &&
-                             std::mismatch(macro_prefix.begin(), macro_prefix.end(), image_desc.begin()).first ==
-                                 macro_prefix.end());
-            bool is_label = (image_desc.size() >= label_prefix.size() &&
-                             std::mismatch(label_prefix.begin(), label_prefix.end(), image_desc.begin()).first ==
-                                 label_prefix.end());
+            bool is_macro = (std::mismatch(macro_prefix.begin(), macro_prefix.end(), image_desc.begin()).first == macro_prefix.end());
+            bool is_label = (std::mismatch(label_prefix.begin(), label_prefix.end(), image_desc.begin()).first == label_prefix.end());
 
             if (is_macro || is_label)
             {
@@ -842,6 +1153,9 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s
                 double downsample = std::round((pixel_spacings[spacing_index].first / spacing_x_l0 +
                                                 pixel_spacings[spacing_index].second / spacing_y_l0) /
                                                2);
+                // Fix width and height of IFD. Guard against bogus pixel spacing metadata that
+                // rounds to zero (or worse, negative): the division is floating-point, so the
+                // result would be inf/negative and the narrowing to uint32_t undefined.
                 if (downsample > 0)
                 {
                     ifd->width_ = width_l0 / downsample;
@@ -888,6 +1202,11 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s
                     size_t image_desc_len = first_ifd->image_description().size();
                     if (node_offset >= 0 && static_cast<size_t>(node_offset) < image_desc_len)
                     {
+                        // `image_desc_cstr[node_offset]` would point to the following text:
+                        //   Attribute Element="0x1004" Group="0x301D" Name="PIM_DP_IMAGE_DATA" PMSVR="IString">
+                        //     (base64-encoded JPEG image)
+                        //   </Attribute>
+                        //
                         const char* data_ptr = image_desc_cstr + node_offset;
                         const char* desc_end = image_desc_cstr + image_desc_len;
 
@@ -901,6 +1220,7 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s
                         {
                             ++data_ptr; // start of base64-encoded data
                             const char* data_end_ptr = data_ptr;
+                            // Seek until it finds '<' for '</Attribute>'
                             while (data_end_ptr < desc_end && *data_end_ptr != '<' && *data_end_ptr != '\0')
                             {
                                 ++data_end_ptr;
@@ -1039,7 +1359,8 @@ bool TIFF::read(const cucim::io::format::ImageMetadataDesc* metadata,
             "Invalid level ({}) in the request! (Should be < {})", request->level, level_to_ifd_idx_.size()));
     }
     auto main_ifd = ifds_[level_to_ifd_idx_[0]];
-    auto ifd = ifds_[level_to_ifd_idx_[request->level]];
+    size_t selected_ifd_idx = level_to_ifd_idx_[request->level];
+    auto ifd = ifds_[selected_ifd_idx];
     auto original_img_width = main_ifd->width();
     auto original_img_height = main_ifd->height();
 
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h
index cb8e8b5ca..2d34aedd2 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -17,6 +17,8 @@
 #include <map>
 #include <memory>
 #include <mutex>
+#include <string>
+#include <tuple>
 #include <vector>
 
 #include "ifd.h"
@@ -98,6 +100,26 @@ class EXPORT_VISIBLE TIFF : public std::enable_shared_from_this<TIFF>
     friend class IFD;
 
 private:
+    std::vector<std::unique_ptr<cuslide2::nvimgcodec::TiffFileParser>> companion_parsers_;
+    std::vector<std::string> companion_parser_paths_;
+    std::vector<uint64_t> parser_file_hashes_;
+    std::map<std::pair<size_t, uint32_t>, size_t> parser_local_to_global_ifd_;
+    std::vector<std::pair<size_t, uint32_t>> global_to_parser_local_ifd_;
+    bool has_ome_plane_index_ = false;
+    int64_t ome_size_c_ = 1;
+    int64_t ome_size_z_ = 1;
+    int64_t ome_size_t_ = 1;
+    std::vector<std::string> ome_channel_names_;
+    std::map<std::string, int64_t> ome_channel_name_to_index_;
+    std::map<std::tuple<int64_t, int64_t, int64_t, uint16_t>, size_t> ome_plane_to_ifd_;
+    void _populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr<IFD>& first_ifd);
+    size_t _ensure_global_ifd(size_t parser_idx, uint32_t local_ifd_idx);
+    size_t _find_or_add_parser_context(const std::string& abs_path);
+    static uint64_t _hash_path(const std::string& path);
+    const cuslide2::nvimgcodec::TiffFileParser& parser_for_global_ifd(size_t global_ifd_idx) const;
+    uint32_t local_ifd_for_global_ifd(size_t global_ifd_idx) const;
+    uint64_t file_hash_for_global_ifd(size_t global_ifd_idx) const;
+
     // UPDATED: These now use nvImageCodec TiffFileParser instead of libtiff
     void _populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr<IFD>& first_ifd);
     void _populate_aperio_svs_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr<IFD>& first_ifd);
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff_constants.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff_constants.h
index 45bed33de..c8c9971de 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff_constants.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff_constants.h
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -38,6 +38,8 @@ constexpr uint16_t COMPRESSION_APERIO_JP2K_YCBCR = 33003;
 constexpr uint16_t COMPRESSION_APERIO_JP2K_RGB = 33005;
 
 // TIFF Photometric Interpretation
+constexpr uint16_t PHOTOMETRIC_MINISWHITE = 0;
+constexpr uint16_t PHOTOMETRIC_MINISBLACK = 1;
 constexpr uint16_t PHOTOMETRIC_RGB = 2;
 constexpr uint16_t PHOTOMETRIC_YCBCR = 6;
 
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h
index 42d5ea24b..08f8f6b00 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h
@@ -25,7 +25,7 @@ enum class TiffType : uint32_t
     Ventana = 5,
     Trestle = 6,
     OmeTiff = 7,
-    QpTiff = 8,
+    Qptiff = 8,
 };
 
 enum class AssociatedImageBufferType : uint8_t
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec.h b/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec.h
index 2cc0003ca..d005a8c86 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec.h
@@ -1,6 +1,7 @@
 /*
  * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
+ *
  */
 
 /**
@@ -385,6 +386,16 @@ extern "C"
         NVIMGCODEC_SAMPLEFORMAT_UNKNOWN = 0,
         NVIMGCODEC_SAMPLEFORMAT_P_UNCHANGED = 1, /**< Unchanged planar. */
         NVIMGCODEC_SAMPLEFORMAT_I_UNCHANGED = 2, /**< Unchanged interleaved. */
+        /** Alias of P_UNCHANGED for callers that want to pin only the channel
+         * interleaving (planar) and let color_spec choose the components.
+         * Preferred over P_UNCHANGED whenever color_spec is concrete; the two
+         * names produce identical behaviour. */
+        NVIMGCODEC_SAMPLEFORMAT_P_AUTO_COMPONENTS = NVIMGCODEC_SAMPLEFORMAT_P_UNCHANGED,
+        /** Alias of I_UNCHANGED for callers that want to pin only the channel
+         * interleaving (interleaved) and let color_spec choose the components.
+         * Preferred over I_UNCHANGED whenever color_spec is concrete; the two
+         * names produce identical behaviour. */
+        NVIMGCODEC_SAMPLEFORMAT_I_AUTO_COMPONENTS = NVIMGCODEC_SAMPLEFORMAT_I_UNCHANGED,
         NVIMGCODEC_SAMPLEFORMAT_P_Y = 3,         /**< Planar Y component only. For 3-dimensional shape is defined as (1, H, W) and this is only difference with I_Y.*/
         NVIMGCODEC_SAMPLEFORMAT_I_Y = 4,         /**< Interleaved Y component only. For 3-dimensional shape is defined as (H, W, 1) and this is only difference with P_Y.*/
         NVIMGCODEC_SAMPLEFORMAT_P_YA = 5,        /**< Planar Y component with alpha. */
@@ -436,9 +447,11 @@ extern "C"
         size_t struct_size;                    /**< The size of the structure, in bytes. */
         void* struct_next;                     /**< Is NULL or a pointer to an extension structure type. */
 
-        int rotated; /**< Rotation angle in degrees (clockwise). Only multiples of 90 are allowed. */
-        int flip_x;  /**< Flip horizontal 0 or 1*/
-        int flip_y;  /**< Flip vertical 0 or 1*/
+        int rotated; /**< Counter-clockwise rotation in degrees applied to the raw codestream
+                          pixels to produce the displayed image. Only multiples of 90 are
+                          allowed (0, 90, 180, 270). */
+        int flip_x;  /**< Mirror across the vertical axis, applied after the rotation. 0 or 1. */
+        int flip_y;  /**< Mirror across the horizontal axis, applied after the rotation. 0 or 1. */
     } nvimgcodecOrientation_t;
 
     /**
@@ -486,6 +499,18 @@ extern "C"
 
     /**
      * @brief Defines region of an image.
+     *
+     * The coordinate system depends on `nvimgcodecDecodeParams_t::apply_exif_orientation`:
+     * - When `apply_exif_orientation == 1`, the region is interpreted in **display**
+     *   (post-orientation) coordinates. For a codestream whose EXIF orientation swaps width
+     *   and height (rotated 90 or 270 degrees), this means the region's axes match the rotated
+     *   image, not the raw codestream layout.
+     * - When `apply_exif_orientation == 0`, the region is interpreted in **codestream**
+     *   coordinates, matching `nvimgcodecImagePlaneInfo_t::width`/`height` of the parsed code
+     *   stream and `nvimgcodec::CodeStream.height`/`width` in the Python bindings.
+     *
+     * `start[0]`/`end[0]` is the y (height) axis; `start[1]`/`end[1]` is the x (width) axis.
+     * Half-open ranges: a pixel is included when `start <= idx < end`.
      */
     typedef struct
     {
@@ -556,8 +581,10 @@ extern "C"
 
        size_t image_idx;                     /**< Image index starts from 0. */
        nvimgcodecRegion_t region;            /**< Region of interest. */
-       size_t bitstream_offset;             /**< Byte offset to start parsing from relative to file start. 0 = start at default position. */
-       uint32_t limit_images;               /**< Maximum number of images to parse. 0 = no limit (parse all). */
+       size_t bitstream_offset;              /**< TIFF IFD byte offset selecting one image view; 0 selects the default view.
+                                                 Note that this field cannot be used in conjunction with `image_idx`. Currently, this field
+                                                 is only used by the TIFF extension, and offsets should normally come from `ifd_offset`,
+                                                 `next_ifd_offset`, or `subifd_offsets` of `nvimgcodecCodeStreamInfoTiffExt_t`. */
    } nvimgcodecCodeStreamView_t;
 
     /**
@@ -573,7 +600,7 @@ extern "C"
         char codec_name[NVIMGCODEC_MAX_CODEC_NAME_SIZE]; /**< Information about codec used. Only valid when used with code stream. */
 
         size_t num_images;                  /**< Number of images in CodeStream. */
-        size_t size;              /**< Size of bitstream in bytes. */
+        size_t size;                        /**< Size of underlying bitstream in bytes. */
     } nvimgcodecCodeStreamInfo_t;
 
     /** Maximum number of SubIFD offsets that can be stored in CodeStreamInfoTiffExt */
@@ -583,7 +610,7 @@ extern "C"
      * @brief TIFF-specific extension for CodeStreamInfo.
      *
      * Chain this struct via struct_next of nvimgcodecCodeStreamInfo_t to receive
-     * pagination and SubIFD information for TIFF files.
+     * IFD traversal and SubIFD information for TIFF files.
      */
     typedef struct
     {
@@ -591,10 +618,10 @@ extern "C"
         size_t struct_size;                    /**< The size of the structure, in bytes. */
         void* struct_next;                     /**< Is NULL or a pointer to an extension structure type. */
 
-        size_t next_bitstream_offset;  /**< Offset to next IFD for pagination. 0 = no more images. */
-        size_t target_ifd_offset;      /**< Byte offset of the target image's IFD. Used internally to avoid redundant IFD walks. */
+        size_t ifd_offset;             /**< Byte offset of the selected image's IFD. 0 = unknown or not applicable. */
+        size_t next_ifd_offset;        /**< Byte offset of the next sibling IFD. 0 = no next sibling. */
 
-        uint32_t subifd_count;         /**< Number of SubIFDs for the current image (Tag 330). */
+        uint32_t subifd_count;         /**< Number of stored SubIFD offsets for the current image (Tag 330), capped at `NVIMGCODEC_MAX_SUBIFD_OFFSETS`. */
         size_t subifd_offsets[NVIMGCODEC_MAX_SUBIFD_OFFSETS]; /**< SubIFD byte offsets. Only first subifd_count entries are valid. */
     } nvimgcodecCodeStreamInfoTiffExt_t;
 
@@ -698,7 +725,7 @@ extern "C"
         float load_hint;
 
         /**
-         * If true, the backend load will be adapted on every iteration to minimize idle time of the threads.
+         * If true, the backend load will be adapted on every iteration to minize idle time of the threads.
          */
         nvimgcodecLoadHintPolicy_t load_hint_policy;
     } nvimgcodecBackendParams_t;
@@ -773,19 +800,20 @@ extern "C"
     } nvimgcodecDecodeParams_t;
 
     /**
-     * @brief Supported quality types (algorithms), which determines how `quality_value` is interpreted.
+     * @brief Supported quality types (algorithms), which determines how ``quality_value`` is interpreted.
      */
     typedef enum
     {
-        NVIMGCODEC_QUALITY_TYPE_DEFAULT = 0,           /**< Each plugin decides what is best quality setting to use. `quality_value` is ignored.*/
-        NVIMGCODEC_QUALITY_TYPE_LOSSLESS = 1,          /**< Image encoding is reversible and keeps original image quality. `quality_value` is ignored except for the CUDA tiff encoder backend,
-                                                            for which `quality_value=0` means no compression, and `quality_value=1` means LZW compression. */
-        NVIMGCODEC_QUALITY_TYPE_QUALITY = 2,           /**< `quality_value` is interpreted as JPEG-like quality in range from 1 (worst) to 100 (best). */
-        NVIMGCODEC_QUALITY_TYPE_QUANTIZATION_STEP = 3, /**< `quality_value` is interpreted as quantization step (by how much pixel data will be divided).
+        NVIMGCODEC_QUALITY_TYPE_DEFAULT = 0,           /**< Each plugin decides what is best quality setting to use. ``quality_value``is ignored.*/
+        NVIMGCODEC_QUALITY_TYPE_LOSSLESS = 1,          /**< Image encoding is reversible and keeps original image quality. ``quality_value`` is ignored except for:
+                                                            - the CUDA TIFF encoder backend, for which ``quality_value = 0`` means no compression, and ``quality_value = 1`` means LZW compression,
+                                                            - the CPU PNG encoder backend, for which ``quality_value`` can be in range from 0 (lowest compression ratio, fastest) to 9 (highest compression ratio, slowest). */
+        NVIMGCODEC_QUALITY_TYPE_QUALITY = 2,           /**< ``quality_value`` is interpreted as JPEG-like quality in range from 1 (worst) to 100 (best). */
+        NVIMGCODEC_QUALITY_TYPE_QUANTIZATION_STEP = 3, /**< ``quality_value`` is interpreted as quantization step (by how much pixel data will be divided).
                                                             The higher the value, the worse quality image is produced.*/
-        NVIMGCODEC_QUALITY_TYPE_PSNR = 4,              /**< `quality_value` is interpreted as desired Peak Signal-to-Noise Ratio (PSNR) target for the encoded image.
+        NVIMGCODEC_QUALITY_TYPE_PSNR = 4,              /**< ``quality_value`` is interpreted as desired Peak Signal-to-Noise Ratio (PSNR) target for the encoded image.
                                                             The higher the value, the better quality image is produced. Value should be positive. */
-        NVIMGCODEC_QUALITY_TYPE_SIZE_RATIO = 5,         /**< `quality_value` is interpreted as desired encoded image size ratio compared to original size, should be floating in range (0.0, 1.0).
+        NVIMGCODEC_QUALITY_TYPE_SIZE_RATIO = 5,         /**< ``quality_value`` is interpreted as desired encoded image size ratio compared to original size, should be floating in range (0.0, 1.0).
                                                             E.g. value 0.1 means target size of 10% of original image. */
         NVIMGCODEC_QUALITY_TYPE_ENUM_FORCE_INT = INT32_MAX
     } nvimgcodecQualityType_t;
@@ -915,9 +943,10 @@ extern "C"
         NVIMGCODEC_METADATA_VALUE_TYPE_SRATIONAL  = 10,   /**< Two SLONGs: numerator and denominator */
         NVIMGCODEC_METADATA_VALUE_TYPE_FLOAT      = 11,   /**< 4-byte IEEE floating point value */
         NVIMGCODEC_METADATA_VALUE_TYPE_DOUBLE     = 12,   /**< 8-byte IEEE floating point value */
+        NVIMGCODEC_METADATA_VALUE_TYPE_IFD        = 13,   /**< 4-byte (32-bit) unsigned integer used for IFD offsets */
         NVIMGCODEC_METADATA_VALUE_TYPE_LONG8      = 16,   /**< 8-byte (64-bit) unsigned integer (BigTIFF) */
         NVIMGCODEC_METADATA_VALUE_TYPE_SLONG8     = 17,   /**< 8-byte (64-bit) signed integer (BigTIFF) */
-        NVIMGCODEC_METADATA_VALUE_TYPE_IFD8       = 18,   /**< 8-byte (64-bit) unsigned integer used for offsets (BigTIFF) */
+        NVIMGCODEC_METADATA_VALUE_TYPE_IFD8       = 18,   /**< 8-byte (64-bit) unsigned integer used for IFD offsets (BigTIFF) */
         NVIMGCODEC_METADATA_VALUE_TYPE_ENUM_FORCE_INT = INT32_MAX
     } nvimgcodecMetadataValueType_t;
 
@@ -1835,7 +1864,9 @@ extern "C"
      *        If *code_stream is not NULL, the existing code stream instance will be reused instead of creating a new one.
      * @param file_name [in] File name with compressed image data to wrap.
      * @param code_stream_view [in] Optional pointer to a nvimgcodecCodeStreamView_t struct specifying parsing parameters
-     *        such as bitstream_offset and limit_images. Can be NULL for default behavior.
+     *        such as bitstream_offset. Can be NULL for default behavior.
+     *        Note: image_idx and region are rejected here; use nvimgcodecCodeStreamGetSubCodeStream to select
+     *        a specific image or region after creation.
      * @return nvimgcodecStatus_t - An error code as specified in {@link nvimgcodecStatus_t API Return Status Codes}
      */
     NVIMGCODECAPI nvimgcodecStatus_t nvimgcodecCodeStreamCreateFromFile(
@@ -1852,7 +1883,9 @@ extern "C"
      * @param data [in] Pointer to buffer with compressed data.
      * @param length [in] Length of compressed data in provided buffer.
      * @param code_stream_view [in] Optional pointer to a nvimgcodecCodeStreamView_t struct specifying parsing parameters
-     *        such as bitstream_offset and limit_images. Can be NULL for default behavior.
+     *        such as bitstream_offset. Can be NULL for default behavior.
+     *        Note: image_idx and region are rejected here; use nvimgcodecCodeStreamGetSubCodeStream to select
+     *        a specific image or region after creation.
      * @return nvimgcodecStatus_t - An error code as specified in {@link nvimgcodecStatus_t API Return Status Codes}
      */
     NVIMGCODECAPI nvimgcodecStatus_t nvimgcodecCodeStreamCreateFromHostMem(
@@ -1880,7 +1913,7 @@ extern "C"
      * @param req_size [in] Requested size of buffer.
      * @return Pointer to requested buffer.
      *
-     * @note This function can be called multiple times and requested size can be lower at the end so buffer can be shrunk.
+     * @note This function can be called multiple times and requested size can be lower at the end so buffer can be shrinked.
      */
     typedef unsigned char* (*nvimgcodecResizeBufferFunc_t)(void* ctx, size_t req_size);
 
@@ -1924,6 +1957,7 @@ extern "C"
      *        If *sub_code_stream is NULL, a new code stream instance will be created.
      *        If *sub_code_stream is not NULL, the existing code stream instance will be reused instead of creating a new one.
      * @param code_stream_view [in] Points to a nvimgcodecCodeStreamView_t struct which describes the view of the code stream to be used for the sub-code stream.
+     *        Views selecting a specific image are resolved during this call, so containers may parse metadata up to the requested image.
      * @return nvimgcodecStatus_t - An error code as specified in {@link nvimgcodecStatus_t API Return Status Codes}
      */
     NVIMGCODECAPI nvimgcodecStatus_t nvimgcodecCodeStreamGetSubCodeStream(nvimgcodecCodeStream_t code_stream, nvimgcodecCodeStream_t* sub_code_stream,
@@ -1945,8 +1979,10 @@ extern "C"
      * @param instance  [in] The library instance handle the decoder will be used with.
      * @param decoder  [in/out] Points a nvimgcodecDecoder_t handle in which the decoder is returned.
      * @param exec_params [in] Points an execution parameters.
-     * @param options [in] String with optional space separated list of parameters for specific decoders in format
-     *                     "<decoder_id>:<parameter_name>=<parameter_value>". For example  "nvjpeg:fancy_upsampling=1"
+     * @param options [in] Optional space-separated list of parameters. Use "<decoder_id>:<name>=<value>" for
+     *                     decoder-specific options, or a leading colon for global options (e.g. ":num_cuda_streams=4")
+     *                     or options that any matching decoder may honor (e.g. ":fancy_upsampling=1"). Pass NULL
+     *                     or empty string for no options. See the documentation for the full list of options per decoder.
      * @return nvimgcodecStatus_t - An error code as specified in {@link nvimgcodecStatus_t API Return Status Codes}
      */
     NVIMGCODECAPI nvimgcodecStatus_t nvimgcodecDecoderCreate(
@@ -2013,10 +2049,11 @@ extern "C"
      * @brief Creates generic image encoder.
      *
      * @param instance [in] The library instance handle the encoder will be used with.
-     * @param encoder [in/out] Points a nvimgcodecEncoder_t handle in which the decoder is returned.
+     * @param encoder [in/out] Points a nvimgcodecEncoder_t handle in which the encoder is returned.
      * @param exec_params [in] Points an execution parameters.
-     * @param options [in] String with optional, space separated, list of parameters for specific encoders, in format
-     *                     "<encoder_id>:<parameter_name>=<parameter_value>."
+     * @param options [in] Optional space-separated list of parameters. Use "<encoder_id>:<name>=<value>" for
+     *                     encoder-specific options, or a leading colon for global options (e.g. ":num_cuda_streams=4").
+     *                     Pass NULL or empty string for no options. See the documentation for encoder options.
      * @return nvimgcodecStatus_t - An error code as specified in {@link nvimgcodecStatus_t API Return Status Codes}
      */
     NVIMGCODECAPI nvimgcodecStatus_t nvimgcodecEncoderCreate(
diff --git a/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec_version.h b/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec_version.h
index 4763bf579..4eda4ade3 100644
--- a/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec_version.h
+++ b/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec_version.h
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 // clang-format off
@@ -7,7 +7,7 @@
 #define NVIMGCODEC_VERSION_H__
 
 #define NVIMGCODEC_VER_MAJOR 0
-#define NVIMGCODEC_VER_MINOR 8
+#define NVIMGCODEC_VER_MINOR 9
 #define NVIMGCODEC_VER_PATCH 0
 #define NVIMGCODEC_VER_BUILD
 
diff --git a/cpp/plugins/cucim.kit.cuslide2/tests/CMakeLists.txt b/cpp/plugins/cucim.kit.cuslide2/tests/CMakeLists.txt
index 2ce2744f3..18d6d9993 100644
--- a/cpp/plugins/cucim.kit.cuslide2/tests/CMakeLists.txt
+++ b/cpp/plugins/cucim.kit.cuslide2/tests/CMakeLists.txt
@@ -17,6 +17,7 @@ add_executable(cuslide2_tests
         test_read_region.cpp
         # test_read_rawtiff.cpp  # Disabled: requires libtiff-specific write_offsets_() method not available in nvImageCodec
         test_philips_tiff.cpp
+        test_ome_xml.cpp
         )
 set_target_properties(cuslide2_tests
     PROPERTIES
diff --git a/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp b/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp
new file mode 100644
index 000000000..3653d294e
--- /dev/null
+++ b/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp
@@ -0,0 +1,73 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <catch2/catch_test_macros.hpp>
+
+#include "cuslide/tiff/ome_xml.h"
+
+TEST_CASE("OME parser reads pixels and channels", "[ome][parser]")
+{
+    constexpr const char* kOmeXml = R"OME(
+<OME xmlns="http://www.openmicroscopy.org/Schemas/OME/2016-06">
+  <Image ID="Image:0" Name="multiplex">
+    <Pixels ID="Pixels:0" DimensionOrder="XYZCT" Type="uint16"
+            SizeX="1000" SizeY="500" SizeZ="2" SizeC="3" SizeT="4"
+            PhysicalSizeX="0.325" PhysicalSizeY="0.325">
+      <Channel ID="Channel:0:0" Name="DAPI" SamplesPerPixel="1"/>
+      <Channel ID="Channel:0:1" Name="CD3" SamplesPerPixel="1"/>
+      <Channel ID="Channel:0:2" Name="CD8" SamplesPerPixel="1"/>
+      <TiffData IFD="0" FirstC="0" FirstZ="0" FirstT="0" PlaneCount="4"/>
+    </Pixels>
+  </Image>
+</OME>
+)OME";
+
+    cuslide::tiff::ome::Model model;
+    std::string error;
+    REQUIRE(cuslide::tiff::ome::parse(kOmeXml, &model, &error));
+    REQUIRE(model.valid);
+    REQUIRE(model.pixels.size_x == 1000);
+    REQUIRE(model.pixels.size_y == 500);
+    REQUIRE(model.pixels.size_c == 3);
+    REQUIRE(model.pixels.size_z == 2);
+    REQUIRE(model.pixels.size_t == 4);
+    REQUIRE(model.pixels.type == "uint16");
+    REQUIRE(model.pixels.dimension_order == "XYZCT");
+    REQUIRE(model.pixels.channels.size() == 3);
+    REQUIRE(model.pixels.channels[0].name == "DAPI");
+    REQUIRE(model.pixels.channels[2].name == "CD8");
+    REQUIRE(model.pixels.tiff_data.size() == 1);
+    REQUIRE(model.pixels.tiff_data[0].plane_count == 4);
+}
+
+TEST_CASE("OME parser reads UUID companion references", "[ome][parser]")
+{
+    constexpr const char* kOmeXml = R"OME(
+<OME xmlns="http://www.openmicroscopy.org/Schemas/OME/2016-06">
+  <Image ID="Image:0">
+    <Pixels ID="Pixels:0" DimensionOrder="XYCZT" Type="uint16"
+            SizeX="256" SizeY="256" SizeZ="1" SizeC="2" SizeT="1">
+      <Channel ID="Channel:0:0" Name="Marker0"/>
+      <Channel ID="Channel:0:1" Name="Marker1"/>
+      <TiffData IFD="0" FirstC="0" FirstZ="0" FirstT="0" PlaneCount="1">
+        <UUID FileName="companion_a.ome.tif">urn:uuid:a</UUID>
+      </TiffData>
+      <TiffData IFD="3" FirstC="1" FirstZ="0" FirstT="0" PlaneCount="1">
+        <UUID FileName="companion_b.ome.tif">urn:uuid:b</UUID>
+      </TiffData>
+    </Pixels>
+  </Image>
+</OME>
+)OME";
+
+    cuslide::tiff::ome::Model model;
+    REQUIRE(cuslide::tiff::ome::parse(kOmeXml, &model));
+    REQUIRE(model.pixels.tiff_data.size() == 2);
+    REQUIRE(model.pixels.tiff_data[0].file_name == "companion_a.ome.tif");
+    REQUIRE(model.pixels.tiff_data[0].uuid == "urn:uuid:a");
+    REQUIRE(model.pixels.tiff_data[1].file_name == "companion_b.ome.tif");
+    REQUIRE(model.pixels.tiff_data[1].ifd == 3);
+    REQUIRE(model.pixels.tiff_data[1].first_c == 1);
+}
diff --git a/dependencies.yaml b/dependencies.yaml
index 26a1b9565..b9912575a 100644
--- a/dependencies.yaml
+++ b/dependencies.yaml
@@ -228,7 +228,7 @@ dependencies:
         packages:
           - cuda-cudart-dev
           - libnvjpeg-dev
-          - libnvimgcodec-dev>=0.8.0,<0.9.0
+          - libnvimgcodec-dev>=0.9.0,<0.10.0
     specific:
       - output_types: conda
         matrices:
@@ -314,19 +314,19 @@ dependencies:
           - scipy>=1.11.2
       - output_types: conda
         packages:
-          - libnvimgcodec>=0.8.0,<0.9.0
-          - nvimgcodec>=0.8.0,<0.9.0
+          - libnvimgcodec>=0.9.0,<0.10.0
+          - nvimgcodec>=0.9.0,<0.10.0
     specific:
       - output_types: [requirements, pyproject]
         matrices:
           - matrix:
               cuda: "12.*"
             packages:
-              - nvidia-nvimgcodec-cu12>=0.8.0,<0.9.0
+              - nvidia-nvimgcodec-cu12>=0.9.0,<0.10.0
           # fallback to CUDA 13 versions if 'cuda' is '13.*' or not provided
           - matrix:
             packages:
-              - nvidia-nvimgcodec-cu13>=0.8.0,<0.9.0
+              - nvidia-nvimgcodec-cu13>=0.9.0,<0.10.0
   depends_on_cupy:
     common:
       - output_types: conda
diff --git a/python/cucim/pyproject.toml b/python/cucim/pyproject.toml
index 53c3c08fe..0844f33a2 100644
--- a/python/cucim/pyproject.toml
+++ b/python/cucim/pyproject.toml
@@ -33,7 +33,7 @@ dependencies = [
     "cupy-cuda13x[ctk]>=14.0.1,!=14.1.0",
     "lazy-loader>=0.4",
     "numpy>=2.0,<3.0",
-    "nvidia-nvimgcodec-cu13>=0.8.0,<0.9.0",
+    "nvidia-nvimgcodec-cu13>=0.9.0,<0.10.0",
     "scikit-image>=0.23.2,<0.27.0",
     "scipy>=1.11.2",
 ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`.
@@ -199,7 +199,8 @@ order-by-type = true
 #   codespell --toml python/cucim/pyproject.toml . -i 3 -w
 skip = "build*,dist,.cache,html,_build,_deps,3rdparty/*,_static,generated,latex,.git,*.ipynb,test_data/input/LICENSE-3rdparty,jitify_testing"
 # ignore-regex = ""
-ignore-words-list = "ans,coo,boun,bu,bui,gool,hart,lond,manuel,nd,paeth,unser,wronly,thirdparty,burnin,tge"
+# minize/shrinked: upstream typos in vendored nvImageCodec headers
+ignore-words-list = "ans,coo,boun,bu,bui,gool,hart,lond,manuel,nd,paeth,unser,wronly,thirdparty,burnin,tge,minize,shrinked"
 quiet-level = 3
 
 # to undo: ./test_data/input/LICENSE-3rdparty
diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py
new file mode 100644
index 000000000..fe0daf37b
--- /dev/null
+++ b/scripts/test_pyramidal_ome_tiff.py
@@ -0,0 +1,708 @@
+#!/usr/bin/env python3
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+"""Validation script for pyramidal OME-TIFF with cuslide2.
+
+This script focuses on:
+  - OME metadata presence and basic semantic checks
+  - multi-level pyramid reads
+  - channel-plane selection via read_region kwargs (C/Z/T)
+  - CPU vs GPU decode consistency (where GPU is available)
+  - synthetic uint16 / higher bit-depth decode coverage
+  - optional tile-level caching smoke test
+"""
+
+import json
+import sys
+import tempfile
+import time
+import traceback
+from pathlib import Path
+
+import numpy as np
+from test_common import setup_environment, test_tile_level_caching
+
+
+def _to_numpy(arr):
+    """Convert CuImage/CuPy/NumPy-like objects to NumPy arrays."""
+    if hasattr(arr, "get"):
+        # CuPy array
+        return arr.get()
+    if hasattr(arr, "__cuda_array_interface__"):
+        # Device-resident CuImage. NumPy cannot consume the CUDA array interface and
+        # would silently produce a 0-d object array, which turns every GPU comparison
+        # into a skipped check, so copy to host through CuPy instead.
+        import cupy
+
+        return cupy.asarray(arr).get()
+    return np.asarray(arr)
+
+
+def _max_read_size(level_dims: tuple[int, int], cap: int = 512):
+    return [min(cap, int(level_dims[0])), min(cap, int(level_dims[1]))]
+
+
+def _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count):
+    """Assert pyramid levels shrink and match reported downsample factors.
+
+    cuslide2 reports a single scalar downsample per level, computed as the mean
+    of the per-axis ratios ((w0/wi) + (h0/hi)) / 2. Levels are free to use
+    different X and Y ratios, so dims cannot be recovered from the scalar.
+    Multi-plane files (OME C/Z/T) also expose several IFDs per resolution, so
+    dims are required to be non-increasing rather than strictly decreasing.
+    """
+    print("\n📐 Pyramid geometry checks")
+    if level_count < 1:
+        raise RuntimeError("level_count must be >= 1")
+
+    base_w, base_h = int(level_dimensions[0][0]), int(level_dimensions[0][1])
+    if float(level_downsamples[0]) != 1.0:
+        raise RuntimeError(
+            f"Level 0 downsample should be 1.0, got {level_downsamples[0]}"
+        )
+
+    resolutions = 1
+    for level in range(1, level_count):
+        w, h = int(level_dimensions[level][0]), int(level_dimensions[level][1])
+        ds = float(level_downsamples[level])
+        prev_w = int(level_dimensions[level - 1][0])
+        prev_h = int(level_dimensions[level - 1][1])
+        prev_ds = float(level_downsamples[level - 1])
+
+        if w > prev_w or h > prev_h:
+            raise RuntimeError(
+                f"Level {level} dims {w}x{h} are not smaller than "
+                f"level {level - 1} {prev_w}x{prev_h}"
+            )
+
+        if w == prev_w and h == prev_h:
+            if abs(ds - prev_ds) > 1e-3 * max(1.0, prev_ds):
+                raise RuntimeError(
+                    f"Level {level} repeats dims {w}x{h} from level "
+                    f"{level - 1} but reports downsample {ds} != {prev_ds}"
+                )
+        else:
+            resolutions += 1
+            if ds <= prev_ds:
+                raise RuntimeError(
+                    f"Level {level} downsample {ds} should be > "
+                    f"level {level - 1} downsample {prev_ds}"
+                )
+
+        expected_ds = ((base_w / w) + (base_h / h)) / 2.0
+        if abs(ds - expected_ds) > max(0.01, 0.01 * expected_ds):
+            raise RuntimeError(
+                f"Level {level}: downsample {ds} != mean axis ratio "
+                f"{expected_ds:.4f} for dims {w}x{h} "
+                f"(base {base_w}x{base_h})"
+            )
+
+    print(
+        f"  ✅ Pyramid dims are non-increasing and match reported downsamples "
+        f"({resolutions} distinct resolution(s) over {level_count} level(s))"
+    )
+    if resolutions < level_count:
+        print(
+            f"  ⚠️  {level_count - resolutions} level(s) repeat an existing "
+            f"resolution — likely multi-plane (C/Z/T) IFDs exposed as levels"
+        )
+
+
+def _dtype_bits(dtype) -> int | None:
+    """Best-effort bit depth from CuImage/DLDataType or NumPy dtype."""
+    bits = getattr(dtype, "bits", None)
+    if bits is not None:
+        return int(bits)
+    try:
+        return int(np.dtype(dtype).itemsize * 8)
+    except Exception:
+        return None
+
+
+def _write_synthetic_uint16_ome_tiff(path: Path, height: int = 128, width: int = 128):
+    """Write a small tiled pyramidal uint16 OME-TIFF for decode coverage."""
+    import tifffile
+
+    ome = f"""<?xml version="1.0" encoding="UTF-8"?>
+<OME xmlns="http://www.openmicroscopy.org/Schemas/OME/2016-06">
+  <Image ID="Image:0" Name="uint16-synth">
+    <Pixels ID="Pixels:0" DimensionOrder="XYCZT" Type="uint16"
+            SizeX="{width}" SizeY="{height}" SizeZ="1" SizeC="1" SizeT="1"
+            PhysicalSizeX="1.0" PhysicalSizeY="1.0"
+            PhysicalSizeXUnit="um" PhysicalSizeYUnit="um">
+      <Channel ID="Channel:0:0" Name="C0" SamplesPerPixel="1"/>
+      <TiffData IFD="0" PlaneCount="1"/>
+    </Pixels>
+  </Image>
+</OME>"""
+
+    rng = np.random.default_rng(0)
+    level0 = rng.integers(0, 4096, size=(height, width), dtype=np.uint16)
+    # Plant a recognizable high-bit pattern (>255) so uint8 truncation would fail.
+    level0[10:20, 10:20] = 12345
+    level1 = level0[::2, ::2].copy()
+
+    path.parent.mkdir(parents=True, exist_ok=True)
+    with tifffile.TiffWriter(path, bigtiff=True) as tif:
+        tif.write(
+            level0,
+            description=ome,
+            tile=(64, 64),
+            compression="deflate",
+            photometric="minisblack",
+            metadata=None,
+        )
+        tif.write(
+            level1,
+            tile=(64, 64),
+            compression="deflate",
+            photometric="minisblack",
+            subfiletype=1,
+            metadata=None,
+        )
+    return level0, level1
+
+
+def _validate_uint16_decode(plugin_lib: str):
+    """Exercise the uint16 / higher bit-depth decode path on a synthetic OME-TIFF."""
+    print("\n🧪 Synthetic uint16 decode checks")
+    from cucim import CuImage
+    from cucim.clara import _set_plugin_root
+
+    _set_plugin_root(str(plugin_lib))
+
+    with tempfile.TemporaryDirectory(prefix="cucim_uint16_") as tmp:
+        path = Path(tmp) / "uint16_pyramid.ome.tif"
+        level0, level1 = _write_synthetic_uint16_ome_tiff(path)
+        img = CuImage(str(path))
+
+        bits = _dtype_bits(img.dtype)
+        if bits != 16 and np.dtype(img.typestr) != np.uint16:
+            raise RuntimeError(
+                f"Expected uint16 image dtype, got dtype={img.dtype}, typestr={img.typestr}"
+            )
+        print(f"  ✅ Metadata dtype is uint16 (bits={bits}, typestr={img.typestr})")
+
+        region = _to_numpy(img.read_region((10, 10), (10, 10), level=0, device="cpu"))
+        if region.dtype != np.uint16:
+            raise RuntimeError(f"Expected uint16 region dtype, got {region.dtype}")
+        region2d = region[..., 0] if region.ndim == 3 else region
+        expected = level0[10:20, 10:20]
+        if not np.array_equal(region2d, expected):
+            raise RuntimeError(
+                "uint16 level-0 patch mismatch "
+                f"(max_diff={int(np.max(np.abs(region2d.astype(np.int64) - expected.astype(np.int64))))})"
+            )
+        print(
+            f"  ✅ Level-0 uint16 patch matches source ({region2d.shape}, dtype={region.dtype})"
+        )
+
+        # Values above 255 prove we did not silently narrow to uint8.
+        if int(region2d.max()) <= 255:
+            raise RuntimeError(
+                f"Expected high-bit values in uint16 patch, max={int(region2d.max())}"
+            )
+        print(f"  ✅ High-bit values preserved (max={int(region2d.max())})")
+
+        level1_region = _to_numpy(
+            img.read_region((0, 0), list(level1.shape[::-1]), level=1, device="cpu")
+        )
+        if level1_region.dtype != np.uint16:
+            raise RuntimeError(
+                f"Expected uint16 level-1 dtype, got {level1_region.dtype}"
+            )
+        level1_2d = level1_region[..., 0] if level1_region.ndim == 3 else level1_region
+        if not np.array_equal(level1_2d, level1):
+            raise RuntimeError(
+                "uint16 level-1 decode does not match source pyramid plane"
+            )
+        print(
+            f"  ✅ Level-1 uint16 plane matches source "
+            f"({level1_2d.shape}, dtype={level1_region.dtype})"
+        )
+
+    print("  ✅ Synthetic uint16 decode path validated")
+
+
+def _print_public_dataset_references():
+    print("\n📋 Public Cell DIVE datasets used in this effort:")
+    print("=" * 70)
+    print("Heart sample (5 markers):")
+    print(
+        "  https://portal.hubmapconsortium.org/browse/dataset/e4263715a087881e46ea4d11f49139aa"
+    )
+    print("Skin sample (19 markers):")
+    print(
+        "  https://portal.hubmapconsortium.org/browse/dataset/1b8121539ff16f53681de6108069be24"
+    )
+
+
+def _extract_ome_meta(metadata: dict):
+    """Extract OME metadata from CuImage metadata dict."""
+    ome = metadata.get("ome", {}) if isinstance(metadata, dict) else {}
+    size_c = int(ome.get("size_c", -1)) if isinstance(ome, dict) else -1
+    size_z = int(ome.get("size_z", -1)) if isinstance(ome, dict) else -1
+    size_t = int(ome.get("size_t", -1)) if isinstance(ome, dict) else -1
+    channel_names = ome.get("channel_names", []) if isinstance(ome, dict) else []
+    return ome, size_c, size_z, size_t, channel_names
+
+
+def _load_cell_dive_tsv(tsv_path: str, hubmap_id: str | None = None):
+    """Load Cell DIVE assay keys from HuBMAP-style TSV metadata export."""
+    path = Path(tsv_path)
+    if not path.exists():
+        raise FileNotFoundError(f"TSV not found: {tsv_path}")
+
+    lines = path.read_text(encoding="utf-8").splitlines()
+    if not lines:
+        raise RuntimeError(f"TSV is empty: {tsv_path}")
+
+    # Header expected: HuBMAP ID, Entity, Key, Value, Description
+    records = []
+    for line in lines[1:]:
+        parts = line.split("\t")
+        if len(parts) < 5:
+            continue
+        record = {
+            "hubmap_id": parts[0].strip(),
+            "entity": parts[1].strip(),
+            "key": parts[2].strip(),
+            "value": parts[3].strip(),
+            "description": parts[4].strip(),
+        }
+        records.append(record)
+
+    cell_dive_records = [r for r in records if r["entity"] == "Cell DIVE"]
+    if not cell_dive_records:
+        raise RuntimeError("No 'Cell DIVE' entity rows found in TSV")
+
+    if hubmap_id is None:
+        hubmap_id = cell_dive_records[0]["hubmap_id"]
+    selected = [r for r in cell_dive_records if r["hubmap_id"] == hubmap_id]
+    if not selected:
+        raise RuntimeError(
+            f"No Cell DIVE rows found for HuBMAP ID '{hubmap_id}'. "
+            f"Available IDs: {sorted({r['hubmap_id'] for r in cell_dive_records})}"
+        )
+
+    kv = {r["key"]: r["value"] for r in selected}
+    return hubmap_id, kv
+
+
+def _as_float_or_none(v):
+    if v is None or v == "":
+        return None
+    try:
+        return float(v)
+    except Exception:
+        return None
+
+
+def _as_int_or_none(v):
+    if v is None or v == "":
+        return None
+    try:
+        return int(v)
+    except Exception:
+        return None
+
+
+def _validate_tsv_mapping(
+    tsv_hubmap_id: str, tsv_meta: dict, ome: dict, size_c: int, channel_names
+):
+    """Map selected TSV keys to OME/cuslide2 outputs and validate consistency."""
+    print("\n🧾 TSV → OME mapping checks")
+    print("-" * 50)
+    print(f"  HuBMAP ID: {tsv_hubmap_id}")
+
+    # 1) Pixel spacing mapping:
+    # TSV resolution_x/y_value (um) ↔ OME physical_size_x/y (um)
+    tsv_rx = _as_float_or_none(tsv_meta.get("resolution_x_value"))
+    tsv_ry = _as_float_or_none(tsv_meta.get("resolution_y_value"))
+    ome_px = _as_float_or_none(ome.get("physical_size_x"))
+    ome_py = _as_float_or_none(ome.get("physical_size_y"))
+
+    if tsv_rx is not None and ome_px is not None:
+        dx = abs(tsv_rx - ome_px)
+        print(
+            f"  resolution_x_value ({tsv_rx}) ↔ physical_size_x ({ome_px}), Δ={dx:.6f}"
+        )
+        if dx > 1e-3:
+            raise RuntimeError(
+                f"TSV/OME X spacing mismatch is too large: TSV={tsv_rx}, OME={ome_px}, delta={dx}"
+            )
+    else:
+        print("  ⚠️  Skipping X-spacing strict check (missing TSV or OME value)")
+
+    if tsv_ry is not None and ome_py is not None:
+        dy = abs(tsv_ry - ome_py)
+        print(
+            f"  resolution_y_value ({tsv_ry}) ↔ physical_size_y ({ome_py}), Δ={dy:.6f}"
+        )
+        if dy > 1e-3:
+            raise RuntimeError(
+                f"TSV/OME Y spacing mismatch is too large: TSV={tsv_ry}, OME={ome_py}, delta={dy}"
+            )
+    else:
+        print("  ⚠️  Skipping Y-spacing strict check (missing TSV or OME value)")
+
+    # 2) Channel-count heuristics:
+    # - number_of_antibodies is expected to be <= decoded size_c in most assembled datasets.
+    # - number_of_channels is per-round and not equal to final SizeC, so only soft check.
+    n_antibodies = _as_int_or_none(tsv_meta.get("number_of_antibodies"))
+    n_channels_per_round = _as_int_or_none(tsv_meta.get("number_of_channels"))
+    n_rounds = _as_int_or_none(tsv_meta.get("number_of_total_imaging_rounds"))
+
+    print(f"  OME size_c={size_c}")
+    if n_antibodies is not None:
+        print(f"  TSV number_of_antibodies={n_antibodies}")
+        if size_c > 0 and size_c < n_antibodies:
+            raise RuntimeError(
+                f"Decoded SizeC ({size_c}) is smaller than TSV antibody count ({n_antibodies})"
+            )
+    if n_channels_per_round is not None:
+        print(f"  TSV number_of_channels (per round)={n_channels_per_round}")
+    if n_rounds is not None:
+        print(f"  TSV number_of_total_imaging_rounds={n_rounds}")
+
+    # 3) Nuclear marker mapping:
+    # TSV nuclear_marker_or_stain often expected among OME channel names.
+    nuc = (tsv_meta.get("nuclear_marker_or_stain") or "").strip()
+    if nuc and channel_names:
+        lower_names = {str(x).strip().lower() for x in channel_names}
+        print(f"  TSV nuclear_marker_or_stain={nuc}")
+        if nuc.lower() in lower_names:
+            print("  ✅ nuclear marker found in OME channel names")
+        else:
+            print("  ⚠️  nuclear marker not found in OME channel names (non-fatal)")
+    elif nuc:
+        print("  ⚠️  channel names missing, cannot verify nuclear marker mapping")
+
+    # 4) Instrument/model metadata (soft checks)
+    vendor = (tsv_meta.get("acquisition_instrument_vendor") or "").strip()
+    model = (tsv_meta.get("acquisition_instrument_model") or "").strip()
+    if vendor:
+        print(f"  TSV acquisition_instrument_vendor={vendor}")
+    if model:
+        print(f"  TSV acquisition_instrument_model={model}")
+
+
+def _validate_channel_selection(img, level_dims, level, c_index, z_index=0, t_index=0):
+    """Validate per-plane selection using read_region kwargs."""
+    read_size = _max_read_size(level_dims)
+    kwargs = {"C": int(c_index), "Z": int(z_index), "T": int(t_index)}
+
+    # CPU read
+    start = time.time()
+    cpu_region = img.read_region((0, 0), read_size, level=level, device="cpu", **kwargs)
+    cpu_time = time.time() - start
+    cpu_np = _to_numpy(cpu_region)
+
+    print(
+        f"    CPU plane read C={c_index},Z={z_index},T={t_index}: "
+        f"{cpu_np.shape}, {cpu_np.dtype}, {cpu_time:.4f}s"
+    )
+
+    # GPU read (optional, skip if unavailable)
+    try:
+        start = time.time()
+        gpu_region = img.read_region(
+            (0, 0), read_size, level=level, device="cuda", **kwargs
+        )
+        gpu_time = time.time() - start
+        gpu_np = _to_numpy(gpu_region)
+        print(
+            f"    GPU plane read C={c_index},Z={z_index},T={t_index}: "
+            f"{gpu_np.shape}, {gpu_np.dtype}, {gpu_time:.4f}s"
+        )
+
+        if gpu_np.shape != cpu_np.shape:
+            raise RuntimeError(
+                f"Shape mismatch for plane C={c_index},Z={z_index},T={t_index}: "
+                f"GPU={gpu_np.shape}, CPU={cpu_np.shape}"
+            )
+        if not np.array_equal(gpu_np, cpu_np):
+            max_diff = int(
+                np.max(np.abs(gpu_np.astype(np.int64) - cpu_np.astype(np.int64)))
+            )
+            raise RuntimeError(
+                f"Pixel mismatch for plane C={c_index},Z={z_index},T={t_index}: max_diff={max_diff}"
+            )
+        print("    ✅ GPU and CPU plane decode are identical")
+    except Exception as e:
+        print(f"    ⚠️  GPU plane validation skipped/failed: {e}")
+
+
+def _validate_batch_decode(img, level_dims):
+    """Validate batch decode path (multiple tile reads)."""
+    print("\n🔄 Batch decode validation")
+    print("-" * 50)
+
+    tile_w = min(256, int(level_dims[0]))
+    tile_h = min(256, int(level_dims[1]))
+    if tile_w <= 0 or tile_h <= 0:
+        print("  ⚠️  Invalid level dimensions; skipping batch check")
+        return
+
+    max_x = max(0, int(level_dims[0]) - tile_w)
+    max_y = max(0, int(level_dims[1]) - tile_h)
+    locations = []
+    for y_off in range(0, min(max_y + 1, tile_h * 4), tile_h):
+        for x_off in range(0, min(max_x + 1, tile_w * 4), tile_w):
+            locations.append([x_off, y_off])
+
+    if len(locations) < 2:
+        print("  ⚠️  Not enough locations for batch test")
+        return
+
+    batch_size = min(8, len(locations))
+    print(
+        f"  Locations={len(locations)}, tile={tile_w}x{tile_h}, batch_size={batch_size}"
+    )
+
+    # CPU ground truth
+    start = time.time()
+    cpu_tiles = list(
+        img.read_region(
+            location=locations,
+            size=[tile_w, tile_h],
+            level=0,
+            device="cpu",
+            batch_size=batch_size,
+            num_workers=1,
+        )
+    )
+    cpu_time = time.time() - start
+    print(f"  CPU batch decode: {cpu_time:.4f}s ({len(cpu_tiles)} tiles)")
+
+    # GPU batch compare (optional)
+    try:
+        start = time.time()
+        gpu_tiles = list(
+            img.read_region(
+                location=locations,
+                size=[tile_w, tile_h],
+                level=0,
+                device="cuda",
+                batch_size=batch_size,
+                num_workers=1,
+            )
+        )
+        gpu_time = time.time() - start
+        print(f"  GPU batch decode: {gpu_time:.4f}s ({len(gpu_tiles)} tiles)")
+        if gpu_time > 0:
+            print(f"  🎯 Batch speedup: {cpu_time / gpu_time:.2f}x")
+
+        if len(cpu_tiles) != len(gpu_tiles):
+            raise RuntimeError(
+                f"Tile count mismatch CPU={len(cpu_tiles)} GPU={len(gpu_tiles)}"
+            )
+
+        mismatch = 0
+        for idx, (cpu_tile, gpu_tile) in enumerate(zip(cpu_tiles, gpu_tiles)):
+            c = _to_numpy(cpu_tile)
+            g = _to_numpy(gpu_tile)
+            if c.shape != g.shape or not np.array_equal(c, g):
+                mismatch += 1
+                if mismatch <= 3:
+                    print(
+                        f"    ❌ mismatch at tile idx={idx}, location={locations[idx]}"
+                    )
+        if mismatch:
+            raise RuntimeError(
+                f"Batch decode mismatch count: {mismatch}/{len(cpu_tiles)}"
+            )
+        print("  ✅ Batch decode GPU and CPU are identical")
+    except Exception as e:
+        print(f"  ⚠️  GPU batch validation skipped/failed: {e}")
+
+
+def test_pyramidal_ome_tiff(
+    file_path,
+    plugin_lib,
+    run_cache=True,
+    tsv_path: str | None = None,
+    hubmap_id: str | None = None,
+):
+    print("=" * 70)
+    print("🔬 Testing pyramidal OME-TIFF with cuslide2")
+    print("=" * 70)
+    print(f"📁 File: {file_path}")
+
+    if not Path(file_path).exists():
+        raise FileNotFoundError(f"OME-TIFF file not found: {file_path}")
+
+    from cucim import CuImage
+    from cucim.clara import _set_plugin_root
+
+    _set_plugin_root(str(plugin_lib))
+    print(f"✅ Plugin root set: {plugin_lib}")
+
+    print("\n📂 Loading image...")
+    start = time.time()
+    img = CuImage(file_path)
+    load_time = time.time() - start
+    print(f"✅ Loaded in {load_time:.3f}s")
+
+    print("\n📊 Image summary")
+    print(f"  Shape: {img.shape}")
+    print(f"  Dtype: {img.dtype}")
+    print(f"  Device: {img.device}")
+    level_count = img.resolutions["level_count"]
+    print(f"  Levels: {level_count}")
+    level_dimensions = img.resolutions["level_dimensions"]
+    level_downsamples = img.resolutions["level_downsamples"]
+    for level in range(level_count):
+        dims = level_dimensions[level]
+        ds = level_downsamples[level]
+        print(f"  Level {level}: {dims[0]}x{dims[1]} (downsample: {ds:.3f}x)")
+
+    image_bits = _dtype_bits(img.dtype)
+    expect_uint16 = image_bits == 16 or np.dtype(img.typestr) == np.uint16
+    if expect_uint16:
+        print("  ✅ Input reported as uint16 — decode checks will enforce dtype")
+
+    print("\n🧬 OME metadata checks")
+    metadata = img.metadata
+    if isinstance(metadata, str):
+        try:
+            metadata = json.loads(metadata)
+        except Exception:
+            metadata = {}
+    ome, size_c, size_z, size_t, channel_names = _extract_ome_meta(metadata)
+    if ome:
+        print("  ✅ Found OME metadata block")
+        print(f"  size_c={size_c}, size_z={size_z}, size_t={size_t}")
+        print(f"  channel_names_count={len(channel_names)}")
+    else:
+        print("  ⚠️  OME metadata block not found in img.metadata")
+
+    if tsv_path:
+        tsv_hubmap_id, tsv_meta = _load_cell_dive_tsv(tsv_path, hubmap_id)
+        _validate_tsv_mapping(tsv_hubmap_id, tsv_meta, ome, size_c, channel_names)
+
+    _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count)
+
+    # Multi-level single read smoke checks
+    print("\n🧪 Multi-level decode checks")
+    test_levels = sorted(set([0, min(level_count - 1, 1), level_count - 1]))
+    for level in test_levels:
+        dims = level_dimensions[level]
+        read_size = _max_read_size(dims)
+        start = time.time()
+        cpu_region = img.read_region((0, 0), read_size, level=level, device="cpu")
+        cpu_time = time.time() - start
+        cpu_np = _to_numpy(cpu_region)
+        print(f"  CPU level {level}: {cpu_np.shape}, {cpu_np.dtype}, {cpu_time:.4f}s")
+        if expect_uint16 and cpu_np.dtype != np.uint16:
+            raise RuntimeError(
+                f"Expected uint16 CPU decode at level {level}, got {cpu_np.dtype}"
+            )
+        try:
+            start = time.time()
+            gpu_region = img.read_region((0, 0), read_size, level=level, device="cuda")
+            gpu_time = time.time() - start
+            gpu_np = _to_numpy(gpu_region)
+            print(
+                f"  GPU level {level}: {gpu_np.shape}, {gpu_np.dtype}, {gpu_time:.4f}s"
+            )
+            if expect_uint16 and gpu_np.dtype != np.uint16:
+                raise RuntimeError(
+                    f"Expected uint16 GPU decode at level {level}, got {gpu_np.dtype}"
+                )
+            if gpu_np.shape != cpu_np.shape:
+                raise RuntimeError(
+                    f"Shape mismatch at level {level}: CPU={cpu_np.shape}, GPU={gpu_np.shape}"
+                )
+            if not np.array_equal(gpu_np, cpu_np):
+                max_diff = int(
+                    np.max(np.abs(gpu_np.astype(np.int64) - cpu_np.astype(np.int64)))
+                )
+                raise RuntimeError(
+                    f"Pixel mismatch at level {level}: max_diff={max_diff}"
+                )
+            print(f"  ✅ GPU and CPU level {level} decode are identical")
+        except Exception as e:
+            print(f"  ⚠️  GPU level {level} validation skipped/failed: {e}")
+
+    # Plane-selection checks (if OME C/Z/T metadata is available)
+    if size_c > 0:
+        print("\n🎯 Plane selection checks (C/Z/T kwargs)")
+        dims0 = level_dimensions[0]
+        c_candidates = [0]
+        if size_c > 1:
+            c_candidates.append(1)
+        z_idx = 0 if size_z <= 0 else min(size_z - 1, 0)
+        t_idx = 0 if size_t <= 0 else min(size_t - 1, 0)
+        for c_idx in c_candidates:
+            _validate_channel_selection(
+                img, dims0, level=0, c_index=c_idx, z_index=z_idx, t_index=t_idx
+            )
+    else:
+        print("\n⚠️  Skipping C/Z/T plane checks (size_c not available)")
+
+    _validate_batch_decode(img, level_dimensions[0])
+
+    if run_cache:
+        print("\n💾 Tile-level cache check")
+        test_tile_level_caching(img, file_path, CuImage)
+
+    print("\n✅ Pyramidal OME-TIFF test completed")
+
+
+def main():
+    if len(sys.argv) < 2:
+        print(
+            "Usage: python test_pyramidal_ome_tiff.py <path_to_ome_tiff> "
+            "[--no-cache] [--tsv <path>] [--hubmap-id <ID>]"
+        )
+        print("")
+        print("Examples:")
+        print("  python test_pyramidal_ome_tiff.py /data/sample.ome.tif")
+        print("  python test_pyramidal_ome_tiff.py /data/sample.ome.tiff --no-cache")
+        print(
+            "  python test_pyramidal_ome_tiff.py /data/sample.ome.tif "
+            "--tsv /path/HBM388.RPTF.754.tsv --hubmap-id HBM388.RPTF.754"
+        )
+        _print_public_dataset_references()
+        return 1
+
+    file_path = sys.argv[1]
+    run_cache = "--no-cache" not in sys.argv[2:]
+    tsv_path = None
+    hubmap_id = None
+
+    argv = sys.argv[2:]
+    i = 0
+    while i < len(argv):
+        if argv[i] == "--tsv" and i + 1 < len(argv):
+            tsv_path = argv[i + 1]
+            i += 2
+            continue
+        if argv[i] == "--hubmap-id" and i + 1 < len(argv):
+            hubmap_id = argv[i + 1]
+            i += 2
+            continue
+        i += 1
+
+    plugin_lib = setup_environment("cucim_ome_tiff_test")
+    try:
+        # Always cover the uint16 decode path, independent of the user-provided file.
+        _validate_uint16_decode(plugin_lib)
+        test_pyramidal_ome_tiff(
+            file_path,
+            plugin_lib,
+            run_cache=run_cache,
+            tsv_path=tsv_path,
+            hubmap_id=hubmap_id,
+        )
+        return 0
+    except Exception as e:
+        print(f"\n❌ Test failed: {e}")
+        traceback.print_exc()
+        return 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())