diff --git a/cmake/discover_tests.cmake b/cmake/discover_tests.cmake index 35cf702..8e9785a 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/include/fast_npp.h b/include/fast_npp.h index 4c673d1..58d5f09 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. + // 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) { \ + 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, 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); @@ -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. 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) ===== -#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()); \ + // 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) @@ -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& pSrc1, const fk::Ptr2D& pSrc2) { \ - return fk::DualSourceRead::build(pSrc2, pSrc1) \ - .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, Add) diff --git a/tests/arithmetic/fastNPP_absdiff_shift_test.cu b/tests/arithmetic/fastNPP_absdiff_shift_test.cu index 24f5e7e..7b1f2a4 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,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); - 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 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(s1, s2), + fk::executeOperations>(st, + 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); @@ -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& a, const fk::Ptr2D& 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& a, const fk::Ptr2D& 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& a, const fk::Ptr2D& 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; } diff --git a/tests/arithmetic/fastNPP_two_image_test.cu b/tests/arithmetic/fastNPP_two_image_test.cu index de8bd8e..f01a7e7 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,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 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(s1, s2), + 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); @@ -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& a, const fk::Ptr2D& 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& a, const fk::Ptr2D& 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& a, const fk::Ptr2D& 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& a, const fk::Ptr2D& 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; } diff --git a/tests/morphology/fastNPP_morphology_test.cu b/tests/morphology/fastNPP_morphology_test.cu new file mode 100644 index 0000000..aced0be --- /dev/null +++ b/tests/morphology/fastNPP_morphology_test.cu @@ -0,0 +1,175 @@ +/* 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); + + // Allocate device memory with a pitch aligned by CUDA + const int rb = W; + 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, pitchBytes * H); + cudaMemset(dFastDst, 0, pitchBytes * H); + + NppiSize srcSize{ W, H }; + NppiPoint srcOffset{ 0, 0 }; + NppiSize roiSize{ W, H }; + NppiSize maskSize{ mW, mH }; + NppiPoint anchor{ aX, aY }; + + // --- NPP reference --- + nppFn(dSrc, pitch, srcSize, srcOffset, + dNppDst, pitch, roiSize, + dmask, maskSize, anchor, + NPP_BORDER_REPLICATE, makeCtx()); + cudaDeviceSynchronize(); + cudaMemcpy2D(ref.data(), rb, dNppDst, pitchBytes, rb, H, cudaMemcpyDeviceToHost); + + // --- FastNPP (accepts NPP parameters; converts to fk:: internally) --- + fastNppFn(dSrc, pitch, srcSize, dFastDst, pitch, maskSize, anchor, makeCtx()); + cudaDeviceSynchronize(); + 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(dSrc); cudaFree(dNppDst); cudaFree(dFastDst); 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 Npp8u* s, Npp32s sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, + NppiSize mSz, NppiPoint anc, NppStreamContext 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 sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, + NppiSize mSz, NppiPoint anc, NppStreamContext 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 sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, + NppiSize mSz, NppiPoint anc, NppStreamContext 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 sStep, NppiSize sz, + Npp8u* d, Npp32s dStep, + NppiSize mSz, NppiPoint anc, NppStreamContext 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; +}