Fix CI build failure: update fkl submodule and add morphology operations

cmake/discover_tests.cmake

@@ -42,7 +42,7 @@ function (discover_tests DIR)
 
         string(FIND ${cuda_source} "npp" is_npp)    	
 
-        target_link_libraries(${cuda_target} PRIVATE CUDA::nppc CUDA::nppial CUDA::nppidei CUDA::nppig) 								              
+        target_link_libraries(${cuda_target} PRIVATE CUDA::nppc CUDA::nppial CUDA::nppidei CUDA::nppig CUDA::nppim) 								              
 
     endforeach()
 endfunction()

include/fast_npp.h

@@ -23,6 +23,7 @@
 #include <fused_kernel/algorithms/image_processing/resize.h>
 #include <fused_kernel/algorithms/basic_ops/vector_ops.h>
 #include <fused_kernel/algorithms/basic_ops/arithmetic.h>
+#include <fused_kernel/algorithms/image_processing/morphology.h>
 #include <fused_kernel/algorithms/basic_ops/bitwise.h>
 #include <fused_kernel/algorithms/basic_ops/math.h>
 #include <fused_kernel/algorithms/basic_ops/memory_operations.h>
@@ -35,6 +36,39 @@
 
 namespace fastNPP {
 
+    // ===== Morphology: Erode (min) / Dilate (max) with REPLICATE border =====
+    // Implements rectangular (all-active) structuring elements over a
+    // configurable mask size and anchor. Border semantics match
+    // nppiErodeBorder / nppiDilateBorder with NPP_BORDER_REPLICATE when
+    // a full all-ones mask is used.
+    // Because these functions execute a GPU kernel directly they accept a
+    // caller-supplied ReadIOp and WriteIOp instead of hard-coding
+    // PerThreadRead / PerThreadWrite; the remaining parameters are standard
+    // NPP types (NppiSize, NppiPoint, NppStreamContext).
+#define FASTNPP_DEFINE_MORPH(NPPNAME, T, FKL_EXEC_FN)                               \
+    template <typename ReadIOp, typename WriteIOp>                                   \
+    inline void NPPNAME(const ReadIOp& input, const WriteIOp& output,               \
+                        NppiSize oSrcSize, NppiSize oMaskSize, NppiPoint oAnchor,   \
+                        NppStreamContext nppStreamCtx) {                              \
+        fk::MorphologyDPPDetails<T> details{};                                       \
+        details.width   = oSrcSize.width;                                            \
+        details.height  = oSrcSize.height;                                           \
+        details.maskW   = oMaskSize.width;                                           \
+        details.maskH   = oMaskSize.height;                                          \
+        details.anchorX = oAnchor.x;                                                 \
+        details.anchorY = oAnchor.y;                                                 \
+        fk::Stream stream(nppStreamCtx.hStream);                                     \
+        fk::FKL_EXEC_FN(details, input, output, stream);                             \
+    }
+
+    FASTNPP_DEFINE_MORPH(ErodeBorder_8u_C1R_Ctx,   uchar,  executeErode)
+    FASTNPP_DEFINE_MORPH(ErodeBorder_8u_C3R_Ctx,   uchar3, executeErode)
+    FASTNPP_DEFINE_MORPH(ErodeBorder_16u_C1R_Ctx,  ushort, executeErode)
+    FASTNPP_DEFINE_MORPH(ErodeBorder_32f_C1R_Ctx,  float,  executeErode)
+    FASTNPP_DEFINE_MORPH(DilateBorder_8u_C1R_Ctx,  uchar,  executeDilate)
+    FASTNPP_DEFINE_MORPH(DilateBorder_8u_C3R_Ctx,  uchar3, executeDilate)
+    FASTNPP_DEFINE_MORPH(DilateBorder_16u_C1R_Ctx, ushort, executeDilate)
+    FASTNPP_DEFINE_MORPH(DilateBorder_32f_C1R_Ctx, float,  executeDilate)
     // ===== AbsDiff with constant: |src - C| =====
     constexpr inline auto AbsDiffC_8u_C1R_Ctx(const uchar& nConstant) {
         return fk::AbsDiff<uchar>::build(nConstant);
@@ -59,11 +93,65 @@ namespace fastNPP {
     FASTNPP_DEFINE_SHIFT(RShiftC_16u_C1R_Ctx, ushort, ShiftRight)
     FASTNPP_DEFINE_SHIFT(RShiftC_32s_C1R_Ctx, int,    ShiftRight)
 
+    // ===== Dual-source Read composition for two-image operations =====
+    // Reads from two caller-supplied Read IOps at the same thread coordinates
+    // and returns their results as fk::Tuple<O1, O2>. This lets two-image
+    // fastNPP entry points take Read IOps (instead of raw NPP pointers) for
+    // both sources while still feeding a Tuple to the two-input Unary compute
+    // Operation that follows (e.g. Add<T,T,T,UnaryType>).
+    namespace detail {
+        template <typename BackIOp_>
+        struct DualSourceReadBack {
+            static_assert(fk::isTuple_v<BackIOp_> && BackIOp_::size == 2,
+                          "DualSourceReadBack expects an fk::Tuple with exactly two Read IOps");
+        private:
+            using Point = fk::Point;
+            using IOp1 = fk::get_t<0, BackIOp_>;
+            using IOp2 = fk::get_t<1, BackIOp_>;
+            static_assert(fk::isAnyCompleteReadType<IOp1> && fk::isAnyCompleteReadType<IOp2>,
+                          "Both elements of the BackIOp Tuple must be complete Read IOps");
+            using SelfType = DualSourceReadBack<BackIOp_>;
+        public:
+            FK_STATIC_STRUCT(DualSourceReadBack, SelfType)
+            using Parent = fk::ReadBackOperation<typename IOp1::Operation::ReadDataType, NullType, BackIOp_,
+                fk::Tuple<typename IOp1::Operation::OutputType, typename IOp2::Operation::OutputType>,
+                DualSourceReadBack<BackIOp_>>;
+            DECLARE_READBACK_PARENT_BASIC
+            FK_HOST_DEVICE_FUSE OutputType exec(const Point thread, const ParamsType&, const BackIOp& backIOp) {
+                return { IOp1::Operation::exec(thread, fk::get<0>(backIOp)),
+                         IOp2::Operation::exec(thread, fk::get<1>(backIOp)) };
+            }
+            FK_HOST_DEVICE_FUSE uint num_elems_x(const Point thread, const OperationDataType& opData) {
+                return IOp1::Operation::num_elems_x(thread, fk::get<0>(opData.backIOp));
+            }
+            FK_HOST_DEVICE_FUSE uint num_elems_y(const Point thread, const OperationDataType& opData) {
+                return IOp1::Operation::num_elems_y(thread, fk::get<0>(opData.backIOp));
+            }
+            FK_HOST_DEVICE_FUSE uint num_elems_z(const Point thread, const OperationDataType& opData) {
+                return IOp1::Operation::num_elems_z(thread, fk::get<0>(opData.backIOp));
+            }
+            FK_HOST_DEVICE_FUSE fk::ActiveThreads getActiveThreads(const OperationDataType& opData) {
+                return { num_elems_x(Point{0,0,0}, opData), num_elems_y(Point{0,0,0}, opData),
+                         num_elems_z(Point{0,0,0}, opData) };
+            }
+            FK_HOST_FUSE InstantiableType build(const IOp1& iop1, const IOp2& iop2) {
+                return { { NullType{}, BackIOp_{ iop1, iop2 } } };
+            }
+        };
+    } // namespace detail
+
     // ===== Two-image bitwise (And/Or/Xor) =====
-#define FASTNPP_DEFINE_TWO_IMAGE_BW(NPPNAME, T, FKLOP)                          \
-    inline auto NPPNAME(const fk::Ptr2D<T>& pSrc1, const fk::Ptr2D<T>& pSrc2) { \
-        return fk::DualSourceRead<fk::ND::_2D, T>::build(pSrc2, pSrc1)           \
-               .then(fk::FKLOP<T, T, T, fk::UnaryType>::build());               \
+    // Accepts two caller-supplied Read IOps instead of raw NPP pointers.
+    // detail::DualSourceReadBack reads both sources into an fk::Tuple<T, T>,
+    // consumed directly by the two-input Unary form of the bitwise Operation.
+    // NPP computes dst = pSrc2 OP pSrc1, so src2/src1 are fed in that order to
+    // reproduce NPP's operand order exactly. Returns a composed Read+Unary
+    // IOp; the caller appends a Write IOp to execute it via executeOperations.
+#define FASTNPP_DEFINE_TWO_IMAGE_BW(NPPNAME, T, FKLOP)                                      \
+    template <typename ReadIOp1, typename ReadIOp2>                                          \
+    inline auto NPPNAME(const ReadIOp1& src1, const ReadIOp2& src2) {                       \
+        return detail::DualSourceReadBack<fk::Tuple<ReadIOp2, ReadIOp1>>::build(src2, src1) \
+               .then(fk::FKLOP<T, T, T, fk::UnaryType>::build());                            \
     }
     FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C1R_Ctx, uchar,  BwAnd)
     FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C3R_Ctx, uchar3, BwAnd)
@@ -72,13 +160,12 @@ namespace fastNPP {
     FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C1R_Ctx, uchar,  BwXor)
     FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C3R_Ctx, uchar3, BwXor)
     // ===== Two-image element-wise arithmetic (32f) =====
-    // NPP computes dst = pSrc2 OP pSrc1; we feed (src2, src1) into DualSourceRead
-    // so the fused Unary operator reproduces NPP's operand order exactly.
-    // These return a complete read->op chain; the caller appends the write.
-#define FASTNPP_DEFINE_TWO_IMAGE(NPPNAME, T, FKLOP)                              \
-    inline auto NPPNAME(const fk::Ptr2D<T>& pSrc1, const fk::Ptr2D<T>& pSrc2) {  \
-        return fk::DualSourceRead<fk::ND::_2D, T>::build(pSrc2, pSrc1)            \
-               .then(fk::FKLOP<T, T, T, fk::UnaryType>::build());                \
+    // Same IOp-based composition as the bitwise two-image ops above.
+#define FASTNPP_DEFINE_TWO_IMAGE(NPPNAME, T, FKLOP)                                          \
+    template <typename ReadIOp1, typename ReadIOp2>                                          \
+    inline auto NPPNAME(const ReadIOp1& src1, const ReadIOp2& src2) {                       \
+        return detail::DualSourceReadBack<fk::Tuple<ReadIOp2, ReadIOp1>>::build(src2, src1) \
+               .then(fk::FKLOP<T, T, T, fk::UnaryType>::build());                            \
     }
 
     FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C1R_Ctx,  float,  Add)

tests/arithmetic/fastNPP_absdiff_shift_test.cu

@@ -103,19 +103,25 @@ int bwTwoImg(const char* label, NppFn nppFn, FKLChainFn chainFn) {
     const size_t N = (size_t)W * H;
     std::vector<Npp8u> h1(N), h2(N), ref(N), fkl(N);
     for (size_t i = 0; i < N; ++i) { h1[i] = (Npp8u)((i * 53) % 256); h2[i] = (Npp8u)((i * 97) % 256); }
-    fk::Ptr2D<uchar> s1(W, H), s2(W, H), d(W, H);
-    const int pitch = (int)s1.ptr().dims.pitch, rb = W;
+    fk::Ptr2D<uchar> d(W, H);
+    const int pitch = (int)d.ptr().dims.pitch, rb = W;
     Npp8u *d1, *d2, *dd;
     cudaMalloc(&d1, (size_t)pitch * H); cudaMalloc(&d2, (size_t)pitch * H); cudaMalloc(&dd, (size_t)pitch * H);
     cudaMemcpy2D(d1, pitch, h1.data(), rb, rb, H, cudaMemcpyHostToDevice);
     cudaMemcpy2D(d2, pitch, h2.data(), rb, rb, H, cudaMemcpyHostToDevice);
     nppFn(d1, pitch, d2, pitch, dd, pitch, {W, H}, makeCtx());
     cudaDeviceSynchronize();
     cudaMemcpy2D(ref.data(), rb, dd, pitch, rb, H, cudaMemcpyDeviceToHost);
-    cudaMemcpy2D(s1.ptr().data, pitch, h1.data(), rb, rb, H, cudaMemcpyHostToDevice);
-    cudaMemcpy2D(s2.ptr().data, pitch, h2.data(), rb, rb, H, cudaMemcpyHostToDevice);
+
+    int devID = 0;
+    cudaGetDevice(&devID);
+    fk::Ptr2D<uchar> f1(d1, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID);
+    fk::Ptr2D<uchar> f2(d2, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID);
+
     fk::Stream st;
-    fk::executeOperations<fk::TransformDPP<>>(st, chainFn(s1, s2),
+    fk::executeOperations<fk::TransformDPP<>>(st,
+        chainFn(fk::PerThreadRead<fk::ND::_2D, uchar>::build(f1),
+                fk::PerThreadRead<fk::ND::_2D, uchar>::build(f2)),
         fk::PerThreadWrite<fk::ND::_2D, uchar>::build(d));
     st.sync();
     cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost);
@@ -133,11 +139,11 @@ int launch() {
     bad += shiftC("LShiftC_8u_C1R", 1, nppiLShiftC_8u_C1R_Ctx, fastNPP::LShiftC_8u_C1R_Ctx(1));
     bad += shiftC("RShiftC_8u_C1R", 2, nppiRShiftC_8u_C1R_Ctx, fastNPP::RShiftC_8u_C1R_Ctx(2));
     bad += bwTwoImg("And_8u_C1R", nppiAnd_8u_C1R_Ctx,
-                    [](const fk::Ptr2D<uchar>& a, const fk::Ptr2D<uchar>& b){ return fastNPP::And_8u_C1R_Ctx(a, b); });
+                    [](const auto& a, const auto& b) { return fastNPP::And_8u_C1R_Ctx(a, b); });
     bad += bwTwoImg("Or_8u_C1R", nppiOr_8u_C1R_Ctx,
-                    [](const fk::Ptr2D<uchar>& a, const fk::Ptr2D<uchar>& b){ return fastNPP::Or_8u_C1R_Ctx(a, b); });
+                    [](const auto& a, const auto& b) { return fastNPP::Or_8u_C1R_Ctx(a, b); });
     bad += bwTwoImg("Xor_8u_C1R", nppiXor_8u_C1R_Ctx,
-                    [](const fk::Ptr2D<uchar>& a, const fk::Ptr2D<uchar>& b){ return fastNPP::Xor_8u_C1R_Ctx(a, b); });
+                    [](const auto& a, const auto& b) { return fastNPP::Xor_8u_C1R_Ctx(a, b); });
     printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED");
     return bad == 0 ? 0 : 1;
 }

tests/arithmetic/fastNPP_two_image_test.cu

@@ -52,8 +52,8 @@ int twoImgC1(const char* label, NppFn nppFn, FKLChainFn chainFn, bool avoidZero)
         h1[i] = (float)(i % 200) + (avoidZero ? 1.f : -50.f);
         h2[i] = (float)((i * 7) % 300) + 1.f;
     }
-    fk::Ptr2D<float> s1(W, H), s2(W, H), d(W, H);
-    const int pitch = (int)s1.ptr().dims.pitch, rb = W * sizeof(float);
+    fk::Ptr2D<float> d(W, H);
+    const int pitch = (int)d.ptr().dims.pitch, rb = W * sizeof(float);
     Npp32f *d1, *d2, *dd;
     cudaMalloc(&d1, (size_t)pitch * H);
     cudaMalloc(&d2, (size_t)pitch * H);
@@ -64,11 +64,15 @@ int twoImgC1(const char* label, NppFn nppFn, FKLChainFn chainFn, bool avoidZero)
     cudaDeviceSynchronize();
     cudaMemcpy2D(ref.data(), rb, dd, pitch, rb, H, cudaMemcpyDeviceToHost);
 
-    cudaMemcpy2D(s1.ptr().data, pitch, h1.data(), rb, rb, H, cudaMemcpyHostToDevice);
-    cudaMemcpy2D(s2.ptr().data, pitch, h2.data(), rb, rb, H, cudaMemcpyHostToDevice);
+    int devID = 0;
+    cudaGetDevice(&devID);
+    fk::Ptr2D<float> f1(d1, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID);
+    fk::Ptr2D<float> f2(d2, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID);
+
     fk::Stream st;
     fk::executeOperations<fk::TransformDPP<>>(st,
-        chainFn(s1, s2),
+        chainFn(fk::PerThreadRead<fk::ND::_2D, float>::build(f1),
+                fk::PerThreadRead<fk::ND::_2D, float>::build(f2)),
         fk::PerThreadWrite<fk::ND::_2D, float>::build(d));
     st.sync();
     cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost);
@@ -89,13 +93,13 @@ int twoImgC1(const char* label, NppFn nppFn, FKLChainFn chainFn, bool avoidZero)
 int launch() {
     int bad = 0;
     bad += twoImgC1("Add_32f_C1R", nppiAdd_32f_C1R_Ctx,
-                    [](const fk::Ptr2D<float>& a, const fk::Ptr2D<float>& b){ return fastNPP::Add_32f_C1R_Ctx(a, b); }, false);
+                    [](const auto& a, const auto& b) { return fastNPP::Add_32f_C1R_Ctx(a, b); }, false);
     bad += twoImgC1("Sub_32f_C1R", nppiSub_32f_C1R_Ctx,
-                    [](const fk::Ptr2D<float>& a, const fk::Ptr2D<float>& b){ return fastNPP::Sub_32f_C1R_Ctx(a, b); }, false);
+                    [](const auto& a, const auto& b) { return fastNPP::Sub_32f_C1R_Ctx(a, b); }, false);
     bad += twoImgC1("Mul_32f_C1R", nppiMul_32f_C1R_Ctx,
-                    [](const fk::Ptr2D<float>& a, const fk::Ptr2D<float>& b){ return fastNPP::Mul_32f_C1R_Ctx(a, b); }, false);
+                    [](const auto& a, const auto& b) { return fastNPP::Mul_32f_C1R_Ctx(a, b); }, false);
     bad += twoImgC1("Div_32f_C1R", nppiDiv_32f_C1R_Ctx,
-                    [](const fk::Ptr2D<float>& a, const fk::Ptr2D<float>& b){ return fastNPP::Div_32f_C1R_Ctx(a, b); }, true);
+                    [](const auto& a, const auto& b) { return fastNPP::Div_32f_C1R_Ctx(a, b); }, true);
     printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED");
     return bad == 0 ? 0 : 1;
 }