Fixed image enhance module.

src/image_process_v5/CMakeLists.txt

@@ -1,8 +1,12 @@
 target_sources(${PROJECT_NAME} PRIVATE
+        image_enhance.cpp
         image_viewer.cpp
         image_process.cpp
         osg_helper.cpp
         sp_image.cpp
         process_python/fast_sam.cpp
         process_python/rotate_registration.cpp
-)
+)
+
+add_subdirectory(cuda_impl)
+target_link_libraries(${PROJECT_NAME} ImageProcessCudaImpl)

src/image_process_v5/cuda_impl/CMakeLists.txt

@@ -0,0 +1,20 @@
+cmake_minimum_required(VERSION 3.25)
+project(ImageProcessCudaImpl LANGUAGES CXX CUDA)
+
+set(CMAKE_CXX_STANDARD 20)
+
+add_library(${PROJECT_NAME}
+        image_enhance.cu)
+
+target_include_directories(${PROJECT_NAME} PRIVATE ${PROJECT_SOURCE_PATH})
+
+# CUDA config
+find_package(CUDAToolkit REQUIRED)
+target_link_directories(${PROJECT_NAME} PRIVATE /usr/local/cuda/lib64)
+target_link_libraries(${PROJECT_NAME} CUDA::cudart CUDA::cuda_driver)
+set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_ARCHITECTURES "61;75;86")
+
+target_compile_options(${PROJECT_NAME} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:
+        -Xptxas -v # show kernel info
+        #        -g -G      # debug options
+        >)

src/image_process_v5/cuda_impl/image_enhance.cu

@@ -0,0 +1,104 @@
+#include "image_enhance.cuh"
+
+#include <thrust/reduce.h>
+#include <thrust/transform.h>
+
+struct extract_lum_op {
+    float eps_factor;
+
+    __device__
+    float operator()(const uchar3 pix) const {
+        const float lum = max(pix.x, max(pix.y, pix.z)) / 255.f;
+        // return logf(lum + eps_factor);
+        return lum;
+    }
+};
+
+// https://moontree.github.io/2020/08/30/tone-mapping/
+struct aces_enhance_op {
+    float lum_factor;
+
+    __device__
+    static float tone_mapping(const float lum) {
+        static constexpr float
+                A = 2.51, B = 0.03,
+                C = 2.43, D = 0.59, E = 0.14;
+        return lum * (A * lum + B) / (lum * (C * lum + D) + E);
+    }
+
+    __device__
+    uchar3 operator()(const uchar3 pix) const {
+        const uint8_t c_max = max(max(pix.x, pix.y), pix.z);
+        const uint8_t c_min = min(min(pix.x, pix.y), pix.z);
+        const uint8_t delta = c_max - c_min;
+
+        float h; // 60 is eliminated
+        if (delta == 0) {
+            h = 0;
+        } else {
+            const float delta_inv = 1.0f / delta;
+            // @formatter:off
+            if (c_max == pix.x) { // c_max == r
+                h = delta_inv * (pix.y - pix.z); // (g-b)/delta % 6
+                if (h < 0) h += 6;
+            } else if (c_max == pix.y) { // c_max == g
+                h = delta_inv * (pix.z - pix.x) + 2; // (b-r)/delta + 2
+            } else { // c_max == b
+                h = delta_inv * (pix.x - pix.y) + 4; // (r-g)/delta + 2
+            }
+            // @formatter:on
+        }
+
+        float s;
+        if (c_max == 0) {
+            s = 0;
+        } else {
+            s = (float) delta / c_max;
+        }
+
+        // enhance V channel
+        const float lum = c_max / 255.f;
+        const float v = __saturatef(tone_mapping(lum * lum_factor));
+
+        // convert HSV to RGB
+        // https://www.rapidtables.com/convert/color/hsv-to-rgb.html
+        const float c = v * s;
+        const float x = c * (1 - fabsf(fmodf(h, 2) - 1)); // c * (1 - |h % 2 - 1|)
+        const float m = v - c;
+        float r, g, b;
+        switch ((uint8_t) h) {
+            // @formatter:off
+            case 0: { r = c; g = x; b = 0; break; }
+            case 1: { r = x; g = c; b = 0; break; }
+            case 2: { r = 0; g = c; b = x; break; }
+            case 3: { r = 0; g = x; b = c; break; }
+            case 4: { r = x; g = 0; b = c; break; }
+            case 5: { r = c; g = 0; b = x; break; }
+            default: { r = 0; g = 0; b = 0; assert(false); }
+            // @formatter:on
+        }
+
+        return make_uchar3(255 * (r + m), 255 * (g + m), 255 * (b + m));
+    }
+};
+
+namespace image_enhance_impl {
+    float call_reduce_lum_mean(const uchar3 *in, float *buf, const uint32_t num,
+                               const float eps_factor, cudaStream_t stream) {
+        const auto op = extract_lum_op(eps_factor);
+        thrust::transform(thrust::cuda::par.on(stream),
+                          in, in + num, buf, op);
+        const auto ret = thrust::reduce(thrust::cuda::par.on(stream),
+                                        buf, buf + num, 0.f,
+                                        thrust::plus<float>());
+        // return expf(ret / num);
+        return ret / num;
+    }
+
+    void call_aces_enhance(const uchar3 *in, uchar3 *out, const uint32_t num,
+                           const float lum_factor, cudaStream_t stream) {
+        const auto op = aces_enhance_op(lum_factor);
+        thrust::transform(thrust::cuda::par.on(stream),
+                          in, in + num, out, op);
+    }
+}

src/image_process_v5/cuda_impl/image_enhance.cuh

@@ -0,0 +1,13 @@
+#ifndef IMAGE_ENHANCE_CUH
+#define IMAGE_ENHANCE_CUH
+
+#include <cuda.h>
+
+namespace image_enhance_impl {
+    float call_reduce_lum_mean(const uchar3 *in, float *buf, uint32_t num,
+                               float eps_factor, cudaStream_t stream);
+    void call_aces_enhance(const uchar3 *in, uchar3 *out, uint32_t num,
+                           float lum_factor, cudaStream_t stream);
+}
+
+#endif //IMAGE_ENHANCE_CUH

src/image_process_v5/image_enhance.cpp

@@ -0,0 +1,98 @@
+#include "image_enhance.h"
+#include "cuda_impl/image_enhance.cuh"
+
+namespace impl = image_enhance_impl;
+
+sp_image image_enhance(sp_image img, const image_enhance_options opts) {
+    assert(img.cv_type() == CV_8UC3);
+    img = create_dense(img);
+    const auto read_helper = read_access_helper(img.cuda());
+    float lum_mean = 0;
+    if (true) {
+        const auto buf = sp_image::create_like(img, CV_32FC1);
+        const auto write_helper = write_access_helper(buf.cuda());
+        lum_mean = impl::call_reduce_lum_mean(
+            (uchar3 *) read_helper.ptr(), (float *) write_helper.ptr(),
+            img.elem_count(), opts.black_eps_factor, current_cuda_stream());
+    }
+
+    const auto lum_factor = opts.luminance_target / lum_mean;
+    auto ret = sp_image::create_like(img);
+    const auto write_helper = write_access_helper(ret.cuda());
+    impl::call_aces_enhance(
+        (uchar3 *) read_helper.ptr(), (uchar3 *) write_helper.ptr(),
+        img.elem_count(), lum_factor, current_cuda_stream());
+    ret.merge_meta(img);
+    return ret;
+}
+
+#include "core/imgui_utility.hpp"
+
+struct image_enhance_ui::impl {
+    create_config conf;
+    obj_conn_type conn;
+
+    struct ui_config {
+        image_enhance_options opts;
+        bool passthrough = false;
+    };
+
+    std::shared_ptr<ui_config> ui_conf;
+
+    static void image_callback_impl(create_config conf) {
+        const auto img = OBJ_QUERY(sp_image, conf.in_name);
+        const auto ret_img = image_enhance(img, conf.opts);
+        OBJ_SAVE(conf.out_name, ret_img);
+    }
+
+    void image_callback() {
+        conf.opts = ui_conf->opts;
+        if (ui_conf->passthrough) {
+            OBJ_SAVE(conf.out_name, OBJ_QUERY(sp_image, conf.in_name));
+            return;
+        }
+
+        auto task = [conf=conf] {
+            try {
+                image_callback_impl(conf);
+            } catch (...) { (void) 0; }
+        };
+        TP_DETACH(task);
+    }
+
+    void show_ui() const {
+        ImGui::Checkbox("Passthrough", &ui_conf->passthrough);
+        auto disable_guard = imgui_disable_guard(ui_conf->passthrough);
+        ImGui::DragFloat("Luminance Target", &ui_conf->opts.luminance_target,
+                         0.002, 0, 2, "%.3f");
+        float eps_db = 10 * log10f(ui_conf->opts.black_eps_factor);
+        if (ImGui::DragFloat("Black Level (dB)", &eps_db,
+                             0.5, -70, -10, "%.1f")) {
+            ui_conf->opts.black_eps_factor = exp10f(eps_db / 10);
+        }
+    }
+
+    explicit impl(const create_config _conf) : conf(_conf) {
+        ui_conf = std::make_shared<ui_config>();
+        ui_conf->opts = conf.opts;
+        conn = OBJ_SIG(conf.in_name)->connect(
+            [this](auto _) { image_callback(); });
+    }
+
+    ~impl() {
+        conn.disconnect();
+    }
+};
+
+image_enhance_ui::image_enhance_ui(create_config conf)
+    : pimpl(std::make_unique<impl>(conf)) { (void) 0; }
+
+image_enhance_ui::~image_enhance_ui() = default;
+
+void image_enhance_ui::show_ui() const {
+    pimpl->show_ui();
+}
+
+void image_enhance_ui::sync_with(const image_enhance_ui *other) const {
+    pimpl->ui_conf = other->pimpl->ui_conf;
+}

src/image_process_v5/image_enhance.h

@@ -0,0 +1,35 @@
+#ifndef IMAGE_ENHANCE_H
+#define IMAGE_ENHANCE_H
+
+#include "sp_image.h"
+
+struct image_enhance_options {
+    float black_eps_factor = 0.01;
+    float luminance_target = 0.36;
+};
+
+sp_image image_enhance(sp_image img, image_enhance_options opts);
+
+#include "core_v2/object_manager.h"
+
+class image_enhance_ui {
+public:
+    struct create_config {
+        obj_name_type in_name = 0, out_name = 0;
+        image_enhance_options opts = {};
+    };
+
+    explicit image_enhance_ui(create_config conf);
+
+    ~image_enhance_ui();
+
+    void show_ui() const;
+
+    void sync_with(const image_enhance_ui *other) const;
+
+private:
+    struct impl;
+    std::unique_ptr<impl> pimpl;
+};
+
+#endif //IMAGE_ENHANCE_H

src/impl/apps/debug/app_debug.cpp

@@ -5,27 +5,30 @@
 
 #include <GLFW/glfw3.h>
 
-app_debug::app_debug(const create_config &conf) {
-    if (true) {
-        const auto img_fixed = sp_image::from_file(
-            "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_1.jpg");
-        const auto img_move = sp_image::from_file(
-            "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_51.jpg");
-        image_rotate_registration(img_fixed, img_move, [](rotate_registration_result_type ret) {
-            SPDLOG_DEBUG("angle: {}, center: ({}, {}), error: {}",
-                         ret.angle, ret.center.x, ret.center.y, ret.error);
-        });
-    }
+#include <opencv2/core/ocl.hpp>
 
-    if (true) {
-        const auto img_fixed = sp_image::from_file(
-            "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_51.jpg");
-        const auto img_move = sp_image::from_file(
-            "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_81.jpg");
-        image_rotate_registration(img_fixed, img_move, [](rotate_registration_result_type ret) {
-            SPDLOG_DEBUG("angle: {}, center: ({}, {}), error: {}",
-                         ret.angle, ret.center.x, ret.center.y, ret.error);
-            MAIN_DETACH([] { glfwSetWindowShouldClose(glfwGetCurrentContext(), true); });
-        });
-    }
+app_debug::app_debug(const create_config &conf) {
+    SPDLOG_DEBUG("OpenCL: {}", cv::ocl::useOpenCL());
+    // if (true) {
+    //     const auto img_fixed = sp_image::from_file(
+    //         "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_1.jpg");
+    //     const auto img_move = sp_image::from_file(
+    //         "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_51.jpg");
+    //     image_rotate_registration(img_fixed, img_move, [](rotate_registration_result_type ret) {
+    //         SPDLOG_DEBUG("angle: {}, center: ({}, {}), error: {}",
+    //                      ret.angle, ret.center.x, ret.center.y, ret.error);
+    //     });
+    // }
+    //
+    // if (true) {
+    //     const auto img_fixed = sp_image::from_file(
+    //         "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_51.jpg");
+    //     const auto img_move = sp_image::from_file(
+    //         "/home/tpx/ext/project/DepthGuide/cmake-build-debug/capture/rotation/Endoscope_81.jpg");
+    //     image_rotate_registration(img_fixed, img_move, [](rotate_registration_result_type ret) {
+    //         SPDLOG_DEBUG("angle: {}, center: ({}, {}), error: {}",
+    //                      ret.angle, ret.center.x, ret.center.y, ret.error);
+    //         MAIN_DETACH([] { glfwSetWindowShouldClose(glfwGetCurrentContext(), true); });
+    //     });
+    // }
 }

src/impl/apps/remote_ar/remote_ar_v2.cpp

@@ -7,6 +7,8 @@
 #include "device_v5/ndi_stray_point_tracker.h"
 
 // from sophiar
+#include <image_process_v5/cuda_impl/image_enhance.cuh>
+
 #include "core/local_connection.h"
 
 app_remote_ar_v2::app_remote_ar_v2(create_config _conf)
@@ -34,6 +36,13 @@ app_remote_ar_v2::app_remote_ar_v2(create_config _conf)
         uvc_cam.emplace(uvc_cam_conf);
     }
 
+    if (true) {
+        auto sub_conf = image_enhance_ui::create_config{
+            .in_name = left_img_id, .out_name = left_img_fix_id
+        };
+        img_isp.emplace(sub_conf);
+    }
+
     if (true) {
         auto sub_conf = stereo_output_helper::create_config();
         sub_conf.left_name = left_img_id;
@@ -47,7 +56,7 @@ app_remote_ar_v2::app_remote_ar_v2(create_config _conf)
     if (true) {
         auto sub_conf = image_viewer_v2::create_config();
         sub_conf.items.emplace_back(uvc_img_id, "Endoscope", true);
-        sub_conf.items.emplace_back(left_img_id, "Left", true);
+        sub_conf.items.emplace_back(left_img_fix_id, "Left", true);
         sub_conf.items.emplace_back(right_img_id, "Right", true);
         bg_viewer.emplace(sub_conf);
     }
@@ -61,21 +70,21 @@ app_remote_ar_v2::app_remote_ar_v2(create_config _conf)
 
     if (true) {
         auto sub_conf = versatile_saver::create_config();
-        // sub_conf.items.emplace_back(left_img_id, "Left", true);
-        // sub_conf.items.emplace_back(right_img_id, "Right", true);
+        sub_conf.items.emplace_back(left_img_id, "Left", true);
+        sub_conf.items.emplace_back(right_img_id, "Right", true);
         sub_conf.items.emplace_back(uvc_img_id, "Endoscope", true);
         saver.emplace(sub_conf);
     }
 
-    start_sophiar(LOAD_STR("sophiar_config"),
-                  LOAD_STR("sophiar_start_name"));
+    // start_sophiar(LOAD_STR("sophiar_config"),
+    //               LOAD_STR("sophiar_start_name"));
 
-    if (true) {
-        auto sub_conf = ndi_stray_point_tracker::create_config();
-        sub_conf.ip_addr = LOAD_STR("ndi_ip");
-        sub_conf.port = LOAD_NUMBER(uint16_t, "ndi_port");
-        create_ndi_stray_points_tracker(sub_conf);
-    }
+    // if (true) {
+    //     auto sub_conf = ndi_stray_point_tracker::create_config();
+    //     sub_conf.ip_addr = LOAD_STR("ndi_ip");
+    //     sub_conf.port = LOAD_NUMBER(uint16_t, "ndi_port");
+    //     create_ndi_stray_points_tracker(sub_conf);
+    // }
 
     // if (true) {
     //     auto sub_conf = oblique_calibrator::create_config();
@@ -97,6 +106,8 @@ void app_remote_ar_v2::show_ui() {
         if (ImGui::CollapsingHeader("MVS Camera")) {
             auto id_guard = imgui_id_guard("mvs_camera");
             mvs_cam->show();
+            ImGui::SeparatorText("Enhance");
+            img_isp->show_ui();
         }
 
         if (ImGui::CollapsingHeader("UVC Camera")) {

src/impl/apps/remote_ar/remote_ar_v2.h

@@ -6,6 +6,7 @@
 #include "device/uvc_camera_ui.h"
 #include "image_process_v5/image_viewer.h"
 #include "image_process_v5/image_process.h"
+#include "image_process_v5/image_enhance.h"
 #include "module/image_streamer.h"
 #include "module_v5/versatile_saver.h"
 #include "module_v5/oblique_calibrator.h"
@@ -30,10 +31,12 @@ private:
             left_img_id = 1,
             right_img_id = 2,
             output_img_id = 3,
-            uvc_img_id = 4;
+            uvc_img_id = 4,
+            left_img_fix_id = 5;
 
     std::optional<mvs_camera_ui> mvs_cam;
     std::optional<uvc_camera_ui> uvc_cam;
+    std::optional<image_enhance_ui> img_isp;
     std::optional<stereo_output_helper> output_helper;
     std::optional<image_viewer_v2> bg_viewer;
     std::optional<image_streamer> streamer;