From 70e0e73371c13534981a791923ce42a941287c8c Mon Sep 17 00:00:00 2001 From: cdinea Date: Tue, 21 Jul 2026 21:52:29 -0700 Subject: [PATCH 01/17] cuslide2: add OME-TIFF semantics and nvImageCodec 0.9 support --- cpp/cmake/deps/nvimgcodec.cmake | 57 +- cpp/plugins/cucim.kit.cuslide2/CMakeLists.txt | 2 + .../cmake/deps/nvimgcodec.cmake | 74 +-- .../src/cuslide/cuslide.cpp | 126 ++++- .../cuslide/nvimgcodec/nvimgcodec_decoder.cpp | 91 +++- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 112 +--- .../src/cuslide/tiff/ifd.cpp | 98 ++-- .../cucim.kit.cuslide2/src/cuslide/tiff/ifd.h | 9 +- .../src/cuslide/tiff/ome_xml.cpp | 171 ++++++ .../src/cuslide/tiff/ome_xml.h | 64 +++ .../src/cuslide/tiff/tiff.cpp | 495 ++++++++++++++---- .../src/cuslide/tiff/tiff.h | 22 + .../src/cuslide/tiff/tiff_constants.h | 2 + .../src/cuslide/tiff/types.h | 4 +- .../nvimgcodec_dynlink/nvimgcodec_version.h | 2 +- .../cucim.kit.cuslide2/tests/CMakeLists.txt | 1 + .../cucim.kit.cuslide2/tests/test_ome_xml.cpp | 74 +++ 17 files changed, 1073 insertions(+), 331 deletions(-) create mode 100644 cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp create mode 100644 cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h create mode 100644 cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp diff --git a/cpp/cmake/deps/nvimgcodec.cmake b/cpp/cmake/deps/nvimgcodec.cmake index 38e5689f9..97a44008c 100644 --- a/cpp/cmake/deps/nvimgcodec.cmake +++ b/cpp/cmake/deps/nvimgcodec.cmake @@ -6,6 +6,22 @@ # if (NOT TARGET deps::nvimgcodec) + set(NVIMGCODEC_SONAME_CANDIDATES + "libnvimgcodec.so.0.9.0" + "libnvimgcodec.so.0.8.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() + # First try to find it as a package find_package(nvimgcodec QUIET) @@ -22,31 +38,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 +75,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 +85,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() 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..4759fa8dd 100644 --- a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake +++ b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake @@ -9,6 +9,22 @@ if (NOT TARGET deps::nvimgcodec) set(NVIMGCODEC_LIB_PATH "") set(NVIMGCODEC_INCLUDE_PATH "") + set(NVIMGCODEC_SONAME_CANDIDATES + "libnvimgcodec.so.0.9.0" + "libnvimgcodec.so.0.8.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() + # Try find_package first (works in conda and system installations) find_package(nvimgcodec QUIET) @@ -23,11 +39,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 +48,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 +86,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 +98,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,14 +109,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) - 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() diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp index 790609a4e..3f6c0bdf6 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 #include #include +#include #include using json = nlohmann::json; @@ -98,27 +99,24 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo size_t level_count = tif->level_count(); // Detect if this is an Aperio SVS file - // Try ImageDescription first + // Try ImageDescription first (works with nvImageCodec 0.7.0+) bool is_aperio_svs = (tif->ifd(0)->image_description().rfind("Aperio", 0) == 0); // 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); - // 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 detection for nvImageCodec 0.6.0: check for multiple resolution levels + // Aperio SVS files typically have 3-6 IFDs with multiple resolution levels + // If we have multiple IFDs and they look like a pyramid, treat 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 +127,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; } @@ -144,7 +142,7 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo 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) { @@ -181,32 +179,109 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo std::pmr::vector 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 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(); + } + 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()); + } + } + } + if (ome.contains("physical_size_x") && ome["physical_size_x"].is_number()) + { + ome_physical_size_x = ome["physical_size_x"].get(); + 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(); + 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(std::max(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(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(ome_size_c); + shape[2] = n_ch; } + else + { + shape[2] = n_ch; + } + std::pmr::vector channel_names(&resource); channel_names.reserve(n_ch); - if (n_ch == 3) + auto add_owned_channel_name = [&resource, &channel_names](const std::string& name) { + void* raw = resource.allocate(name.size() + 1, alignof(char)); + auto* buf = static_cast(raw); + memcpy(buf, name.c_str(), name.size() + 1); + channel_names.emplace_back(buf, name.size()); + }; + if (!ome_channel_names.empty()) + { + for (size_t i = 0; i < static_cast(n_ch); ++i) + { + if (i < ome_channel_names.size()) + { + add_owned_channel_name(ome_channel_names[i]); + } + else + { + add_owned_channel_name(fmt::format("C{}", i)); + } + } + } + else if (n_ch == 3) { channel_names.emplace_back(std::string_view{ "R" }); channel_names.emplace_back(std::string_view{ "G" }); channel_names.emplace_back(std::string_view{ "B" }); } - else + else if (n_ch == 4) { 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" }); } + else + { + for (uint8_t i = 0; i < n_ch; ++i) + { + add_owned_channel_name(fmt::format("C{}", i)); + } + } // Spacing units static constexpr std::string_view empty_unit{""}; @@ -222,7 +297,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(ome_physical_size_y)); + spacing.emplace_back(static_cast(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 +388,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(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..47ad830f6 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 @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include @@ -215,6 +216,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; + if constexpr (std::is_integral_v) + { + photometric = static_cast(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 +385,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 = 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; + // Note: buffer_size removed in nvImageCodec v0.7.0 - size is inferred from plane_info output_image_info.cuda_stream = cuda_stream; // Step 4: Provide output buffer @@ -475,7 +529,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, } // ============================================================================ -// Batch ROI Decoding +// Batch ROI Decoding (nvImageCodec v0.7.0+) // ============================================================================ std::vector decode_batch_regions_nvimgcodec( @@ -623,9 +677,9 @@ std::vector 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 = 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 +695,9 @@ std::vector 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 +706,7 @@ std::vector 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 +980,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 = region.width * num_channels * decode_spec.bytes_per_sample; size_t buffer_size = row_stride * region.height; void* buffer_ptr = nullptr; @@ -960,7 +1015,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 +1024,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_tiff_parser.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/nvimgcodec/nvimgcodec_tiff_parser.cpp index d9352b0dc..38b917486 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 @@ -13,7 +13,6 @@ // - 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 // (COMPRESSION, SUBFILETYPE, IMAGEDESCRIPTION, JPEGTABLES, ...) // // ============================================================================ @@ -21,10 +20,8 @@ #include "nvimgcodec_tiff_parser.h" #include // for std::transform -#include // for stderr -#include // for std::atexit, std::getenv, std::strtol +#include // for std::atexit #include // for strlen -#include // for std::thread::hardware_concurrency #include #ifdef CUCIM_HAS_NVIMGCODEC @@ -54,7 +51,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 +63,8 @@ namespace cuslide2::nvimgcodec // NvimgcodecLoadSymbol() + dladdr() (not the stub address) // 3. Probes /extensions/ and /../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() @@ -146,83 +143,9 @@ static std::string detect_nvimgcodec_extensions_path() } -// ============================================================================ -// Decoder thread count heuristic -// ============================================================================ -// -// Computes a reasonable max_num_cpu_threads for nvImageCodec's internal -// threadpool. On multi-worker deployments the total thread count across -// all workers should not exceed the physical core count to avoid -// context-switch overhead. -// -// Heuristic: min(hardware_concurrency / 4, kHeuristicMaxThreads) -// -// The CUDA stream count is left at nvImageCodec's default because -// empirically it has limited performance impact and over-allocating -// streams can itself cause oversubscription. -// -// Overridable via CUCIM_MAX_DECODER_THREADS environment variable. -// - value > 0 : use that exact thread count -// - value == 0: fall back to nvImageCodec default (= num cpu cores) -// - malformed : warn and fall through to heuristic -// -static constexpr int kHeuristicMaxThreads = 8; - -static int compute_max_decoder_threads() -{ - const char* env_val = std::getenv("CUCIM_MAX_DECODER_THREADS"); - if (env_val) - { - char* end = nullptr; - long val = std::strtol(env_val, &end, 10); - if (end == env_val || *end != '\0') - { - fmt::print(stderr, - "[cuslide2] WARNING: CUCIM_MAX_DECODER_THREADS='{}' is not a " - "valid integer — ignoring and using heuristic\n", env_val); - } - else if (val > 0) - { - #ifdef DEBUG - fmt::print("[cuslide2] CUCIM_MAX_DECODER_THREADS={}\n", val); - #endif - return static_cast(val); - } - else if (val == 0) - { - return 0; // explicit 0 → nvImageCodec default - } - else - { - fmt::print(stderr, - "[cuslide2] WARNING: CUCIM_MAX_DECODER_THREADS={} is negative " - "— ignoring and using heuristic\n", val); - } - } - - unsigned int hw_threads = std::thread::hardware_concurrency(); - if (hw_threads == 0) return 0; // unknown → let nvImageCodec decide - - // Fair share: assume this process is one of potentially several workers. - // Cap at kHeuristicMaxThreads to keep the threadpool bounded even on - // high-core-count machines (e.g. 64 cores / 4 = 16 is still too many). - int fair_share = std::min( - std::max(1, static_cast(hw_threads) / 4), - kHeuristicMaxThreads); - - #ifdef DEBUG - fmt::print("[nvimgcodec] max_decoder_threads heuristic: hw_threads={}, " - "fair_share={}, cap={}\n", - hw_threads, fair_share, kHeuristicMaxThreads); - #endif - - return fair_share; -} - - // nvimgcodec API compatibility // -// TIFF-tag retrieval via nvimgcodecDecoderGetMetadata. +// TIFF-tag retrieval via nvimgcodecDecoderGetMetadata (nvImageCodec >= 0.7.0). // 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. @@ -351,7 +274,7 @@ NvImageCodecTiffParserManager::NvImageCodecTiffParserManager() exec_params.struct_next = nullptr; exec_params.device_allocator = nullptr; exec_params.pinned_allocator = nullptr; - exec_params.max_num_cpu_threads = compute_max_decoder_threads(); + exec_params.max_num_cpu_threads = 0; exec_params.executor = nullptr; exec_params.device_id = NVIMGCODEC_DEVICE_CURRENT; // GPU-enabled for decode + metadata exec_params.pre_init = 0; @@ -456,8 +379,7 @@ TiffFileParser::TiffFileParser(const std::string& file_path) nvimgcodecStatus_t status = nvimgcodecCodeStreamCreateFromFile( manager.get_instance(), &main_code_stream_, - file_path.c_str(), - nullptr + file_path.c_str() ); if (status != NVIMGCODEC_STATUS_SUCCESS) @@ -631,9 +553,11 @@ 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+. + // Current limitation (nvImageCodec v0.6.0): + // - codec_name returns "tiff" (container format) not compression type + // - Individual TIFF tags not exposed through metadata API + // - Only vendor-specific metadata blobs available (MED_APERIO, MED_PHILIPS, etc.) + // if (ifd_info.codec == "tiff") { @@ -958,6 +882,13 @@ void TiffFileParser::extract_ifd_metadata(IfdInfo& ifd_info) } } + // WORKAROUND for nvImageCodec 0.6.0: Philips TIFF metadata limitation + // ======================================================================== + // nvImageCodec 0.6.0 does NOT expose: + // 1. Individual TIFF tags (SOFTWARE, ImageDescription, etc.) + // 2. Philips format detection for some files + // + } const IfdInfo& TiffFileParser::get_ifd(uint32_t index) const @@ -1005,7 +936,7 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info) } // ======================================================================== - // Direct TIFF Tag Retrieval by ID (nvImageCodec 0.7.0+) + // 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. @@ -1153,7 +1084,8 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info) } #endif // CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA - // Fallback: file extension heuristics when COMPRESSION tag is not present in file + // Fallback: file extension heuristics when COMPRESSION tag is not available + // (either nvImageCodec < 0.7.0 or tag not present in file) #ifdef DEBUG fmt::print(" ℹ️ COMPRESSION tag not available, using file extension heuristics\n"); #endif // DEBUG 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..6d77d15c6 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp @@ -112,7 +112,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 +131,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 +149,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(subfile_type); #ifdef DEBUG @@ -158,7 +166,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 +174,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,11 +200,27 @@ 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(std::stoul(photo_str)); + } + catch (...) + { + // Keep fallback below. + } + } } // 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) @@ -204,7 +229,8 @@ IFD::IFD(TIFF* tiff, uint16_t index, const ::cuslide2::nvimgcodec::IfdInfo& ifd_ 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 +280,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 +298,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(std::max(1, samples_per_pixel_)); int ndim = 3; + const size_t bytes_per_sample = std::max(1, bits_per_sample_ / 8); - size_t one_raster_size = static_cast(w) * static_cast(h) * samples_per_pixel_; + size_t one_raster_size = static_cast(w) * static_cast(h) * n_ch * bytes_per_sample; size_t raster_size = one_raster_size; void* raster = nullptr; @@ -419,7 +446,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(raster_type), out_device.index() }; - out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 }; + out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast(bits_per_sample_), 1 }; out_image_data->container.ndim = ndim; out_image_data->container.shape = static_cast(cucim_malloc(ndim * sizeof(int64_t))); if (ndim == 4) @@ -456,9 +483,10 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff, output_buffer = static_cast(out_buf->data); } - // Get IFD info from TiffFileParser - const auto& ifd_info = tiff->nvimgcodec_parser_->get_ifd(static_cast(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 @@ -473,9 +501,10 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff, // 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(width_) && ey < static_cast(height_)); @@ -497,7 +526,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(1, samples_per_pixel_); + const uint32_t bytes_per_px = static_cast(std::max(1, bits_per_sample_ / 8)); const uint8_t background_value = tiff->background_value_; // Tile grid offsets that the ROI overlaps @@ -542,7 +572,7 @@ bool IFD::read([[maybe_unused]] const TIFF* tiff, host_raster = host_raster_owner.get(); } - const uint32_t dest_row_stride = static_cast(w) * samples; + const uint32_t dest_row_stride = static_cast(w) * samples * bytes_per_px; uint8_t* dest_row_ptr = host_raster; #ifdef DEBUG @@ -567,15 +597,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(actual_tw) * actual_th * samples; - const uint32_t tile_row_stride = actual_tw * samples; + const size_t tile_buf_size = static_cast(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 +707,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 +763,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(raster_type), out_device.index() }; - out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 }; + out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast(bits_per_sample_), 1 }; out_image_data->container.ndim = 3; out_image_data->container.shape = static_cast(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(1, samples_per_pixel_); out_image_data->container.strides = nullptr; out_image_data->container.byte_offset = 0; @@ -776,12 +806,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(out_device.type()), out_device.index() }; - out_image_data->container.dtype = DLDataType{ kDLUInt, 8, 1 }; + out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast(bits_per_sample_), 1 }; out_image_data->container.ndim = 3; out_image_data->container.shape = static_cast(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(1, samples_per_pixel_); out_image_data->container.strides = nullptr; out_image_data->container.byte_offset = 0; @@ -1002,9 +1032,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..c11f487b6 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.h @@ -34,7 +34,12 @@ class EXPORT_VISIBLE IFD : public std::enable_shared_from_this 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 // 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..6e5d1ded1 --- /dev/null +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp @@ -0,0 +1,171 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "ome_xml.h" + +#include +#include + +#include + +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(1, pixels_node.attribute("SizeC").as_uint(1)); + pixels.size_z = std::max(1, pixels_node.attribute("SizeZ").as_uint(1)); + pixels.size_t = std::max(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(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(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..290737173 --- /dev/null +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h @@ -0,0 +1,64 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include + +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 channels; + std::vector 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..17ca5367e 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 #include +#include +#include #include #include -#include #include #include #include +#include "ome_xml.h" #include "tiff_constants.h" #include @@ -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& 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 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(this, ifd_index, ifd_info); + auto ifd = std::make_shared( + 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(std::hash{}(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(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( + this, + static_cast(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 (kind=5 in v0.8.0) 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 @@ -576,18 +662,21 @@ void TIFF::resolve_vendor_format() } // Detect Philips TIFF - // Check for Philips TIFF using SOFTWARE tag, ImageDescription XML, or nvImageCodec metadata kind + // NOTE: nvImageCodec 0.6.0 doesn't expose individual TIFF tags (like SOFTWARE) + // Workaround: Check for Philips XML in ImageDescription or use nvImageCodec metadata kind { bool is_philips = false; + // Method 1: Check SOFTWARE tag (available in nvImageCodec 0.7.0+) 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 + // (Workaround for nvImageCodec 0.6.0 where SOFTWARE tag is not available) if (!is_philips) { auto& image_desc = first_ifd->image_description(); @@ -598,11 +687,11 @@ void TIFF::resolve_vendor_format() } } - // Method 3: Check metadata_blobs for Philips (MED_PHILIPS) + // Method 3: Check metadata_blobs for Philips (kind=6 in v0.8.0) 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 +731,28 @@ void TIFF::resolve_vendor_format() } } - const std::remove_reference_tget_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 +761,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 +805,208 @@ void TIFF::resolve_vendor_format() } } +void TIFF::_populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr& first_ifd) +{ + (void)ifd_count; + json* json_metadata = reinterpret_cast(metadata); + + ome::Model model; + std::string parse_error; + if (!ome::parse(first_ifd->image_description(), &model, &parse_error) || !model.valid) + { + return; + } + + ome_size_c_ = std::max(1, model.pixels.size_c); + ome_size_z_ = std::max(1, model.pixels.size_z); + ome_size_t_ = std::max(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(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 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> 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(); + if (!pending_planes.empty()) + { + std::vector> level_dims; + level_dims.reserve(ifds_.size()); + for (const auto& ifd : ifds_) + { + std::pair 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 dim = { ifd->width(), ifd->height() }; + auto it = std::find(level_dims.begin(), level_dims.end(), dim); + return (it == level_dims.end()) ? 0 : static_cast(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 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; + } + } + } + if (!level_representative_ifd.empty()) + { + 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& first_ifd) { json* json_metadata = reinterpret_cast(metadata); @@ -785,12 +1077,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 +1092,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,11 +1124,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); - if (downsample > 0) - { - ifd->width_ = width_l0 / downsample; - ifd->height_ = height_l0 / downsample; - } + // Fix width and height of IFD + ifd->width_ = width_l0 / downsample; + ifd->height_ = height_l0 / downsample; ++spacing_index; } else @@ -885,23 +1165,27 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s { auto node_offset = associated_node.node().offset_debug(); - size_t image_desc_len = first_ifd->image_description().size(); - if (node_offset >= 0 && static_cast(node_offset) < image_desc_len) + if (node_offset >= 0) { - const char* data_ptr = image_desc_cstr + node_offset; - const char* desc_end = image_desc_cstr + image_desc_len; - - // Find '>' that closes the opening tag - while (data_ptr < desc_end && *data_ptr != '>') + // `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) + // + // + + // 34 is from `Attribute Name="PIM_DP_IMAGE_DATA"` + char* data_ptr = const_cast(image_desc_cstr) + node_offset + 34; + uint32_t data_len = 0; + while (*data_ptr != '>' && *data_ptr != '\0') { ++data_ptr; } - uint32_t data_len = 0; - if (data_ptr < desc_end && *data_ptr != '\0') + if (*data_ptr != '\0') { ++data_ptr; // start of base64-encoded data - const char* data_end_ptr = data_ptr; - while (data_end_ptr < desc_end && *data_end_ptr != '<' && *data_end_ptr != '\0') + char* data_end_ptr = data_ptr; + // Seek until it finds '<' for '' + while (*data_end_ptr != '<' && *data_end_ptr != '\0') { ++data_end_ptr; } @@ -1039,7 +1323,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..3fdb71542 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h @@ -17,6 +17,8 @@ #include #include #include +#include +#include #include #include "ifd.h" @@ -98,6 +100,26 @@ class EXPORT_VISIBLE TIFF : public std::enable_shared_from_this friend class IFD; private: + std::vector> companion_parsers_; + std::vector companion_parser_paths_; + std::vector parser_file_hashes_; + std::map, size_t> parser_local_to_global_ifd_; + std::vector> 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 ome_channel_names_; + std::map ome_channel_name_to_index_; + std::map, size_t> ome_plane_to_ifd_; + void _populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr& 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& first_ifd); void _populate_aperio_svs_metadata(uint16_t ifd_count, void* metadata, std::shared_ptr& 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..fb3512833 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 @@ -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..a81706f8b 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2021, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -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_version.h b/cpp/plugins/cucim.kit.cuslide2/src/nvimgcodec_dynlink/nvimgcodec_version.h index 4763bf579..cd2aa2c65 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 @@ -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..41c6a597b --- /dev/null +++ b/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp @@ -0,0 +1,74 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include "cuslide/tiff/ome_xml.h" + +TEST_CASE("OME parser reads pixels and channels", "[ome][parser]") +{ + constexpr const char* kOmeXml = R"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( + + + + + + + urn:uuid:a + + + urn:uuid:b + + + + +)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); +} + From 57aa61598203ec068971feb89ba37ac97547bfdd Mon Sep 17 00:00:00 2001 From: cdinea Date: Mon, 27 Jul 2026 14:22:02 -0700 Subject: [PATCH 02/17] cuslide2: harden nvimgcodec 0.9 OME parsing path Update nvImageCodec dependency constraints to 0.9.x, add parser/decoder compatibility fallbacks for 0.9 runtime behavior, and relax legacy SubfileType assumptions so stacked OME-style TIFF pages load successfully in cuslide2 validation. --- .../all_cuda-129_arch-aarch64.yaml | 6 +- .../all_cuda-129_arch-x86_64.yaml | 6 +- .../all_cuda-133_arch-aarch64.yaml | 6 +- .../all_cuda-133_arch-x86_64.yaml | 6 +- .../recipes/libcucim/conda_build_config.yaml | 2 +- .../src/cuslide/cuslide.cpp | 14 +- .../cuslide/nvimgcodec/nvimgcodec_decoder.cpp | 7 +- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 73 ++- dependencies.yaml | 10 +- python/cucim/pyproject.toml | 2 +- scripts/test_pyramidal_ome_tiff.py | 471 ++++++++++++++++++ 11 files changed, 565 insertions(+), 38 deletions(-) create mode 100644 scripts/test_pyramidal_ome_tiff.py 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/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp index 3f6c0bdf6..b88972365 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp @@ -105,6 +105,9 @@ 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("tiff_type() == cuslide::tiff::TiffType::Hamamatsu || tif->tiff_type() == cuslide::tiff::TiffType::Leica || tif->tiff_type() == cuslide::tiff::TiffType::Ventana || @@ -157,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.")); } } 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 47ad830f6..7cb5a921b 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,6 +12,7 @@ #include #include +#include #include #include #include @@ -339,7 +340,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = sizeof(nvimgcodecCodeStreamView_t); + view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); view.struct_next = nullptr; view.image_idx = ifd_info.index; view.region = region; @@ -636,7 +637,7 @@ std::vector decode_batch_regions_nvimgcodec( nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = sizeof(nvimgcodecCodeStreamView_t); + view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); view.struct_next = nullptr; view.image_idx = ifd_info.index; // Use IFD index for nvImageCodec page selection view.region = nvregion; @@ -935,7 +936,7 @@ BatchDecodeState schedule_batch_decode( nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = sizeof(nvimgcodecCodeStreamView_t); + view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); view.struct_next = nullptr; view.image_idx = ifd_info.index; view.region = nvregion; 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 38b917486..ae586e444 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 @@ -20,6 +20,7 @@ #include "nvimgcodec_tiff_parser.h" #include // for std::transform +#include // for offsetof #include // for std::atexit #include // for strlen #include @@ -283,21 +284,32 @@ 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(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(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 @@ -379,7 +391,8 @@ TiffFileParser::TiffFileParser(const std::string& file_path) nvimgcodecStatus_t status = nvimgcodecCodeStreamCreateFromFile( manager.get_instance(), &main_code_stream_, - file_path.c_str() + file_path.c_str(), + nullptr ); if (status != NVIMGCODEC_STATUS_SUCCESS) @@ -469,6 +482,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) { @@ -478,9 +496,17 @@ bool TiffFileParser::parse_tiff_structure() // Create view for this IFD nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = sizeof(nvimgcodecCodeStreamView_t); + // ABI compatibility: nvImageCodec builds prior to the `limit_images` + // field expect a smaller CodeStreamView payload. + view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); view.struct_next = nullptr; view.image_idx = i; // Note: nvImageCodec uses 'image_idx' not 'image_index' + view.bitstream_offset = 0; + view.limit_images = 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_, @@ -489,6 +515,11 @@ bool TiffFileParser::parse_tiff_structure() if (status != NVIMGCODEC_STATUS_SUCCESS) { + ++substream_failures; + if (first_substream_status == 0) + { + first_substream_status = static_cast(status); + } #ifdef DEBUG fmt::print("❌ Failed to get sub-code stream for IFD {} (status: {})\n", i, static_cast(status)); @@ -511,6 +542,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(status); + } #ifdef DEBUG fmt::print("❌ Failed to get image info for IFD {} (status: {})\n", i, static_cast(status)); @@ -674,6 +710,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; } 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..d8da2806e 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`. diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py new file mode 100644 index 000000000..301575f40 --- /dev/null +++ b/scripts/test_pyramidal_ome_tiff.py @@ -0,0 +1,471 @@ +#!/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) + - optional tile-level caching smoke test +""" + +import json +import sys +import time +import traceback +from pathlib import Path +from typing import Dict, Optional, Tuple + +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() + 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 _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: Optional[str] = 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}, " + f"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: Optional[str] = None, + hubmap_id: Optional[str] = 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.clara import _set_plugin_root + from cucim import CuImage + + _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)") + + 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) + + # 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") + 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 gpu_np.shape != cpu_np.shape: + raise RuntimeError( + f"Shape mismatch at level {level}: CPU={cpu_np.shape}, GPU={gpu_np.shape}" + ) + 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 " + "[--no-cache] [--tsv ] [--hubmap-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: + 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()) + From 9bd44b8a2a8bfa71dcee3403bcfc8be88da70940 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 12:26:27 -0700 Subject: [PATCH 03/17] fix(cuslide2): restore decoder thread-count heuristic max_num_cpu_threads = 0 requests one CPU thread per core from nvImageCodec's default executor. That was harmless when this decoder was CPU-only and used solely for metadata extraction, but it now backs the primary decode path (NVIMGCODEC_DEVICE_CURRENT, all backends enabled), so every worker process allocated a full-core threadpool and multi-worker runs regressed into the thread over-contention this heuristic was originally added to prevent. Restore compute_max_decoder_threads(), which bounds the pool to min(max(1, hardware_concurrency / 4), 8) and stays overridable via CUCIM_MAX_DECODER_THREADS. Also drop the stale nvImageCodec v0.6.0 limitation comments: individual TIFF tags are queried unconditionally through NVIMGCODEC_METADATA_KIND_TIFF_TAG by extract_tiff_tags(), so the claim that they are unavailable contradicted the adjacent code. --- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 93 ++++++++++++++++--- 1 file changed, 79 insertions(+), 14 deletions(-) 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 ae586e444..9518c5842 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 @@ -21,8 +21,9 @@ #include // for std::transform #include // for offsetof -#include // for std::atexit +#include // for std::atexit, std::getenv, std::strtol #include // for strlen +#include // for std::thread::hardware_concurrency #include #ifdef CUCIM_HAS_NVIMGCODEC @@ -144,6 +145,80 @@ static std::string detect_nvimgcodec_extensions_path() } +// ============================================================================ +// Decoder thread count heuristic +// ============================================================================ +// +// Computes a reasonable max_num_cpu_threads for nvImageCodec's internal +// threadpool. On multi-worker deployments the total thread count across +// all workers should not exceed the physical core count to avoid +// context-switch overhead. +// +// Heuristic: min(hardware_concurrency / 4, kHeuristicMaxThreads) +// +// The CUDA stream count is left at nvImageCodec's default because +// empirically it has limited performance impact and over-allocating +// streams can itself cause oversubscription. +// +// Overridable via CUCIM_MAX_DECODER_THREADS environment variable. +// - value > 0 : use that exact thread count +// - value == 0: fall back to nvImageCodec default (= num cpu cores) +// - malformed : warn and fall through to heuristic +// +static constexpr int kHeuristicMaxThreads = 8; + +static int compute_max_decoder_threads() +{ + const char* env_val = std::getenv("CUCIM_MAX_DECODER_THREADS"); + if (env_val) + { + char* end = nullptr; + long val = std::strtol(env_val, &end, 10); + if (end == env_val || *end != '\0') + { + fmt::print(stderr, + "[cuslide2] WARNING: CUCIM_MAX_DECODER_THREADS='{}' is not a " + "valid integer — ignoring and using heuristic\n", env_val); + } + else if (val > 0) + { + #ifdef DEBUG + fmt::print("[cuslide2] CUCIM_MAX_DECODER_THREADS={}\n", val); + #endif + return static_cast(val); + } + else if (val == 0) + { + return 0; // explicit 0 → nvImageCodec default + } + else + { + fmt::print(stderr, + "[cuslide2] WARNING: CUCIM_MAX_DECODER_THREADS={} is negative " + "— ignoring and using heuristic\n", val); + } + } + + unsigned int hw_threads = std::thread::hardware_concurrency(); + if (hw_threads == 0) return 0; // unknown → let nvImageCodec decide + + // Fair share: assume this process is one of potentially several workers. + // Cap at kHeuristicMaxThreads to keep the threadpool bounded even on + // high-core-count machines (e.g. 64 cores / 4 = 16 is still too many). + int fair_share = std::min( + std::max(1, static_cast(hw_threads) / 4), + kHeuristicMaxThreads); + + #ifdef DEBUG + fmt::print("[nvimgcodec] max_decoder_threads heuristic: hw_threads={}, " + "fair_share={}, cap={}\n", + hw_threads, fair_share, kHeuristicMaxThreads); + #endif + + return fair_share; +} + + // nvimgcodec API compatibility // // TIFF-tag retrieval via nvimgcodecDecoderGetMetadata (nvImageCodec >= 0.7.0). @@ -275,7 +350,7 @@ NvImageCodecTiffParserManager::NvImageCodecTiffParserManager() exec_params.struct_next = nullptr; exec_params.device_allocator = nullptr; exec_params.pinned_allocator = nullptr; - exec_params.max_num_cpu_threads = 0; + exec_params.max_num_cpu_threads = compute_max_decoder_threads(); exec_params.executor = nullptr; exec_params.device_id = NVIMGCODEC_DEVICE_CURRENT; // GPU-enabled for decode + metadata exec_params.pre_init = 0; @@ -589,11 +664,8 @@ bool TiffFileParser::parse_tiff_structure() // Extract TIFF metadata using available methods extract_tiff_tags(ifd_info); - // Current limitation (nvImageCodec v0.6.0): - // - codec_name returns "tiff" (container format) not compression type - // - Individual TIFF tags not exposed through metadata API - // - Only vendor-specific metadata blobs available (MED_APERIO, MED_PHILIPS, etc.) - // + // 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") { @@ -931,13 +1003,6 @@ void TiffFileParser::extract_ifd_metadata(IfdInfo& ifd_info) } } - // WORKAROUND for nvImageCodec 0.6.0: Philips TIFF metadata limitation - // ======================================================================== - // nvImageCodec 0.6.0 does NOT expose: - // 1. Individual TIFF tags (SOFTWARE, ImageDescription, etc.) - // 2. Philips format detection for some files - // - } const IfdInfo& TiffFileParser::get_ifd(uint32_t index) const From 091263674b670cf69b79144df7dd4a6e1869e286 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 12:46:33 -0700 Subject: [PATCH 04/17] fix(cuslide2): require nvImageCodec 0.9 in dependency discovery --- cpp/cmake/deps/nvimgcodec.cmake | 53 ++++++++++++++++++- .../cmake/deps/nvimgcodec.cmake | 53 ++++++++++++++++++- 2 files changed, 102 insertions(+), 4 deletions(-) diff --git a/cpp/cmake/deps/nvimgcodec.cmake b/cpp/cmake/deps/nvimgcodec.cmake index 97a44008c..a7fcc5053 100644 --- a/cpp/cmake/deps/nvimgcodec.cmake +++ b/cpp/cmake/deps/nvimgcodec.cmake @@ -6,9 +6,11 @@ # 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.8.0" "libnvimgcodec.so.0" "libnvimgcodec.so") @@ -22,14 +24,60 @@ if (NOT TARGET deps::nvimgcodec) set(${out_var} "" PARENT_SCOPE) endfunction() + # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above + # resolve to whatever major-0 build is installed, so an unsupported 0.8 would + # otherwise be picked up silently. The 0.9 API is not backward compatible + # (limit_images was removed) 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 "") @@ -104,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/cmake/deps/nvimgcodec.cmake b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake index 4759fa8dd..ac054543b 100644 --- a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake +++ b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake @@ -9,9 +9,11 @@ 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.8.0" "libnvimgcodec.so.0" "libnvimgcodec.so") @@ -25,13 +27,59 @@ if (NOT TARGET deps::nvimgcodec) set(${out_var} "" PARENT_SCOPE) endfunction() + # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above + # resolve to whatever major-0 build is installed, so an unsupported 0.8 would + # otherwise be picked up silently. The 0.9 API is not backward compatible + # (limit_images was removed) 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}) @@ -127,6 +175,7 @@ if (NOT TARGET deps::nvimgcodec) # 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}" From ce00b564c98f5d209cfd0a1679ed81c608f3039a Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 14:17:59 -0700 Subject: [PATCH 05/17] fix(cuslide2): sync vendored nvImageCodec headers with 0.9.0 Replace the checked-in nvimgcodec.h / nvimgcodec_version.h copies with the real 0.9.0 package headers. The previous vendored files claimed 0.9 but still described the pre-0.9 CodeStreamView layout that included limit_images, so call sites could compile against an API the linked library no longer provides. With the real headers in place, restore struct_size to sizeof(nvimgcodecCodeStreamView_t) at all four CodeStreamView sites and drop the leftover limit_images assignment and offsetof ABI shrink. Also bump SPDX copyright years on files touched by this PR. --- cpp/cmake/deps/nvimgcodec.cmake | 2 +- .../cmake/deps/nvimgcodec.cmake | 2 +- .../cuslide/nvimgcodec/nvimgcodec_decoder.cpp | 6 +- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 7 +- .../src/cuslide/tiff/types.h | 2 +- .../src/nvimgcodec_dynlink/nvimgcodec.h | 107 +++++++++++++----- .../nvimgcodec_dynlink/nvimgcodec_version.h | 14 ++- 7 files changed, 99 insertions(+), 41 deletions(-) diff --git a/cpp/cmake/deps/nvimgcodec.cmake b/cpp/cmake/deps/nvimgcodec.cmake index a7fcc5053..72c8e453d 100644 --- a/cpp/cmake/deps/nvimgcodec.cmake +++ b/cpp/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 # diff --git a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake index ac054543b..6e6c207ad 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 # 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 7cb5a921b..273e9b85c 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 @@ -340,7 +340,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); + view.struct_size = sizeof(nvimgcodecCodeStreamView_t); view.struct_next = nullptr; view.image_idx = ifd_info.index; view.region = region; @@ -637,7 +637,7 @@ std::vector decode_batch_regions_nvimgcodec( nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); + view.struct_size = sizeof(nvimgcodecCodeStreamView_t); view.struct_next = nullptr; view.image_idx = ifd_info.index; // Use IFD index for nvImageCodec page selection view.region = nvregion; @@ -936,7 +936,7 @@ BatchDecodeState schedule_batch_decode( nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); + view.struct_size = sizeof(nvimgcodecCodeStreamView_t); view.struct_next = nullptr; view.image_idx = ifd_info.index; view.region = nvregion; 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 9518c5842..00b435cd9 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 @@ -20,7 +20,7 @@ #include "nvimgcodec_tiff_parser.h" #include // for std::transform -#include // for offsetof +#include // for size_t #include // for std::atexit, std::getenv, std::strtol #include // for strlen #include // for std::thread::hardware_concurrency @@ -571,13 +571,10 @@ bool TiffFileParser::parse_tiff_structure() // Create view for this IFD nvimgcodecCodeStreamView_t view{}; view.struct_type = NVIMGCODEC_STRUCTURE_TYPE_CODE_STREAM_VIEW; - // ABI compatibility: nvImageCodec builds prior to the `limit_images` - // field expect a smaller CodeStreamView payload. - view.struct_size = offsetof(nvimgcodecCodeStreamView_t, limit_images); + 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.limit_images = 0; view.region.struct_type = NVIMGCODEC_STRUCTURE_TYPE_REGION; view.region.struct_size = sizeof(nvimgcodecRegion_t); view.region.struct_next = nullptr; 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 a81706f8b..08f8f6b00 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/types.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2021, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ 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..9e84ca15c 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,19 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * */ /** @@ -385,6 +398,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 +459,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 +511,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 +593,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 +612,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 +622,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 +630,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 +737,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 +812,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 +955,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 +1876,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 +1895,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 +1925,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 +1969,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 +1991,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 - * ":=". For example "nvjpeg:fancy_upsampling=1" + * @param options [in] Optional space-separated list of parameters. Use ":=" 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 +2061,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 - * ":=." + * @param options [in] Optional space-separated list of parameters. Use ":=" 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 cd2aa2c65..174684a8f 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,6 +1,18 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. */ // clang-format off #ifndef NVIMGCODEC_VERSION_H__ From d7c002387d54c338515f7564b858c1ac1849f116 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 14:33:06 -0700 Subject: [PATCH 06/17] fix(cuslide2): widen row stride arithmetic to 64-bit width, num_channels and bytes_per_sample are all uint32_t, so the row stride product was evaluated in 32-bit and could wrap before being widened on assignment to size_t. Cast the width operand to size_t so the multiply happens in 64-bit, at all three decode sites. --- .../src/cuslide/nvimgcodec/nvimgcodec_decoder.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) 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 273e9b85c..57dccdeeb 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 @@ -395,7 +395,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, output_image_info.buffer_kind = buffer_kind; uint32_t num_channels = decode_spec.num_channels; - size_t row_stride = width * num_channels * decode_spec.bytes_per_sample; + size_t row_stride = static_cast(width) * num_channels * decode_spec.bytes_per_sample; size_t buffer_size = row_stride * height; output_image_info.plane_info[0].height = height; @@ -680,7 +680,7 @@ std::vector decode_batch_regions_nvimgcodec( const auto& region = regions[i]; const DecodePixelSpec decode_spec = resolve_decode_pixel_spec(ifd_info); uint32_t num_channels = decode_spec.num_channels; - size_t row_stride = region.width * num_channels * decode_spec.bytes_per_sample; + size_t row_stride = static_cast(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)) @@ -983,7 +983,7 @@ BatchDecodeState schedule_batch_decode( const auto& region = regions[i]; const DecodePixelSpec decode_spec = resolve_decode_pixel_spec(ifd_info); uint32_t num_channels = decode_spec.num_channels; - size_t row_stride = region.width * num_channels * decode_spec.bytes_per_sample; + size_t row_stride = static_cast(region.width) * num_channels * decode_spec.bytes_per_sample; size_t buffer_size = row_stride * region.height; void* buffer_ptr = nullptr; From 1328a76412866a6f727043aa03cc0c782576cca4 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 14:43:22 -0700 Subject: [PATCH 07/17] fix(cuslide2): restore zero-downsample guard in Philips metadata --- .../cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) 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 17ca5367e..c487945d3 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp @@ -1124,9 +1124,14 @@ 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 - ifd->width_ = width_l0 / downsample; - ifd->height_ = height_l0 / downsample; + // 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; + ifd->height_ = height_l0 / downsample; + } ++spacing_index; } else From bafbdb5f5746317dcdc0d0a60641b965aae0107b Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 14:51:00 -0700 Subject: [PATCH 08/17] fix(cuslide2): bound associated-image scan to image description The unconditional `+ 34` skip past the PIM_DP_IMAGE_DATA node could land beyond the terminator, leaving the following scan loops unanchored since both stop conditions are byte values inside the buffer. Restore the length check on node_offset and the explicit desc_end sentinel so every loop is bounded by the string, and drop the skip and the const_cast. --- .../src/cuslide/tiff/tiff.cpp | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) 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 c487945d3..4677def2d 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp @@ -1170,27 +1170,29 @@ void TIFF::_populate_philips_tiff_metadata(uint16_t ifd_count, void* metadata, s { auto node_offset = associated_node.node().offset_debug(); - if (node_offset >= 0) + size_t image_desc_len = first_ifd->image_description().size(); + if (node_offset >= 0 && static_cast(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) // // + const char* data_ptr = image_desc_cstr + node_offset; + const char* desc_end = image_desc_cstr + image_desc_len; - // 34 is from `Attribute Name="PIM_DP_IMAGE_DATA"` - char* data_ptr = const_cast(image_desc_cstr) + node_offset + 34; - uint32_t data_len = 0; - while (*data_ptr != '>' && *data_ptr != '\0') + // Find '>' that closes the opening tag + while (data_ptr < desc_end && *data_ptr != '>') { ++data_ptr; } - if (*data_ptr != '\0') + uint32_t data_len = 0; + if (data_ptr < desc_end && *data_ptr != '\0') { ++data_ptr; // start of base64-encoded data - char* data_end_ptr = data_ptr; + const char* data_end_ptr = data_ptr; // Seek until it finds '<' for '' - while (*data_end_ptr != '<' && *data_end_ptr != '\0') + while (data_end_ptr < desc_end && *data_end_ptr != '<' && *data_end_ptr != '\0') { ++data_end_ptr; } From eac5aeec3ac69e9d6c5b6999cbcfe949f72dfec8 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 14:58:59 -0700 Subject: [PATCH 09/17] fix(cuslide2): own channel names as strings while building Replace the allocate-and-view lambda with an owning vector of strings. ImageMetadata::channel_names still takes string_views and may retain pointers into the metadata buffer, so copy into the resource once at the end when constructing the view vector the API requires. --- .../src/cuslide/cuslide.cpp | 38 ++++++++++--------- 1 file changed, 20 insertions(+), 18 deletions(-) diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp index b88972365..d6c1eb429 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp @@ -246,49 +246,51 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo shape[2] = n_ch; } - std::pmr::vector channel_names(&resource); - channel_names.reserve(n_ch); - auto add_owned_channel_name = [&resource, &channel_names](const std::string& name) { - void* raw = resource.allocate(name.size() + 1, alignof(char)); - auto* buf = static_cast(raw); - memcpy(buf, name.c_str(), name.size() + 1); - channel_names.emplace_back(buf, name.size()); - }; + // 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 channel_name_storage; + channel_name_storage.reserve(n_ch); if (!ome_channel_names.empty()) { for (size_t i = 0; i < static_cast(n_ch); ++i) { if (i < ome_channel_names.size()) { - add_owned_channel_name(ome_channel_names[i]); + channel_name_storage.push_back(ome_channel_names[i]); } else { - add_owned_channel_name(fmt::format("C{}", i)); + channel_name_storage.push_back(fmt::format("C{}", i)); } } } else if (n_ch == 3) { - 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_name_storage = { "R", "G", "B" }; } else if (n_ch == 4) { - 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" }); + channel_name_storage = { "R", "G", "B", "A" }; } else { for (uint8_t i = 0; i < n_ch; ++i) { - add_owned_channel_name(fmt::format("C{}", i)); + channel_name_storage.push_back(fmt::format("C{}", i)); } } + std::pmr::vector channel_names(&resource); + channel_names.reserve(channel_name_storage.size()); + for (const auto& name : channel_name_storage) + { + char* buf = static_cast(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 static constexpr std::string_view empty_unit{""}; static constexpr std::string_view micrometer_unit{"micrometer"}; From 84fe4ed196bda0259990024c9156f24d32e0b495 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 15:16:32 -0700 Subject: [PATCH 10/17] test(cuslide2): strengthen pyramidal OME-TIFF validation Add pixel-level CPU/GPU equality checks on multi-level reads, and assert pyramid dimensions shrink and match the reported downsample factors so parsing mistakes cannot pass on shape alone. --- scripts/test_pyramidal_ome_tiff.py | 117 +++++++++++++++++++++++------ 1 file changed, 95 insertions(+), 22 deletions(-) diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py index 301575f40..d895c8b6d 100644 --- a/scripts/test_pyramidal_ome_tiff.py +++ b/scripts/test_pyramidal_ome_tiff.py @@ -17,7 +17,6 @@ import time import traceback from pathlib import Path -from typing import Dict, Optional, Tuple import numpy as np from test_common import setup_environment, test_tile_level_caching @@ -31,20 +30,66 @@ def _to_numpy(arr): return np.asarray(arr) -def _max_read_size(level_dims: Tuple[int, int], cap: int = 512): +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.""" + 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]}" + ) + + 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 ds <= prev_ds: + raise RuntimeError( + f"Level {level} downsample {ds} should be > " + f"level {level - 1} downsample {prev_ds}" + ) + + # Allow ±1 for integer rounding of base / downsample. + expected_w = max(1, int(round(base_w / ds))) + expected_h = max(1, int(round(base_h / ds))) + if abs(w - expected_w) > 1 or abs(h - expected_h) > 1: + raise RuntimeError( + f"Level {level}: dims {w}x{h} != base/ds ~{expected_w}x{expected_h} " + f"(ds={ds})" + ) + + print(" ✅ Pyramid dims decrease and match reported downsamples") + + 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( + " https://portal.hubmapconsortium.org/browse/dataset/e4263715a087881e46ea4d11f49139aa" + ) print("Skin sample (19 markers):") - print(" https://portal.hubmapconsortium.org/browse/dataset/1b8121539ff16f53681de6108069be24") + print( + " https://portal.hubmapconsortium.org/browse/dataset/1b8121539ff16f53681de6108069be24" + ) -def _extract_ome_meta(metadata: Dict): +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 @@ -54,7 +99,7 @@ def _extract_ome_meta(metadata: Dict): return ome, size_c, size_z, size_t, channel_names -def _load_cell_dive_tsv(tsv_path: str, hubmap_id: Optional[str] = None): +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(): @@ -114,7 +159,9 @@ def _as_int_or_none(v): return None -def _validate_tsv_mapping(tsv_hubmap_id: str, tsv_meta: Dict, ome: Dict, size_c: int, channel_names): +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) @@ -129,7 +176,9 @@ def _validate_tsv_mapping(tsv_hubmap_id: str, tsv_meta: Dict, ome: Dict, size_c: 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}") + 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}" @@ -139,7 +188,9 @@ def _validate_tsv_mapping(tsv_hubmap_id: str, tsv_meta: Dict, ome: Dict, size_c: 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}") + 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}" @@ -207,7 +258,9 @@ def _validate_channel_selection(img, level_dims, level, c_index, z_index=0, t_in # 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_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( @@ -221,7 +274,9 @@ def _validate_channel_selection(img, level_dims, level, c_index, z_index=0, t_in 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)))) + 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}" ) @@ -254,8 +309,7 @@ def _validate_batch_decode(img, level_dims): batch_size = min(8, len(locations)) print( - f" Locations={len(locations)}, tile={tile_w}x{tile_h}, " - f"batch_size={batch_size}" + f" Locations={len(locations)}, tile={tile_w}x{tile_h}, batch_size={batch_size}" ) # CPU ground truth @@ -292,7 +346,9 @@ def _validate_batch_decode(img, level_dims): 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)}") + 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)): @@ -301,9 +357,13 @@ def _validate_batch_decode(img, level_dims): 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]}") + print( + f" ❌ mismatch at tile idx={idx}, location={locations[idx]}" + ) if mismatch: - raise RuntimeError(f"Batch decode mismatch count: {mismatch}/{len(cpu_tiles)}") + 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}") @@ -313,8 +373,8 @@ def test_pyramidal_ome_tiff( file_path, plugin_lib, run_cache=True, - tsv_path: Optional[str] = None, - hubmap_id: Optional[str] = None, + tsv_path: str | None = None, + hubmap_id: str | None = None, ): print("=" * 70) print("🔬 Testing pyramidal OME-TIFF with cuslide2") @@ -324,8 +384,8 @@ def test_pyramidal_ome_tiff( if not Path(file_path).exists(): raise FileNotFoundError(f"OME-TIFF file not found: {file_path}") - from cucim.clara import _set_plugin_root from cucim import CuImage + from cucim.clara import _set_plugin_root _set_plugin_root(str(plugin_lib)) print(f"✅ Plugin root set: {plugin_lib}") @@ -368,6 +428,8 @@ def test_pyramidal_ome_tiff( 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])) @@ -384,11 +446,21 @@ def test_pyramidal_ome_tiff( 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") + print( + f" GPU level {level}: {gpu_np.shape}, {gpu_np.dtype}, {gpu_time:.4f}s" + ) 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}") @@ -402,7 +474,9 @@ def test_pyramidal_ome_tiff( 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) + _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)") @@ -468,4 +542,3 @@ def main(): if __name__ == "__main__": sys.exit(main()) - From aef0ca15476ad240a44f46ebb5d4cb222c370565 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 15:23:42 -0700 Subject: [PATCH 11/17] style(cuslide2): satisfy pre-commit copyright and codespell checks Bump SPDX copyright years, normalize trailing newlines, convert the vendored nvImageCodec headers to the SPDX-only form the copyright hook expects, and ignore upstream typos (minize/shrinked) in those headers. --- .../cucim.kit.cuslide2/src/cuslide/tiff/ifd.h | 2 +- .../src/cuslide/tiff/ome_xml.cpp | 1 - .../cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h | 1 - .../cucim.kit.cuslide2/src/cuslide/tiff/tiff.h | 2 +- .../src/cuslide/tiff/tiff_constants.h | 2 +- .../src/nvimgcodec_dynlink/nvimgcodec.h | 14 +------------- .../src/nvimgcodec_dynlink/nvimgcodec_version.h | 14 +------------- .../cucim.kit.cuslide2/tests/test_ome_xml.cpp | 1 - python/cucim/pyproject.toml | 3 ++- 9 files changed, 7 insertions(+), 33 deletions(-) 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 c11f487b6..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 */ 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 index 6e5d1ded1..2e3e52bf4 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.cpp @@ -168,4 +168,3 @@ bool parse(const std::string& xml_text, Model* out_model, std::string* out_error } } // 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 index 290737173..00de44bde 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ome_xml.h @@ -61,4 +61,3 @@ struct Model 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.h b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.h index 3fdb71542..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 */ 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 fb3512833..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 */ 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 9e84ca15c..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,19 +1,7 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - * */ /** 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 174684a8f..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,18 +1,6 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. */ // clang-format off #ifndef NVIMGCODEC_VERSION_H__ diff --git a/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp b/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp index 41c6a597b..3653d294e 100644 --- a/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/tests/test_ome_xml.cpp @@ -71,4 +71,3 @@ TEST_CASE("OME parser reads UUID companion references", "[ome][parser]") REQUIRE(model.pixels.tiff_data[1].ifd == 3); REQUIRE(model.pixels.tiff_data[1].first_c == 1); } - diff --git a/python/cucim/pyproject.toml b/python/cucim/pyproject.toml index d8da2806e..0844f33a2 100644 --- a/python/cucim/pyproject.toml +++ b/python/cucim/pyproject.toml @@ -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 From 15752d3e3a45ebc16ff336cfa32b7db91e3c8342 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 15:36:36 -0700 Subject: [PATCH 12/17] test(cuslide2): cover uint16 pyramidal OME-TIFF decode Add a synthetic uint16 OME-TIFF decode path that asserts metadata and region dtypes stay 16-bit, preserves values above 255, and matches source pixels across pyramid levels. Also enforce uint16 on multi-level CPU/GPU checks when the input file reports that dtype. --- scripts/test_pyramidal_ome_tiff.py | 133 +++++++++++++++++++++++++++++ 1 file changed, 133 insertions(+) diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py index d895c8b6d..ccd55262d 100644 --- a/scripts/test_pyramidal_ome_tiff.py +++ b/scripts/test_pyramidal_ome_tiff.py @@ -9,11 +9,13 @@ - 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 @@ -76,6 +78,122 @@ def _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count) print(" ✅ Pyramid dims decrease and match reported downsamples") +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""" + + + + + + + +""" + + 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) @@ -409,6 +527,11 @@ def test_pyramidal_ome_tiff( 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): @@ -441,6 +564,10 @@ def test_pyramidal_ome_tiff( 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") @@ -449,6 +576,10 @@ def test_pyramidal_ome_tiff( 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}" @@ -526,6 +657,8 @@ def main(): 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, From aeca33b5ea643257248a3d55839ebfbadc1ae24d Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 16:52:56 -0700 Subject: [PATCH 13/17] fix(cuslide2): correct pyramid geometry assertions in OME-TIFF test The geometry check assumed a single-plane WSI pyramid: strictly increasing downsample per level, and level dims recoverable as base / downsample. Both assumptions are false in practice. cuslide2 reports downsample as the mean of the per-axis ratios ((w0/wi) + (h0/hi)) / 2, and levels may use different X and Y ratios, so dims cannot be derived from the scalar. Multi-plane OME files also expose one IFD per C/Z/T plane, so many levels legitimately share a resolution. Validate what cuslide2 actually guarantees instead: dims are non-increasing, downsample strictly increases when the resolution changes and is identical when a resolution repeats, and the reported scalar matches the mean axis ratio. Repeated resolutions now warn that the file looks multi-plane rather than failing the run. --- scripts/test_pyramidal_ome_tiff.py | 50 ++++++++++++++++++++++-------- 1 file changed, 37 insertions(+), 13 deletions(-) diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py index ccd55262d..b0afa5d81 100644 --- a/scripts/test_pyramidal_ome_tiff.py +++ b/scripts/test_pyramidal_ome_tiff.py @@ -37,7 +37,14 @@ def _max_read_size(level_dims: tuple[int, int], cap: int = 512): def _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count): - """Assert pyramid levels shrink and match reported downsample factors.""" + """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") @@ -48,6 +55,7 @@ def _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count) 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]) @@ -60,22 +68,38 @@ def _validate_pyramid_geometry(level_dimensions, level_downsamples, level_count) f"Level {level} dims {w}x{h} are not smaller than " f"level {level - 1} {prev_w}x{prev_h}" ) - if ds <= prev_ds: - raise RuntimeError( - f"Level {level} downsample {ds} should be > " - f"level {level - 1} downsample {prev_ds}" - ) - # Allow ±1 for integer rounding of base / downsample. - expected_w = max(1, int(round(base_w / ds))) - expected_h = max(1, int(round(base_h / ds))) - if abs(w - expected_w) > 1 or abs(h - expected_h) > 1: + 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}: dims {w}x{h} != base/ds ~{expected_w}x{expected_h} " - f"(ds={ds})" + 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(" ✅ Pyramid dims decrease and match reported downsamples") + 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: From 1563e4c949dcf6f02b75e873c6dccd94a68967e9 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 17:58:16 -0700 Subject: [PATCH 14/17] fix(cuslide2): report TIFF resolution and preserve OME pyramid levels Resolution was hardcoded to 1.0 with no unit, so every file reported meaningless pixel spacing. Query tags 282/283/296, decode RATIONAL values, and populate the IFD fields from them. OME TiffData entries only reference full-resolution planes, so mapping planes alone collapsed the level list to level 0 and made pyramid sub-resolutions unreachable through read_region(). Backfill the levels the plane index does not cover. Also surface two previously silent failures: tag retrieval returning nothing (which disables OME/vendor detection) and unusable OME XML. Fix the test helper so device-resident results route through CuPy instead of degrading every GPU comparison to a skipped check. --- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 58 +++++++++++++++++++ .../nvimgcodec/nvimgcodec_tiff_parser.h | 3 + .../src/cuslide/tiff/ifd.cpp | 54 +++++++++++++++-- .../src/cuslide/tiff/tiff.cpp | 32 +++++++++- scripts/test_pyramidal_ome_tiff.py | 7 +++ 5 files changed, 148 insertions(+), 6 deletions(-) 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 00b435cd9..d101e914f 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 @@ -272,6 +272,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 +static bool extract_rational_value(const std::vector& 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(buffer.data()); + auto as_double = [](T numerator, T denominator) -> double { + return (denominator == 0) ? 0.0 : static_cast(numerator) / static_cast(denominator); + }; + + if (pair_count == 1) + { + out_value = as_double(vals[0], vals[1]); + } + else + { + std::vector 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 static bool extract_tag_value(const std::vector& buffer, int value_count, TiffTagValue& out_value) @@ -1063,6 +1098,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"}, @@ -1148,6 +1186,12 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info) case NVIMGCODEC_METADATA_VALUE_TYPE_SLONG: extract_tag_value(buffer, metadata.value_count, tag_value); break; + case NVIMGCODEC_METADATA_VALUE_TYPE_RATIONAL: + extract_rational_value(buffer, tag_value); + break; + case NVIMGCODEC_METADATA_VALUE_TYPE_SRATIONAL: + extract_rational_value(buffer, tag_value); + break; case NVIMGCODEC_METADATA_VALUE_TYPE_LONG8: case NVIMGCODEC_METADATA_VALUE_TYPE_IFD8: extract_tag_value(buffer, metadata.value_count, tag_value); @@ -1181,6 +1225,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()) 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..341b5245e 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 @@ -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; }; /** 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 6d77d15c6..c09e0e0df 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/ifd.cpp @@ -212,6 +212,55 @@ IFD::IFD(TIFF* tiff, // 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(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 @@ -223,11 +272,6 @@ IFD::IFD(TIFF* tiff, } 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_ = (file_hash != 0) ? (file_hash ^ cucim::codec::splitmix64(local_ifd_idx)) : (tiff->file_handle_shared_.get()->hash_value ^ cucim::codec::splitmix64(index)); 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 4677def2d..60ddb42f7 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp @@ -814,6 +814,11 @@ void TIFF::_populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std:: 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; } @@ -914,9 +919,11 @@ void TIFF::_populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std:: } ome_plane_to_ifd_.clear(); + + // Distinct IFD dimensions, largest first: index into this list is the resolution level. + std::vector> level_dims; if (!pending_planes.empty()) { - std::vector> level_dims; level_dims.reserve(ifds_.size()); for (const auto& ifd : ifds_) { @@ -969,8 +976,31 @@ void TIFF::_populate_ome_tiff_metadata(uint16_t ifd_count, void* metadata, std:: } } } + // 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(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) { diff --git a/scripts/test_pyramidal_ome_tiff.py b/scripts/test_pyramidal_ome_tiff.py index b0afa5d81..fe0daf37b 100644 --- a/scripts/test_pyramidal_ome_tiff.py +++ b/scripts/test_pyramidal_ome_tiff.py @@ -29,6 +29,13 @@ def _to_numpy(arr): 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) From 0914c4a416b2f6321649fc0e2b548eb33b5508e4 Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 18:33:00 -0700 Subject: [PATCH 15/17] fix(cuslide2): stop aborting in free() during process exit Reading a multi-plane OME-TIFF aborted with "free(): invalid pointer" after the run completed, roughly one exit in five. The parser manager destroyed its nvImageCodec decoder at process exit. That cannot be ordered against the CUDA driver's own teardown, so nvimgcodecDecoderDestroy intermittently freed pointers the driver had already reclaimed. Neither available ordering was safe: the singleton's static destructor runs before the atexit handler registered in the constructor, because that handler is registered first and exit handlers run in reverse order, and moving teardown to the handler still aborted. Leak the singleton and skip exit-time teardown. These are process- lifetime handles, so the driver and OS reclaim them. shutdown() stays available for deterministic release while CUDA is up. --- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 25 +++++++++++-------- .../nvimgcodec/nvimgcodec_tiff_parser.h | 12 +++++++-- 2 files changed, 24 insertions(+), 13 deletions(-) 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 d101e914f..724b2aaa7 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 @@ -424,14 +424,16 @@ NvImageCodecTiffParserManager::NvImageCodecTiffParserManager() 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) { @@ -475,9 +477,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(); } 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 341b5245e..3aede840e 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 @@ -331,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; } /** From 729d27a0d3848edbacccdf54f49c852b2c3c32cb Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 18:33:21 -0700 Subject: [PATCH 16/17] fix(cuslide2): return background for fully out-of-bounds regions read_region() threw "Failed to decode IFD[0] with nvImageCodec" for a region with no overlap with the image, where the documented behavior and the cuslide behavior is an all-zero raster of the requested size. This failed test_tiff_stripe_outside for all three compression variants. nvImageCodec fills out-of-bounds pixels when a region partially overlaps, which is why boundary reads worked, but it rejects a region that does not intersect the image at all. Detect that case and return a zeroed raster without consulting the decoder, for host and device outputs alike. --- .../src/cuslide/tiff/ifd.cpp | 54 +++++++++++++++++++ 1 file changed, 54 insertions(+) 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 c09e0e0df..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 #include #include +#include #include #include @@ -541,6 +542,59 @@ 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(width_) || sy >= static_cast(height_)) + { + if (!output_buffer) + { + if (out_device.type() == cucim::io::DeviceType::kCUDA) + { + if (cudaMalloc(reinterpret_cast(&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(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(out_device.type()), out_device.index() }; + out_image_data->container.dtype = DLDataType{ kDLUInt, static_cast(bits_per_sample_), 1 }; + out_image_data->container.ndim = 3; + out_image_data->container.shape = static_cast(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) From e011f4429f5d0ebbaed344e4e8764d76989d4a6a Mon Sep 17 00:00:00 2001 From: cdinea Date: Wed, 29 Jul 2026 22:33:42 -0700 Subject: [PATCH 17/17] docs(cuslide2): drop references to unsupported nvImageCodec versions Packaging now pins nvImageCodec to >=0.9.0,<0.10.0, so comments that gate behavior on 0.6.0, 0.7.0, or 0.8.0 no longer describe anything reachable. Reword them to say what the code does rather than which release introduced it, keeping the vendor-detection fallbacks: those still matter for files that carry no SOFTWARE tag or ImageDescription, independently of the codec version. Also remove CUSLIDE2_NVIMGCODEC_HAS_TIFF_TAG_METADATA. It was defined unconditionally to 1 with no alternate branch, so the surrounding preprocessor conditionals were always taken. --- cpp/cmake/deps/nvimgcodec.cmake | 8 +++---- .../cmake/deps/nvimgcodec.cmake | 8 +++---- .../src/cuslide/cuslide.cpp | 8 +++---- .../cuslide/nvimgcodec/nvimgcodec_decoder.cpp | 4 ++-- .../cuslide/nvimgcodec/nvimgcodec_decoder.h | 2 +- .../nvimgcodec/nvimgcodec_tiff_parser.cpp | 21 ++++--------------- .../nvimgcodec/nvimgcodec_tiff_parser.h | 4 ++-- .../src/cuslide/tiff/tiff.cpp | 13 ++++++------ 8 files changed, 27 insertions(+), 41 deletions(-) diff --git a/cpp/cmake/deps/nvimgcodec.cmake b/cpp/cmake/deps/nvimgcodec.cmake index 72c8e453d..9176fb14d 100644 --- a/cpp/cmake/deps/nvimgcodec.cmake +++ b/cpp/cmake/deps/nvimgcodec.cmake @@ -25,10 +25,10 @@ if (NOT TARGET deps::nvimgcodec) endfunction() # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above - # resolve to whatever major-0 build is installed, so an unsupported 0.8 would - # otherwise be picked up silently. The 0.9 API is not backward compatible - # (limit_images was removed) and there are no compile-time version guards, so - # fail at configure time instead of at runtime. + # 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}") diff --git a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake index 6e6c207ad..ac8943ea5 100644 --- a/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake +++ b/cpp/plugins/cucim.kit.cuslide2/cmake/deps/nvimgcodec.cmake @@ -28,10 +28,10 @@ if (NOT TARGET deps::nvimgcodec) endfunction() # The unversioned "libnvimgcodec.so.0" / "libnvimgcodec.so" candidates above - # resolve to whatever major-0 build is installed, so an unsupported 0.8 would - # otherwise be picked up silently. The 0.9 API is not backward compatible - # (limit_images was removed) and there are no compile-time version guards, so - # fail at configure time instead of at runtime. + # 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}") diff --git a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp index d6c1eb429..d3c73e74b 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/cuslide.cpp @@ -99,7 +99,7 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo size_t level_count = tif->level_count(); // Detect if this is an Aperio SVS file - // Try ImageDescription first (works with nvImageCodec 0.7.0+) + // Try ImageDescription first bool is_aperio_svs = (tif->ifd(0)->image_description().rfind("Aperio", 0) == 0); // Detect if this is a Philips TIFF file @@ -109,9 +109,9 @@ static bool CUCIM_ABI parser_parse(CuCIMFileHandle_ptr handle_ptr, cucim::io::fo bool looks_like_ome = (first_desc.find("= 3 && level_count >= 3) { // Check if IFDs form a pyramid structure (decreasing sizes) 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 57dccdeeb..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 @@ -403,7 +403,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, 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 = decode_spec.sample_type; - // Note: buffer_size removed in nvImageCodec v0.7.0 - size is inferred from plane_info + // Buffer size is inferred from plane_info rather than set explicitly. output_image_info.cuda_stream = cuda_stream; // Step 4: Provide output buffer @@ -530,7 +530,7 @@ bool decode_ifd_region_nvimgcodec(const IfdInfo& ifd_info, } // ============================================================================ -// Batch ROI Decoding (nvImageCodec v0.7.0+) +// Batch ROI Decoding // ============================================================================ std::vector decode_batch_regions_nvimgcodec( 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 724b2aaa7..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,7 +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 +// - Direct TIFF tag queries via NVIMGCODEC_METADATA_KIND_TIFF_TAG // (COMPRESSION, SUBFILETYPE, IMAGEDESCRIPTION, JPEGTABLES, ...) // // ============================================================================ @@ -219,14 +219,6 @@ static int compute_max_decoder_threads() } -// nvimgcodec API compatibility -// -// TIFF-tag retrieval via nvimgcodecDecoderGetMetadata (nvImageCodec >= 0.7.0). -// 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) { @@ -1085,11 +1077,10 @@ void TiffFileParser::extract_tiff_tags(IfdInfo& ifd_info) } // ======================================================================== - // nvImageCodec 0.7.0+: Direct TIFF Tag Retrieval by ID + // 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> tiff_tags_to_query = { {254, "SUBFILETYPE"}, {256, "IMAGEWIDTH"}, @@ -1254,10 +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 available - // (either nvImageCodec < 0.7.0 or tag 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 @@ -1304,13 +1293,11 @@ std::vector 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 3aede840e..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 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 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 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 60ddb42f7..6797c6b96 100644 --- a/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp +++ b/cpp/plugins/cucim.kit.cuslide2/src/cuslide/tiff/tiff.cpp @@ -642,7 +642,7 @@ void TIFF::resolve_vendor_format() is_aperio = true; } - // Method 2: Check metadata_blobs for Aperio (kind=5 in v0.8.0) + // Method 2: Check metadata_blobs for Aperio if (!is_aperio && nvimgcodec_parser_) { const auto& metadata_blobs = nvimgcodec_parser_->get_metadata_blobs(0); @@ -661,13 +661,12 @@ void TIFF::resolve_vendor_format() } } - // Detect Philips TIFF - // NOTE: nvImageCodec 0.6.0 doesn't expose individual TIFF tags (like SOFTWARE) - // Workaround: Check for Philips XML in ImageDescription or use 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 (available in nvImageCodec 0.7.0+) + // Method 1: Check SOFTWARE tag std::string_view prefix("Philips"); auto res = std::mismatch(prefix.begin(), prefix.end(), software.begin()); if (res.first == prefix.end()) @@ -676,7 +675,7 @@ void TIFF::resolve_vendor_format() } // Method 2: Check for Philips XML structure in ImageDescription - // (Workaround for nvImageCodec 0.6.0 where SOFTWARE tag is not available) + // (used when the file has no SOFTWARE tag) if (!is_philips) { auto& image_desc = first_ifd->image_description(); @@ -687,7 +686,7 @@ void TIFF::resolve_vendor_format() } } - // Method 3: Check metadata_blobs for Philips (kind=6 in v0.8.0) + // Method 3: Check metadata_blobs for Philips if (!is_philips && nvimgcodec_parser_) { const auto& metadata_blobs = nvimgcodec_parser_->get_metadata_blobs(0);