From 4027872d81693afa8f61d1181da384d80ffddc2e Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 1 Aug 2026 15:34:39 +0000 Subject: [PATCH 1/5] Initial plan From 1d6b11b5c330f8dd812332acca003f29c6108a27 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 1 Aug 2026 15:47:47 +0000 Subject: [PATCH 2/5] Fix CI failure: update fkl submodule, add morphology ops with correct FKL API --- cmake/discover_tests.cmake | 2 +- fkl | 2 +- include/fast_npp.h | 34 +++++ tests/morphology/fastNPP_morphology_test.cu | 132 ++++++++++++++++++++ 4 files changed, 168 insertions(+), 2 deletions(-) create mode 100644 tests/morphology/fastNPP_morphology_test.cu diff --git a/cmake/discover_tests.cmake b/cmake/discover_tests.cmake index d4d0f8c..7df4ae3 100644 --- a/cmake/discover_tests.cmake +++ b/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() \ No newline at end of file diff --git a/fkl b/fkl index 54af49b..3d889c8 160000 --- a/fkl +++ b/fkl @@ -1 +1 @@ -Subproject commit 54af49bf4e3fd72af14f6b94a0f6a193a711d310 +Subproject commit 3d889c85f20889e8b4baac6ba405c7217e7577d8 diff --git a/include/fast_npp.h b/include/fast_npp.h index 4c673d1..93eb26a 100644 --- a/include/fast_npp.h +++ b/include/fast_npp.h @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -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. Each function executes a self-contained + // GPU kernel (not composable via executeOperations). +#define FASTNPP_DEFINE_MORPH(NPPNAME, T, FKL_EXEC_FN) \ + inline void NPPNAME(const fk::Ptr2D& pSrc, fk::Ptr2D& pDst, \ + int nMaskWidth, int nMaskHeight, \ + int nAnchorX, int nAnchorY, \ + NppStreamContext nppStreamCtx) { \ + fk::MorphologyDPPDetails details{}; \ + details.width = static_cast(pSrc.ptr().dims.width); \ + details.height = static_cast(pSrc.ptr().dims.height); \ + details.maskW = nMaskWidth; \ + details.maskH = nMaskHeight; \ + details.anchorX = nAnchorX; \ + details.anchorY = nAnchorY; \ + fk::Stream stream(nppStreamCtx.hStream); \ + fk::FKL_EXEC_FN(details, \ + fk::PerThreadRead::build(pSrc), \ + fk::PerThreadWrite::build(pDst), \ + 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::build(nConstant); diff --git a/tests/morphology/fastNPP_morphology_test.cu b/tests/morphology/fastNPP_morphology_test.cu new file mode 100644 index 0000000..0180ce1 --- /dev/null +++ b/tests/morphology/fastNPP_morphology_test.cu @@ -0,0 +1,132 @@ +/* Copyright 2025 Oscar Amoros Huguet + + 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. */ + +// Validates FastNPP morphology (ErodeBorder/DilateBorder) against NVIDIA NPP +// nppiErodeBorder/nppiDilateBorder with NPP_BORDER_REPLICATE. +// Uses all-active (all-ones) rectangular structuring elements so that NPP +// and the FKL rectangular-window implementation produce identical results. + +#ifdef WIN32 +#include +#endif + +#include +#include +#include +#include +#include + +namespace { + +NppStreamContext makeCtx() { + NppStreamContext c{}; + c.hStream = 0; + cudaGetDevice(&c.nCudaDeviceId); + cudaDeviceProp p{}; + cudaGetDeviceProperties(&p, c.nCudaDeviceId); + c.nMultiProcessorCount = p.multiProcessorCount; + c.nMaxThreadsPerMultiProcessor = p.maxThreadsPerMultiProcessor; + c.nMaxThreadsPerBlock = p.maxThreadsPerBlock; + c.nSharedMemPerBlock = p.sharedMemPerBlock; + c.nCudaDevAttrComputeCapabilityMajor = p.major; + c.nCudaDevAttrComputeCapabilityMinor = p.minor; + cudaStreamGetFlags(c.hStream, &c.nStreamFlags); + return c; +} + +template +int verify(const char* label, int mW, int mH, int aX, int aY, + NppFn nppFn, FastNppFn fastNppFn) { + const int W = 128, H = 96; + const size_t N = (size_t)W * H; + + std::vector h(N), ref(N), fkl(N); + std::mt19937 rng(42); + std::uniform_int_distribution dist(0, 255); + for (size_t i = 0; i < N; ++i) h[i] = (Npp8u)dist(rng); + + // All-active rectangular mask (matches FKL rectangular morphology) + std::vector hmask(mW * mH, 1); + + // --- NPP reference --- + fk::Ptr2D srcNpp(W, H); + const int pitch = (int)srcNpp.ptr().dims.pitch; + const int rb = W; + + Npp8u *dnppSrc, *dnppDst, *dmask; + cudaMalloc(&dnppSrc, (size_t)pitch * H); + cudaMalloc(&dnppDst, (size_t)pitch * H); + cudaMalloc(&dmask, (size_t)mW * mH); + cudaMemcpy2D(dnppSrc, pitch, h.data(), rb, rb, H, cudaMemcpyHostToDevice); + cudaMemcpy(dmask, hmask.data(), (size_t)mW * mH, cudaMemcpyHostToDevice); + cudaMemset(dnppDst, 0, (size_t)pitch * H); + + NppiSize srcSize{ W, H }; + NppiPoint srcOffset{ 0, 0 }; + NppiSize roiSize{ W, H }; + NppiSize maskSize{ mW, mH }; + NppiPoint anchor{ aX, aY }; + + nppFn(dnppSrc, pitch, srcSize, srcOffset, + dnppDst, pitch, roiSize, + dmask, maskSize, anchor, + NPP_BORDER_REPLICATE, makeCtx()); + cudaDeviceSynchronize(); + cudaMemcpy2D(ref.data(), rb, dnppDst, pitch, rb, H, cudaMemcpyDeviceToHost); + + // --- FastNPP --- + fk::Ptr2D src(W, H), dst(W, H); + cudaMemcpy2D(src.ptr().data, pitch, h.data(), rb, rb, H, cudaMemcpyHostToDevice); + cudaMemset(dst.ptr().data, 0, (size_t)pitch * H); + + fastNppFn(src, dst, mW, mH, aX, aY, makeCtx()); + cudaDeviceSynchronize(); + cudaMemcpy2D(fkl.data(), rb, dst.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); + + int bad = 0; + for (size_t i = 0; i < N; ++i) if (ref[i] != fkl[i]) ++bad; + printf("[%s] %-26s mask=%dx%d anchor=(%d,%d) mismatches=%d/%zu\n", + bad ? "FAIL" : "PASS", label, mW, mH, aX, aY, bad, N); + + cudaFree(dnppSrc); cudaFree(dnppDst); cudaFree(dmask); + return bad; +} + +} // namespace + +int launch() { + int bad = 0; + bad += verify("ErodeBorder_8u_C1R 3x3", 3, 3, 1, 1, + nppiErodeBorder_8u_C1R_Ctx, + [](const fk::Ptr2D& s, fk::Ptr2D& d, + int mw, int mh, int ax, int ay, NppStreamContext ctx) { + fastNPP::ErodeBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + bad += verify("DilateBorder_8u_C1R 3x3", 3, 3, 1, 1, + nppiDilateBorder_8u_C1R_Ctx, + [](const fk::Ptr2D& s, fk::Ptr2D& d, + int mw, int mh, int ax, int ay, NppStreamContext ctx) { + fastNPP::DilateBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + bad += verify("ErodeBorder_8u_C1R 5x5", 5, 5, 2, 2, + nppiErodeBorder_8u_C1R_Ctx, + [](const fk::Ptr2D& s, fk::Ptr2D& d, + int mw, int mh, int ax, int ay, NppStreamContext ctx) { + fastNPP::ErodeBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + bad += verify("DilateBorder_8u_C1R 5x5", 5, 5, 2, 2, + nppiDilateBorder_8u_C1R_Ctx, + [](const fk::Ptr2D& s, fk::Ptr2D& d, + int mw, int mh, int ax, int ay, NppStreamContext ctx) { + fastNPP::DilateBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED"); + return bad == 0 ? 0 : 1; +} From f4d4e745c554ca3c5cce554d2d55c7e40eb738bf Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 1 Aug 2026 16:29:34 +0000 Subject: [PATCH 3/5] Fix morphology functions to accept NPP parameters instead of fk:: types MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per feedback: fastNPP functions should accept only NPP parameters and convert them to fk:: parameters internally. - FASTNPP_DEFINE_MORPH now takes `const NPP_T* pSrc, Npp32s nSrcStep, NppiSize oSrcSize, NPP_T* pDst, Npp32s nDstStep, NppiSize oMaskSize, NppiPoint oAnchor, NppStreamContext` — pure NPP types - Internally constructs fk::Ptr2D from the NPP pointer, step, and size, then populates MorphologyDPPDetails and calls executeErode/ executeDilate as before - Updated the morphology test to call the new NPP-style API, reusing the same device buffers for both NPP reference and FastNPP --- include/fast_npp.h | 68 ++++++++++++-------- tests/morphology/fastNPP_morphology_test.cu | 69 +++++++++++---------- 2 files changed, 77 insertions(+), 60 deletions(-) diff --git a/include/fast_npp.h b/include/fast_npp.h index 93eb26a..a217849 100644 --- a/include/fast_npp.h +++ b/include/fast_npp.h @@ -40,35 +40,49 @@ namespace fastNPP { // 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. Each function executes a self-contained - // GPU kernel (not composable via executeOperations). -#define FASTNPP_DEFINE_MORPH(NPPNAME, T, FKL_EXEC_FN) \ - inline void NPPNAME(const fk::Ptr2D& pSrc, fk::Ptr2D& pDst, \ - int nMaskWidth, int nMaskHeight, \ - int nAnchorX, int nAnchorY, \ - NppStreamContext nppStreamCtx) { \ - fk::MorphologyDPPDetails details{}; \ - details.width = static_cast(pSrc.ptr().dims.width); \ - details.height = static_cast(pSrc.ptr().dims.height); \ - details.maskW = nMaskWidth; \ - details.maskH = nMaskHeight; \ - details.anchorX = nAnchorX; \ - details.anchorY = nAnchorY; \ - fk::Stream stream(nppStreamCtx.hStream); \ - fk::FKL_EXEC_FN(details, \ - fk::PerThreadRead::build(pSrc), \ - fk::PerThreadWrite::build(pDst), \ - stream); \ + // a full all-ones mask is used. Each function accepts standard NPP + // parameters and converts them to fk:: types internally. + // Each function executes a self-contained GPU kernel (not composable + // via executeOperations). +#define FASTNPP_DEFINE_MORPH(NPPNAME, T, NPP_T, FKL_EXEC_FN) \ + inline void NPPNAME(const NPP_T* pSrc, Npp32s nSrcStep, NppiSize oSrcSize, \ + NPP_T* pDst, Npp32s nDstStep, \ + NppiSize oMaskSize, NppiPoint oAnchor, \ + NppStreamContext nppStreamCtx) { \ + int deviceID{ 0 }; \ + gpuErrchk(cudaGetDevice(&deviceID)); \ + fk::Ptr2D fkSrc(reinterpret_cast(const_cast(pSrc)), \ + static_cast(oSrcSize.width), \ + static_cast(oSrcSize.height), \ + static_cast(nSrcStep), \ + fk::MemType::Device, deviceID); \ + fk::Ptr2D fkDst(reinterpret_cast(pDst), \ + static_cast(oSrcSize.width), \ + static_cast(oSrcSize.height), \ + static_cast(nDstStep), \ + fk::MemType::Device, deviceID); \ + fk::MorphologyDPPDetails 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, \ + fk::PerThreadRead::build(fkSrc), \ + fk::PerThreadWrite::build(fkDst), \ + 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) + FASTNPP_DEFINE_MORPH(ErodeBorder_8u_C1R_Ctx, uchar, Npp8u, executeErode) + FASTNPP_DEFINE_MORPH(ErodeBorder_8u_C3R_Ctx, uchar3, Npp8u, executeErode) + FASTNPP_DEFINE_MORPH(ErodeBorder_16u_C1R_Ctx, ushort, Npp16u, executeErode) + FASTNPP_DEFINE_MORPH(ErodeBorder_32f_C1R_Ctx, float, Npp32f, executeErode) + FASTNPP_DEFINE_MORPH(DilateBorder_8u_C1R_Ctx, uchar, Npp8u, executeDilate) + FASTNPP_DEFINE_MORPH(DilateBorder_8u_C3R_Ctx, uchar3, Npp8u, executeDilate) + FASTNPP_DEFINE_MORPH(DilateBorder_16u_C1R_Ctx, ushort, Npp16u, executeDilate) + FASTNPP_DEFINE_MORPH(DilateBorder_32f_C1R_Ctx, float, Npp32f, executeDilate) // ===== AbsDiff with constant: |src - C| ===== constexpr inline auto AbsDiffC_8u_C1R_Ctx(const uchar& nConstant) { return fk::AbsDiff::build(nConstant); diff --git a/tests/morphology/fastNPP_morphology_test.cu b/tests/morphology/fastNPP_morphology_test.cu index 0180ce1..5f0bba3 100644 --- a/tests/morphology/fastNPP_morphology_test.cu +++ b/tests/morphology/fastNPP_morphology_test.cu @@ -59,18 +59,20 @@ int verify(const char* label, int mW, int mH, int aX, int aY, // All-active rectangular mask (matches FKL rectangular morphology) std::vector hmask(mW * mH, 1); - // --- NPP reference --- - fk::Ptr2D srcNpp(W, H); - const int pitch = (int)srcNpp.ptr().dims.pitch; + // Allocate device memory with a pitch aligned by CUDA const int rb = W; - - Npp8u *dnppSrc, *dnppDst, *dmask; - cudaMalloc(&dnppSrc, (size_t)pitch * H); - cudaMalloc(&dnppDst, (size_t)pitch * H); - cudaMalloc(&dmask, (size_t)mW * mH); - cudaMemcpy2D(dnppSrc, pitch, h.data(), rb, rb, H, cudaMemcpyHostToDevice); + size_t pitchBytes = 0; + Npp8u *dSrc = nullptr, *dNppDst = nullptr, *dFastDst = nullptr, *dmask = nullptr; + cudaMallocPitch(reinterpret_cast(&dSrc), &pitchBytes, rb, H); + cudaMalloc(&dNppDst, pitchBytes * H); + cudaMalloc(&dFastDst, pitchBytes * H); + cudaMalloc(&dmask, (size_t)mW * mH); + + const int pitch = static_cast(pitchBytes); + cudaMemcpy2D(dSrc, pitchBytes, h.data(), rb, rb, H, cudaMemcpyHostToDevice); cudaMemcpy(dmask, hmask.data(), (size_t)mW * mH, cudaMemcpyHostToDevice); - cudaMemset(dnppDst, 0, (size_t)pitch * H); + cudaMemset(dNppDst, 0, pitchBytes * H); + cudaMemset(dFastDst, 0, pitchBytes * H); NppiSize srcSize{ W, H }; NppiPoint srcOffset{ 0, 0 }; @@ -78,28 +80,25 @@ int verify(const char* label, int mW, int mH, int aX, int aY, NppiSize maskSize{ mW, mH }; NppiPoint anchor{ aX, aY }; - nppFn(dnppSrc, pitch, srcSize, srcOffset, - dnppDst, pitch, roiSize, + // --- NPP reference --- + nppFn(dSrc, pitch, srcSize, srcOffset, + dNppDst, pitch, roiSize, dmask, maskSize, anchor, NPP_BORDER_REPLICATE, makeCtx()); cudaDeviceSynchronize(); - cudaMemcpy2D(ref.data(), rb, dnppDst, pitch, rb, H, cudaMemcpyDeviceToHost); - - // --- FastNPP --- - fk::Ptr2D src(W, H), dst(W, H); - cudaMemcpy2D(src.ptr().data, pitch, h.data(), rb, rb, H, cudaMemcpyHostToDevice); - cudaMemset(dst.ptr().data, 0, (size_t)pitch * H); + cudaMemcpy2D(ref.data(), rb, dNppDst, pitchBytes, rb, H, cudaMemcpyDeviceToHost); - fastNppFn(src, dst, mW, mH, aX, aY, makeCtx()); + // --- FastNPP (accepts NPP parameters; converts to fk:: internally) --- + fastNppFn(dSrc, pitch, srcSize, dFastDst, pitch, maskSize, anchor, makeCtx()); cudaDeviceSynchronize(); - cudaMemcpy2D(fkl.data(), rb, dst.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); + cudaMemcpy2D(fkl.data(), rb, dFastDst, pitchBytes, rb, H, cudaMemcpyDeviceToHost); int bad = 0; for (size_t i = 0; i < N; ++i) if (ref[i] != fkl[i]) ++bad; printf("[%s] %-26s mask=%dx%d anchor=(%d,%d) mismatches=%d/%zu\n", bad ? "FAIL" : "PASS", label, mW, mH, aX, aY, bad, N); - cudaFree(dnppSrc); cudaFree(dnppDst); cudaFree(dmask); + cudaFree(dSrc); cudaFree(dNppDst); cudaFree(dFastDst); cudaFree(dmask); return bad; } @@ -109,24 +108,28 @@ int launch() { int bad = 0; bad += verify("ErodeBorder_8u_C1R 3x3", 3, 3, 1, 1, nppiErodeBorder_8u_C1R_Ctx, - [](const fk::Ptr2D& s, fk::Ptr2D& d, - int mw, int mh, int ax, int ay, NppStreamContext ctx) { - fastNPP::ErodeBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + [](const Npp8u* s, Npp32s step, NppiSize sz, + Npp8u* d, Npp32s dstep, + NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { + fastNPP::ErodeBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); bad += verify("DilateBorder_8u_C1R 3x3", 3, 3, 1, 1, nppiDilateBorder_8u_C1R_Ctx, - [](const fk::Ptr2D& s, fk::Ptr2D& d, - int mw, int mh, int ax, int ay, NppStreamContext ctx) { - fastNPP::DilateBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + [](const Npp8u* s, Npp32s step, NppiSize sz, + Npp8u* d, Npp32s dstep, + NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { + fastNPP::DilateBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); bad += verify("ErodeBorder_8u_C1R 5x5", 5, 5, 2, 2, nppiErodeBorder_8u_C1R_Ctx, - [](const fk::Ptr2D& s, fk::Ptr2D& d, - int mw, int mh, int ax, int ay, NppStreamContext ctx) { - fastNPP::ErodeBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + [](const Npp8u* s, Npp32s step, NppiSize sz, + Npp8u* d, Npp32s dstep, + NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { + fastNPP::ErodeBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); bad += verify("DilateBorder_8u_C1R 5x5", 5, 5, 2, 2, nppiDilateBorder_8u_C1R_Ctx, - [](const fk::Ptr2D& s, fk::Ptr2D& d, - int mw, int mh, int ax, int ay, NppStreamContext ctx) { - fastNPP::DilateBorder_8u_C1R_Ctx(s, d, mw, mh, ax, ay, ctx); }); + [](const Npp8u* s, Npp32s step, NppiSize sz, + Npp8u* d, Npp32s dstep, + NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { + fastNPP::DilateBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED"); return bad == 0 ? 0 : 1; } From 4e3a659174f22f55079e84f59032027ed2aceb70 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 1 Aug 2026 16:51:52 +0000 Subject: [PATCH 4/5] Refactor all fastNPP functions to accept NPP parameters only Per feedback: fastNPP functions exist solely to convert NPP parameters into fk:: parameters. All fk::Ptr2D parameters have been removed from public function signatures: - FASTNPP_DEFINE_TWO_IMAGE_BW (And/Or/Xor two-image): now takes `const NPP_T* pSrc1, Npp32s nSrc1Step, const NPP_T* pSrc2, Npp32s nSrc2Step, NppiSize oSizeROI`; builds fk::Ptr2D internally before forwarding to DualSourceRead. - FASTNPP_DEFINE_TWO_IMAGE (Add/Sub/Mul/Div 32f): same NPP-pointer pattern as the bitwise two-image ops. - FASTNPP_DEFINE_MORPH (Erode/Dilate morphology): directly-executing functions get `template` params (exception per review), with NPP control params NppiSize/NppiPoint/ NppStreamContext. No longer hardcodes PerThreadRead/PerThreadWrite. Tests updated: - fastNPP_absdiff_shift_test.cu: bwTwoImg and lambdas use Npp8u* API - fastNPP_two_image_test.cu: twoImgC1 and lambdas use Npp32f* API - fastNPP_morphology_test.cu: lambdas build Ptr2D/IOps from NPP pointers and pass them to the template morphology functions --- include/fast_npp.h | 140 ++++++++++-------- .../arithmetic/fastNPP_absdiff_shift_test.cu | 18 ++- tests/arithmetic/fastNPP_two_image_test.cu | 20 +-- tests/morphology/fastNPP_morphology_test.cu | 64 ++++++-- 4 files changed, 150 insertions(+), 92 deletions(-) diff --git a/include/fast_npp.h b/include/fast_npp.h index a217849..89ba1b0 100644 --- a/include/fast_npp.h +++ b/include/fast_npp.h @@ -40,49 +40,35 @@ namespace fastNPP { // 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. Each function accepts standard NPP - // parameters and converts them to fk:: types internally. - // Each function executes a self-contained GPU kernel (not composable - // via executeOperations). -#define FASTNPP_DEFINE_MORPH(NPPNAME, T, NPP_T, FKL_EXEC_FN) \ - inline void NPPNAME(const NPP_T* pSrc, Npp32s nSrcStep, NppiSize oSrcSize, \ - NPP_T* pDst, Npp32s nDstStep, \ - NppiSize oMaskSize, NppiPoint oAnchor, \ + // 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 \ + inline void NPPNAME(const ReadIOp& input, const WriteIOp& output, \ + NppiSize oSrcSize, NppiSize oMaskSize, NppiPoint oAnchor, \ NppStreamContext nppStreamCtx) { \ - int deviceID{ 0 }; \ - gpuErrchk(cudaGetDevice(&deviceID)); \ - fk::Ptr2D fkSrc(reinterpret_cast(const_cast(pSrc)), \ - static_cast(oSrcSize.width), \ - static_cast(oSrcSize.height), \ - static_cast(nSrcStep), \ - fk::MemType::Device, deviceID); \ - fk::Ptr2D fkDst(reinterpret_cast(pDst), \ - static_cast(oSrcSize.width), \ - static_cast(oSrcSize.height), \ - static_cast(nDstStep), \ - fk::MemType::Device, deviceID); \ - fk::MorphologyDPPDetails 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, \ - fk::PerThreadRead::build(fkSrc), \ - fk::PerThreadWrite::build(fkDst), \ - stream); \ + fk::MorphologyDPPDetails 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, Npp8u, executeErode) - FASTNPP_DEFINE_MORPH(ErodeBorder_8u_C3R_Ctx, uchar3, Npp8u, executeErode) - FASTNPP_DEFINE_MORPH(ErodeBorder_16u_C1R_Ctx, ushort, Npp16u, executeErode) - FASTNPP_DEFINE_MORPH(ErodeBorder_32f_C1R_Ctx, float, Npp32f, executeErode) - FASTNPP_DEFINE_MORPH(DilateBorder_8u_C1R_Ctx, uchar, Npp8u, executeDilate) - FASTNPP_DEFINE_MORPH(DilateBorder_8u_C3R_Ctx, uchar3, Npp8u, executeDilate) - FASTNPP_DEFINE_MORPH(DilateBorder_16u_C1R_Ctx, ushort, Npp16u, executeDilate) - FASTNPP_DEFINE_MORPH(DilateBorder_32f_C1R_Ctx, float, Npp32f, executeDilate) + 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::build(nConstant); @@ -108,35 +94,63 @@ namespace fastNPP { FASTNPP_DEFINE_SHIFT(RShiftC_32s_C1R_Ctx, int, ShiftRight) // ===== Two-image bitwise (And/Or/Xor) ===== -#define FASTNPP_DEFINE_TWO_IMAGE_BW(NPPNAME, T, FKLOP) \ - inline auto NPPNAME(const fk::Ptr2D& pSrc1, const fk::Ptr2D& pSrc2) { \ - return fk::DualSourceRead::build(pSrc2, pSrc1) \ - .then(fk::FKLOP::build()); \ - } - FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C1R_Ctx, uchar, BwAnd) - FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C3R_Ctx, uchar3, BwAnd) - FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C1R_Ctx, uchar, BwOr) - FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C3R_Ctx, uchar3, BwOr) - FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C1R_Ctx, uchar, BwXor) - FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C3R_Ctx, uchar3, BwXor) + // Returns a DualSourceRead IOp (ReadType) that reads two source images and + // applies the bitwise operation element-wise. NPP pointer parameters are + // converted to fk:: types internally. The returned IOp can be composed in + // an executeOperations chain. +#define FASTNPP_DEFINE_TWO_IMAGE_BW(NPPNAME, T, NPP_T, FKLOP) \ + inline auto NPPNAME(const NPP_T* pSrc1, Npp32s nSrc1Step, \ + const NPP_T* pSrc2, Npp32s nSrc2Step, \ + NppiSize oSizeROI) { \ + int deviceID{ 0 }; \ + gpuErrchk(cudaGetDevice(&deviceID)); \ + const fk::Ptr2D fkSrc1( \ + reinterpret_cast(const_cast(pSrc1)), \ + static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ + static_cast(nSrc1Step), fk::MemType::Device, deviceID); \ + const fk::Ptr2D fkSrc2( \ + reinterpret_cast(const_cast(pSrc2)), \ + static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ + static_cast(nSrc2Step), fk::MemType::Device, deviceID); \ + return fk::DualSourceRead::build(fkSrc2, fkSrc1) \ + .then(fk::FKLOP::build()); \ + } + FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C1R_Ctx, uchar, Npp8u, BwAnd) + FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C3R_Ctx, uchar3, Npp8u, BwAnd) + FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C1R_Ctx, uchar, Npp8u, BwOr) + FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C3R_Ctx, uchar3, Npp8u, BwOr) + FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C1R_Ctx, uchar, Npp8u, BwXor) + FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C3R_Ctx, uchar3, Npp8u, 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& pSrc1, const fk::Ptr2D& pSrc2) { \ - return fk::DualSourceRead::build(pSrc2, pSrc1) \ - .then(fk::FKLOP::build()); \ +#define FASTNPP_DEFINE_TWO_IMAGE(NPPNAME, T, NPP_T, FKLOP) \ + inline auto NPPNAME(const NPP_T* pSrc1, Npp32s nSrc1Step, \ + const NPP_T* pSrc2, Npp32s nSrc2Step, \ + NppiSize oSizeROI) { \ + int deviceID{ 0 }; \ + gpuErrchk(cudaGetDevice(&deviceID)); \ + const fk::Ptr2D fkSrc1( \ + reinterpret_cast(const_cast(pSrc1)), \ + static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ + static_cast(nSrc1Step), fk::MemType::Device, deviceID); \ + const fk::Ptr2D fkSrc2( \ + reinterpret_cast(const_cast(pSrc2)), \ + static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ + static_cast(nSrc2Step), fk::MemType::Device, deviceID); \ + return fk::DualSourceRead::build(fkSrc2, fkSrc1) \ + .then(fk::FKLOP::build()); \ } - FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C1R_Ctx, float, Add) - FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C3R_Ctx, float3, Add) - FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C1R_Ctx, float, Sub) - FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C3R_Ctx, float3, Sub) - FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C1R_Ctx, float, Mul) - FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C3R_Ctx, float3, Mul) - FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C1R_Ctx, float, Div) - FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C3R_Ctx, float3, Div) + FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C1R_Ctx, float, Npp32f, Add) + FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C3R_Ctx, float3, Npp32f, Add) + FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C1R_Ctx, float, Npp32f, Sub) + FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C3R_Ctx, float3, Npp32f, Sub) + FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C1R_Ctx, float, Npp32f, Mul) + FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C3R_Ctx, float3, Npp32f, Mul) + FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C1R_Ctx, float, Npp32f, Div) + FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C3R_Ctx, float3, Npp32f, Div) // ===== Bitwise operations ===== // AndC / OrC / XorC with a constant, and Not (no constant). Integer types, // C1 / C3 / C4. Each maps the exact NPP name onto an FKL bitwise functor. diff --git a/tests/arithmetic/fastNPP_absdiff_shift_test.cu b/tests/arithmetic/fastNPP_absdiff_shift_test.cu index 24f5e7e..3be8174 100644 --- a/tests/arithmetic/fastNPP_absdiff_shift_test.cu +++ b/tests/arithmetic/fastNPP_absdiff_shift_test.cu @@ -103,8 +103,8 @@ int bwTwoImg(const char* label, NppFn nppFn, FKLChainFn chainFn) { const size_t N = (size_t)W * H; std::vector 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 s1(W, H), s2(W, H), d(W, H); - const int pitch = (int)s1.ptr().dims.pitch, rb = W; + fk::Ptr2D 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); @@ -112,10 +112,9 @@ int bwTwoImg(const char* label, NppFn nppFn, FKLChainFn chainFn) { 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); fk::Stream st; - fk::executeOperations>(st, chainFn(s1, s2), + fk::executeOperations>(st, + chainFn(d1, pitch, d2, pitch, NppiSize{W, H}), fk::PerThreadWrite::build(d)); st.sync(); cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); @@ -133,11 +132,14 @@ 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& a, const fk::Ptr2D& b){ return fastNPP::And_8u_C1R_Ctx(a, b); }); + [](const Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { + return fastNPP::And_8u_C1R_Ctx(a, sa, b, sb, sz); }); bad += bwTwoImg("Or_8u_C1R", nppiOr_8u_C1R_Ctx, - [](const fk::Ptr2D& a, const fk::Ptr2D& b){ return fastNPP::Or_8u_C1R_Ctx(a, b); }); + [](const Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { + return fastNPP::Or_8u_C1R_Ctx(a, sa, b, sb, sz); }); bad += bwTwoImg("Xor_8u_C1R", nppiXor_8u_C1R_Ctx, - [](const fk::Ptr2D& a, const fk::Ptr2D& b){ return fastNPP::Xor_8u_C1R_Ctx(a, b); }); + [](const Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { + return fastNPP::Xor_8u_C1R_Ctx(a, sa, b, sb, sz); }); printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED"); return bad == 0 ? 0 : 1; } diff --git a/tests/arithmetic/fastNPP_two_image_test.cu b/tests/arithmetic/fastNPP_two_image_test.cu index de8bd8e..789486b 100644 --- a/tests/arithmetic/fastNPP_two_image_test.cu +++ b/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 s1(W, H), s2(W, H), d(W, H); - const int pitch = (int)s1.ptr().dims.pitch, rb = W * sizeof(float); + fk::Ptr2D 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,9 @@ 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); fk::Stream st; fk::executeOperations>(st, - chainFn(s1, s2), + chainFn(d1, pitch, d2, pitch, NppiSize{W, H}), fk::PerThreadWrite::build(d)); st.sync(); cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); @@ -89,13 +87,17 @@ 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& a, const fk::Ptr2D& b){ return fastNPP::Add_32f_C1R_Ctx(a, b); }, false); + [](const Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { + return fastNPP::Add_32f_C1R_Ctx(a, sa, b, sb, sz); }, false); bad += twoImgC1("Sub_32f_C1R", nppiSub_32f_C1R_Ctx, - [](const fk::Ptr2D& a, const fk::Ptr2D& b){ return fastNPP::Sub_32f_C1R_Ctx(a, b); }, false); + [](const Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { + return fastNPP::Sub_32f_C1R_Ctx(a, sa, b, sb, sz); }, false); bad += twoImgC1("Mul_32f_C1R", nppiMul_32f_C1R_Ctx, - [](const fk::Ptr2D& a, const fk::Ptr2D& b){ return fastNPP::Mul_32f_C1R_Ctx(a, b); }, false); + [](const Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { + return fastNPP::Mul_32f_C1R_Ctx(a, sa, b, sb, sz); }, false); bad += twoImgC1("Div_32f_C1R", nppiDiv_32f_C1R_Ctx, - [](const fk::Ptr2D& a, const fk::Ptr2D& b){ return fastNPP::Div_32f_C1R_Ctx(a, b); }, true); + [](const Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { + return fastNPP::Div_32f_C1R_Ctx(a, sa, b, sb, sz); }, true); printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED"); return bad == 0 ? 0 : 1; } diff --git a/tests/morphology/fastNPP_morphology_test.cu b/tests/morphology/fastNPP_morphology_test.cu index 5f0bba3..aced0be 100644 --- a/tests/morphology/fastNPP_morphology_test.cu +++ b/tests/morphology/fastNPP_morphology_test.cu @@ -108,28 +108,68 @@ int launch() { int bad = 0; bad += verify("ErodeBorder_8u_C1R 3x3", 3, 3, 1, 1, nppiErodeBorder_8u_C1R_Ctx, - [](const Npp8u* s, Npp32s step, NppiSize sz, - Npp8u* d, Npp32s dstep, + [](const Npp8u* s, Npp32s sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { - fastNPP::ErodeBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); + int devID = 0; cudaGetDevice(&devID); + fk::Ptr2D fkSrc(reinterpret_cast(const_cast(s)), + static_cast(sz.width), static_cast(sz.height), + static_cast(sStep), fk::MemType::Device, devID); + fk::Ptr2D fkDst(reinterpret_cast(d), + static_cast(sz.width), static_cast(sz.height), + static_cast(dStep), fk::MemType::Device, devID); + fastNPP::ErodeBorder_8u_C1R_Ctx( + fk::PerThreadRead::build(fkSrc), + fk::PerThreadWrite::build(fkDst), + sz, mSz, anc, ctx); }); bad += verify("DilateBorder_8u_C1R 3x3", 3, 3, 1, 1, nppiDilateBorder_8u_C1R_Ctx, - [](const Npp8u* s, Npp32s step, NppiSize sz, - Npp8u* d, Npp32s dstep, + [](const Npp8u* s, Npp32s sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { - fastNPP::DilateBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); + int devID = 0; cudaGetDevice(&devID); + fk::Ptr2D fkSrc(reinterpret_cast(const_cast(s)), + static_cast(sz.width), static_cast(sz.height), + static_cast(sStep), fk::MemType::Device, devID); + fk::Ptr2D fkDst(reinterpret_cast(d), + static_cast(sz.width), static_cast(sz.height), + static_cast(dStep), fk::MemType::Device, devID); + fastNPP::DilateBorder_8u_C1R_Ctx( + fk::PerThreadRead::build(fkSrc), + fk::PerThreadWrite::build(fkDst), + sz, mSz, anc, ctx); }); bad += verify("ErodeBorder_8u_C1R 5x5", 5, 5, 2, 2, nppiErodeBorder_8u_C1R_Ctx, - [](const Npp8u* s, Npp32s step, NppiSize sz, - Npp8u* d, Npp32s dstep, + [](const Npp8u* s, Npp32s sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { - fastNPP::ErodeBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); + int devID = 0; cudaGetDevice(&devID); + fk::Ptr2D fkSrc(reinterpret_cast(const_cast(s)), + static_cast(sz.width), static_cast(sz.height), + static_cast(sStep), fk::MemType::Device, devID); + fk::Ptr2D fkDst(reinterpret_cast(d), + static_cast(sz.width), static_cast(sz.height), + static_cast(dStep), fk::MemType::Device, devID); + fastNPP::ErodeBorder_8u_C1R_Ctx( + fk::PerThreadRead::build(fkSrc), + fk::PerThreadWrite::build(fkDst), + sz, mSz, anc, ctx); }); bad += verify("DilateBorder_8u_C1R 5x5", 5, 5, 2, 2, nppiDilateBorder_8u_C1R_Ctx, - [](const Npp8u* s, Npp32s step, NppiSize sz, - Npp8u* d, Npp32s dstep, + [](const Npp8u* s, Npp32s sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, NppiSize mSz, NppiPoint anc, NppStreamContext ctx) { - fastNPP::DilateBorder_8u_C1R_Ctx(s, step, sz, d, dstep, mSz, anc, ctx); }); + int devID = 0; cudaGetDevice(&devID); + fk::Ptr2D fkSrc(reinterpret_cast(const_cast(s)), + static_cast(sz.width), static_cast(sz.height), + static_cast(sStep), fk::MemType::Device, devID); + fk::Ptr2D fkDst(reinterpret_cast(d), + static_cast(sz.width), static_cast(sz.height), + static_cast(dStep), fk::MemType::Device, devID); + fastNPP::DilateBorder_8u_C1R_Ctx( + fk::PerThreadRead::build(fkSrc), + fk::PerThreadWrite::build(fkDst), + sz, mSz, anc, ctx); }); printf("%s\n", bad == 0 ? "ALL PASS" : "FAILURES DETECTED"); return bad == 0 ? 0 : 1; } From ef4235e4ba17c786888b5d895ef46e3a15bad829 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 1 Aug 2026 17:57:45 +0000 Subject: [PATCH 5/5] Refactor two-image fastNPP functions to take Read IOps with Tuple-based BackIOp --- include/fast_npp.h | 133 +++++++++++------- .../arithmetic/fastNPP_absdiff_shift_test.cu | 18 ++- tests/arithmetic/fastNPP_two_image_test.cu | 20 +-- 3 files changed, 101 insertions(+), 70 deletions(-) diff --git a/include/fast_npp.h b/include/fast_npp.h index 89ba1b0..58d5f09 100644 --- a/include/fast_npp.h +++ b/include/fast_npp.h @@ -93,64 +93,89 @@ 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. 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). + namespace detail { + template + struct DualSourceReadBack { + static_assert(fk::isTuple_v && 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 && fk::isAnyCompleteReadType, + "Both elements of the BackIOp Tuple must be complete Read IOps"); + using SelfType = DualSourceReadBack; + public: + FK_STATIC_STRUCT(DualSourceReadBack, SelfType) + using Parent = fk::ReadBackOperation, + DualSourceReadBack>; + 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) ===== - // Returns a DualSourceRead IOp (ReadType) that reads two source images and - // applies the bitwise operation element-wise. NPP pointer parameters are - // converted to fk:: types internally. The returned IOp can be composed in - // an executeOperations chain. -#define FASTNPP_DEFINE_TWO_IMAGE_BW(NPPNAME, T, NPP_T, FKLOP) \ - inline auto NPPNAME(const NPP_T* pSrc1, Npp32s nSrc1Step, \ - const NPP_T* pSrc2, Npp32s nSrc2Step, \ - NppiSize oSizeROI) { \ - int deviceID{ 0 }; \ - gpuErrchk(cudaGetDevice(&deviceID)); \ - const fk::Ptr2D fkSrc1( \ - reinterpret_cast(const_cast(pSrc1)), \ - static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ - static_cast(nSrc1Step), fk::MemType::Device, deviceID); \ - const fk::Ptr2D fkSrc2( \ - reinterpret_cast(const_cast(pSrc2)), \ - static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ - static_cast(nSrc2Step), fk::MemType::Device, deviceID); \ - return fk::DualSourceRead::build(fkSrc2, fkSrc1) \ - .then(fk::FKLOP::build()); \ - } - FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C1R_Ctx, uchar, Npp8u, BwAnd) - FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C3R_Ctx, uchar3, Npp8u, BwAnd) - FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C1R_Ctx, uchar, Npp8u, BwOr) - FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C3R_Ctx, uchar3, Npp8u, BwOr) - FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C1R_Ctx, uchar, Npp8u, BwXor) - FASTNPP_DEFINE_TWO_IMAGE_BW(Xor_8u_C3R_Ctx, uchar3, Npp8u, BwXor) + // Accepts two caller-supplied Read IOps instead of raw NPP pointers. + // detail::DualSourceReadBack reads both sources into an fk::Tuple, + // 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 \ + inline auto NPPNAME(const ReadIOp1& src1, const ReadIOp2& src2) { \ + return detail::DualSourceReadBack>::build(src2, src1) \ + .then(fk::FKLOP::build()); \ + } + FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C1R_Ctx, uchar, BwAnd) + FASTNPP_DEFINE_TWO_IMAGE_BW(And_8u_C3R_Ctx, uchar3, BwAnd) + FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C1R_Ctx, uchar, BwOr) + FASTNPP_DEFINE_TWO_IMAGE_BW(Or_8u_C3R_Ctx, uchar3, BwOr) + 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, NPP_T, FKLOP) \ - inline auto NPPNAME(const NPP_T* pSrc1, Npp32s nSrc1Step, \ - const NPP_T* pSrc2, Npp32s nSrc2Step, \ - NppiSize oSizeROI) { \ - int deviceID{ 0 }; \ - gpuErrchk(cudaGetDevice(&deviceID)); \ - const fk::Ptr2D fkSrc1( \ - reinterpret_cast(const_cast(pSrc1)), \ - static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ - static_cast(nSrc1Step), fk::MemType::Device, deviceID); \ - const fk::Ptr2D fkSrc2( \ - reinterpret_cast(const_cast(pSrc2)), \ - static_cast(oSizeROI.width), static_cast(oSizeROI.height), \ - static_cast(nSrc2Step), fk::MemType::Device, deviceID); \ - return fk::DualSourceRead::build(fkSrc2, fkSrc1) \ - .then(fk::FKLOP::build()); \ + // Same IOp-based composition as the bitwise two-image ops above. +#define FASTNPP_DEFINE_TWO_IMAGE(NPPNAME, T, FKLOP) \ + template \ + inline auto NPPNAME(const ReadIOp1& src1, const ReadIOp2& src2) { \ + return detail::DualSourceReadBack>::build(src2, src1) \ + .then(fk::FKLOP::build()); \ } - FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C1R_Ctx, float, Npp32f, Add) - FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C3R_Ctx, float3, Npp32f, Add) - FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C1R_Ctx, float, Npp32f, Sub) - FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C3R_Ctx, float3, Npp32f, Sub) - FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C1R_Ctx, float, Npp32f, Mul) - FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C3R_Ctx, float3, Npp32f, Mul) - FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C1R_Ctx, float, Npp32f, Div) - FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C3R_Ctx, float3, Npp32f, Div) + FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C1R_Ctx, float, Add) + FASTNPP_DEFINE_TWO_IMAGE(Add_32f_C3R_Ctx, float3, Add) + FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C1R_Ctx, float, Sub) + FASTNPP_DEFINE_TWO_IMAGE(Sub_32f_C3R_Ctx, float3, Sub) + FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C1R_Ctx, float, Mul) + FASTNPP_DEFINE_TWO_IMAGE(Mul_32f_C3R_Ctx, float3, Mul) + FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C1R_Ctx, float, Div) + FASTNPP_DEFINE_TWO_IMAGE(Div_32f_C3R_Ctx, float3, Div) // ===== Bitwise operations ===== // AndC / OrC / XorC with a constant, and Not (no constant). Integer types, // C1 / C3 / C4. Each maps the exact NPP name onto an FKL bitwise functor. diff --git a/tests/arithmetic/fastNPP_absdiff_shift_test.cu b/tests/arithmetic/fastNPP_absdiff_shift_test.cu index 3be8174..7b1f2a4 100644 --- a/tests/arithmetic/fastNPP_absdiff_shift_test.cu +++ b/tests/arithmetic/fastNPP_absdiff_shift_test.cu @@ -112,9 +112,16 @@ int bwTwoImg(const char* label, NppFn nppFn, FKLChainFn chainFn) { nppFn(d1, pitch, d2, pitch, dd, pitch, {W, H}, makeCtx()); cudaDeviceSynchronize(); cudaMemcpy2D(ref.data(), rb, dd, pitch, rb, H, cudaMemcpyDeviceToHost); + + int devID = 0; + cudaGetDevice(&devID); + fk::Ptr2D f1(d1, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID); + fk::Ptr2D f2(d2, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID); + fk::Stream st; fk::executeOperations>(st, - chainFn(d1, pitch, d2, pitch, NppiSize{W, H}), + chainFn(fk::PerThreadRead::build(f1), + fk::PerThreadRead::build(f2)), fk::PerThreadWrite::build(d)); st.sync(); cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); @@ -132,14 +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 Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { - return fastNPP::And_8u_C1R_Ctx(a, sa, b, sb, sz); }); + [](const auto& a, const auto& b) { return fastNPP::And_8u_C1R_Ctx(a, b); }); bad += bwTwoImg("Or_8u_C1R", nppiOr_8u_C1R_Ctx, - [](const Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { - return fastNPP::Or_8u_C1R_Ctx(a, sa, b, sb, sz); }); + [](const auto& a, const auto& b) { return fastNPP::Or_8u_C1R_Ctx(a, b); }); bad += bwTwoImg("Xor_8u_C1R", nppiXor_8u_C1R_Ctx, - [](const Npp8u* a, Npp32s sa, const Npp8u* b, Npp32s sb, NppiSize sz) { - return fastNPP::Xor_8u_C1R_Ctx(a, sa, b, sb, sz); }); + [](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; } diff --git a/tests/arithmetic/fastNPP_two_image_test.cu b/tests/arithmetic/fastNPP_two_image_test.cu index 789486b..f01a7e7 100644 --- a/tests/arithmetic/fastNPP_two_image_test.cu +++ b/tests/arithmetic/fastNPP_two_image_test.cu @@ -64,9 +64,15 @@ int twoImgC1(const char* label, NppFn nppFn, FKLChainFn chainFn, bool avoidZero) cudaDeviceSynchronize(); cudaMemcpy2D(ref.data(), rb, dd, pitch, rb, H, cudaMemcpyDeviceToHost); + int devID = 0; + cudaGetDevice(&devID); + fk::Ptr2D f1(d1, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID); + fk::Ptr2D f2(d2, (uint)W, (uint)H, (uint)pitch, fk::MemType::Device, devID); + fk::Stream st; fk::executeOperations>(st, - chainFn(d1, pitch, d2, pitch, NppiSize{W, H}), + chainFn(fk::PerThreadRead::build(f1), + fk::PerThreadRead::build(f2)), fk::PerThreadWrite::build(d)); st.sync(); cudaMemcpy2D(fkl.data(), rb, d.ptr().data, pitch, rb, H, cudaMemcpyDeviceToHost); @@ -87,17 +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 Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { - return fastNPP::Add_32f_C1R_Ctx(a, sa, b, sb, sz); }, 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 Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { - return fastNPP::Sub_32f_C1R_Ctx(a, sa, b, sb, sz); }, 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 Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { - return fastNPP::Mul_32f_C1R_Ctx(a, sa, b, sb, sz); }, 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 Npp32f* a, Npp32s sa, const Npp32f* b, Npp32s sb, NppiSize sz) { - return fastNPP::Div_32f_C1R_Ctx(a, sa, b, sb, sz); }, 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; }