diff --git a/include/fused_kernel/algorithms/image_processing/image_processing.h b/include/fused_kernel/algorithms/image_processing/image_processing.h index be21be3b..28eaacf0 100644 --- a/include/fused_kernel/algorithms/image_processing/image_processing.h +++ b/include/fused_kernel/algorithms/image_processing/image_processing.h @@ -23,7 +23,9 @@ #include #include #include +#include #include +#include #include #endif // FK_IMAGE_PROCESSING diff --git a/include/fused_kernel/algorithms/image_processing/linear_filter.h b/include/fused_kernel/algorithms/image_processing/linear_filter.h index fc903864..1fbc3c19 100644 --- a/include/fused_kernel/algorithms/image_processing/linear_filter.h +++ b/include/fused_kernel/algorithms/image_processing/linear_filter.h @@ -16,6 +16,7 @@ #define FK_LINEAR_FILTER_DPP_H #include +#include #include #include #include @@ -31,19 +32,18 @@ template struct LinearFilterDPPDetails { using ValueType = T; - static constexpr int TILE_WIDTH = TILE_W; - static constexpr int TILE_HEIGHT = TILE_H; - static constexpr int MAX_KERNEL_WIDTH = MAX_KERNEL_W; - static constexpr int MAX_KERNEL_HEIGHT = MAX_KERNEL_H; - static constexpr int MAX_HALO_WIDTH = TILE_W + MAX_KERNEL_W - 1; - static constexpr int MAX_HALO_HEIGHT = TILE_H + MAX_KERNEL_H - 1; - - static_assert(TILE_W > 0 && TILE_H > 0, - "Linear-filter tile dimensions must be positive"); - static_assert(MAX_KERNEL_W > 0 && MAX_KERNEL_H > 0, - "Linear-filter max kernel dimensions must be positive"); - static_assert(TILE_W * TILE_H <= 1024, - "Linear-filter tile exceeds CUDA block size"); + using NeighborhoodPolicy = NeighborhoodDPPPolicy< + T, TILE_W, TILE_H, MAX_KERNEL_W, MAX_KERNEL_H>; + static constexpr int TILE_WIDTH = NeighborhoodPolicy::TILE_WIDTH; + static constexpr int TILE_HEIGHT = NeighborhoodPolicy::TILE_HEIGHT; + static constexpr int MAX_KERNEL_WIDTH = + NeighborhoodPolicy::MAX_WINDOW_WIDTH; + static constexpr int MAX_KERNEL_HEIGHT = + NeighborhoodPolicy::MAX_WINDOW_HEIGHT; + static constexpr int MAX_HALO_WIDTH = + NeighborhoodPolicy::MAX_HALO_WIDTH; + static constexpr int MAX_HALO_HEIGHT = + NeighborhoodPolicy::MAX_HALO_HEIGHT; int width; int height; @@ -54,14 +54,10 @@ struct LinearFilterDPPDetails { FK_HOST_DEVICE_FUSE bool valid( const LinearFilterDPPDetails& details) { - return details.width > 0 && details.height > 0 && - details.kernelWidth > 0 && details.kernelHeight > 0 && - details.kernelWidth <= MAX_KERNEL_W && - details.kernelHeight <= MAX_KERNEL_H && - details.anchorX >= 0 && - details.anchorX < details.kernelWidth && - details.anchorY >= 0 && - details.anchorY < details.kernelHeight; + return NeighborhoodDPPStage::valid( + details.width, details.height, + details.kernelWidth, details.kernelHeight, + details.anchorX, details.anchorY, false); } }; @@ -96,15 +92,8 @@ struct LinearFilterDPP { private: using SelfType = LinearFilterDPP; using T = typename DPPDetails::ValueType; - - template - FK_HOST_STATIC T readReplicate(const DPPDetails& details, - const ImageRead& image, - int x, int y) { - x = x < 0 ? 0 : (x >= details.width ? details.width - 1 : x); - y = y < 0 ? 0 : (y >= details.height ? details.height - 1 : y); - return ImageRead::Operation::exec(Point{x, y, 0}, image); - } + using Stage = NeighborhoodDPPStage< + typename DPPDetails::NeighborhoodPolicy>; public: FK_STATIC_STRUCT(LinearFilterDPP, SelfType) @@ -133,8 +122,8 @@ struct LinearFilterDPP { T accumulator{}; for (int ky = 0; ky < details.kernelHeight; ++ky) { for (int kx = 0; kx < details.kernelWidth; ++kx) { - const T value = readReplicate( - details, image, + const T value = Stage::readReplicate( + details.width, details.height, image, ox + kx - details.anchorX, oy + ky - details.anchorY); const T coefficient = @@ -159,15 +148,8 @@ struct LinearFilterDPP { private: using SelfType = LinearFilterDPP; using T = typename DPPDetails::ValueType; - - template - FK_DEVICE_STATIC T readReplicate(const DPPDetails& details, - const ImageRead& image, - int x, int y) { - x = x < 0 ? 0 : (x >= details.width ? details.width - 1 : x); - y = y < 0 ? 0 : (y >= details.height ? details.height - 1 : y); - return ImageRead::Operation::exec(Point{x, y, 0}, image); - } + using Stage = NeighborhoodDPPStage< + typename DPPDetails::NeighborhoodPolicy>; public: FK_STATIC_STRUCT(LinearFilterDPP, SelfType) @@ -199,20 +181,13 @@ struct LinearFilterDPP { const int threads = blockDim.x * blockDim.y; const int haloWidth = DPPDetails::TILE_WIDTH + details.kernelWidth - 1; - const int haloHeight = DPPDetails::TILE_HEIGHT + - details.kernelHeight - 1; const int tileX = blockIdx.x * DPPDetails::TILE_WIDTH; const int tileY = blockIdx.y * DPPDetails::TILE_HEIGHT; - const int haloX = tileX - details.anchorX; - const int haloY = tileY - details.anchorY; - for (int index = tid; index < haloWidth * haloHeight; - index += threads) { - const int y = index / haloWidth; - const int x = index % haloWidth; - halo[index] = readReplicate( - details, image, haloX + x, haloY + y); - } + Stage::stageReplicate( + details.width, details.height, + details.kernelWidth, details.kernelHeight, + details.anchorX, details.anchorY, image, halo); for (int index = tid; index < details.kernelWidth * details.kernelHeight; index += threads) { diff --git a/include/fused_kernel/algorithms/image_processing/median_filter.h b/include/fused_kernel/algorithms/image_processing/median_filter.h new file mode 100644 index 00000000..e1c70e79 --- /dev/null +++ b/include/fused_kernel/algorithms/image_processing/median_filter.h @@ -0,0 +1,272 @@ +/* Copyright 2025-2026 Oscar Amoros Huguet, Johnny Nunez + + 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. */ + +#ifndef FK_MEDIAN_FILTER_DPP_H +#define FK_MEDIAN_FILTER_DPP_H + +#include +#include +#include +#include +#include +#include + +#include + +namespace fk { + +/** Fixed-capacity register value handed from a neighbourhood DPP to a + * single-thread selection Operation. Only values[0..count) are live. */ +template +struct NeighborhoodWindow { + using ValueType = T; + static constexpr int MAX_CAPACITY = CAPACITY; + T values[CAPACITY]; + int count; +}; + +/** Parameterless Unary Operation selecting the upper median (count/2). */ +template +struct MedianWindowSelect { +private: + using Window = NeighborhoodWindow; + using Parent = UnaryOperation>; + using SelfType = MedianWindowSelect; + +public: + FK_STATIC_STRUCT(MedianWindowSelect, SelfType) + DECLARE_UNARY_PARENT + + FK_HOST_DEVICE_FUSE OutputType exec(InputType window) { + for (int i = 1; i < window.count; ++i) { + const T key = window.values[i]; + int j = i - 1; + while (j >= 0 && window.values[j] > key) { + window.values[j + 1] = window.values[j]; + --j; + } + window.values[j + 1] = key; + } + return window.values[window.count / 2]; + } +}; + +/** Alternate selection Operation used to prove orchestration is generic. */ +template +struct MinimumWindowSelect { +private: + using Window = NeighborhoodWindow; + using Parent = UnaryOperation>; + using SelfType = MinimumWindowSelect; + +public: + FK_STATIC_STRUCT(MinimumWindowSelect, SelfType) + DECLARE_UNARY_PARENT + + FK_HOST_DEVICE_FUSE OutputType exec(const InputType window) { + T result = window.values[0]; + for (int i = 1; i < window.count; ++i) + result = window.values[i] < result ? window.values[i] : result; + return result; + } +}; + +template +struct MedianFilterDPPDetails { + using ValueType = T; + using NeighborhoodPolicy = NeighborhoodDPPPolicy< + T, TILE_W, TILE_H, MAX_WINDOW_W, MAX_WINDOW_H>; + static constexpr int TILE_WIDTH = NeighborhoodPolicy::TILE_WIDTH; + static constexpr int TILE_HEIGHT = NeighborhoodPolicy::TILE_HEIGHT; + static constexpr int MAX_WINDOW_WIDTH = + NeighborhoodPolicy::MAX_WINDOW_WIDTH; + static constexpr int MAX_WINDOW_HEIGHT = + NeighborhoodPolicy::MAX_WINDOW_HEIGHT; + static constexpr int MAX_HALO_WIDTH = + NeighborhoodPolicy::MAX_HALO_WIDTH; + static constexpr int MAX_HALO_HEIGHT = + NeighborhoodPolicy::MAX_HALO_HEIGHT; + static constexpr int WINDOW_CAPACITY = + MAX_WINDOW_WIDTH * MAX_WINDOW_HEIGHT; + + int width; + int height; + int windowWidth; + int windowHeight; + int anchorX; + int anchorY; + + // Even windows are intentionally unsupported: there is no unique median. + FK_HOST_DEVICE_FUSE bool valid(const MedianFilterDPPDetails& details) { + return NeighborhoodDPPStage::valid( + details.width, details.height, + details.windowWidth, details.windowHeight, + details.anchorX, details.anchorY, true); + } +}; + +template +struct MedianFilterDPP; + +template +struct MedianFilterDPP { +private: + using SelfType = MedianFilterDPP; + using T = typename DPPDetails::ValueType; + using Stage = NeighborhoodDPPStage< + typename DPPDetails::NeighborhoodPolicy>; + using Window = NeighborhoodWindow; + +public: + FK_STATIC_STRUCT(MedianFilterDPP, SelfType) + static constexpr ParArch PAR_ARCH = ParArch::CPU; + + template + FK_HOST_STATIC void exec(const DPPDetails& details, + const ReadIOp& input, + const WriteIOp& output, + const SelectionIOp& selection) { + static_assert(isAnyCompleteReadType, + "MedianFilterDPP needs a complete Read IOp"); + static_assert(isAnyWriteType, + "MedianFilterDPP needs a Write IOp"); + static_assert(SelectionIOp::template is, + "Selection must be a Unary IOp"); + static_assert(std::is_same_v, + "Selection input must be this DPP's window type"); + static_assert(std::is_same_v, + "Selection output must match the image value type"); + if (!DPPDetails::valid(details)) return; + + for (int oy = 0; oy < details.height; ++oy) { + for (int ox = 0; ox < details.width; ++ox) { + Window window{}; + for (int wy = 0; wy < details.windowHeight; ++wy) { + for (int wx = 0; wx < details.windowWidth; ++wx) { + window.values[window.count++] = Stage::readReplicate( + details.width, details.height, input, + ox + wx - details.anchorX, + oy + wy - details.anchorY); + } + } + const T selected = window | selection; + WriteIOp::Operation::exec( + Point{ox, oy, 0}, selected, output); + } + } + } +}; + +#if defined(__NVCC__) +template +struct MedianFilterDPP { +private: + using SelfType = MedianFilterDPP; + using T = typename DPPDetails::ValueType; + using Stage = NeighborhoodDPPStage< + typename DPPDetails::NeighborhoodPolicy>; + using Window = NeighborhoodWindow; + +public: + FK_STATIC_STRUCT(MedianFilterDPP, SelfType) + static constexpr ParArch PAR_ARCH = ParArch::GPU_NVIDIA; + + template + FK_DEVICE_STATIC void exec(const DPPDetails& details, + const ReadIOp& input, + const WriteIOp& output, + const SelectionIOp& selection) { + static_assert(isAnyCompleteReadType, + "MedianFilterDPP needs a complete Read IOp"); + static_assert(isAnyWriteType, + "MedianFilterDPP needs a Write IOp"); + static_assert(SelectionIOp::template is, + "Selection must be a Unary IOp"); + static_assert(std::is_same_v, + "Selection input must be this DPP's window type"); + static_assert(std::is_same_v, + "Selection output must match the image value type"); + + __shared__ T halo[DPPDetails::MAX_HALO_WIDTH * + DPPDetails::MAX_HALO_HEIGHT]; + Stage::stageReplicate( + details.width, details.height, + details.windowWidth, details.windowHeight, + details.anchorX, details.anchorY, input, halo); + __syncthreads(); + + const int tileX = blockIdx.x * DPPDetails::TILE_WIDTH; + const int tileY = blockIdx.y * DPPDetails::TILE_HEIGHT; + const int ox = tileX + threadIdx.x; + const int oy = tileY + threadIdx.y; + const int haloWidth = DPPDetails::TILE_WIDTH + + details.windowWidth - 1; + if (threadIdx.x < DPPDetails::TILE_WIDTH && + threadIdx.y < DPPDetails::TILE_HEIGHT && + ox < details.width && oy < details.height) { + Window window{}; + for (int wy = 0; wy < details.windowHeight; ++wy) + for (int wx = 0; wx < details.windowWidth; ++wx) + window.values[window.count++] = halo[ + (threadIdx.y + wy) * haloWidth + + threadIdx.x + wx]; + const T selected = window | selection; + WriteIOp::Operation::exec( + Point{ox, oy, 0}, selected, output); + } + } +}; + +template +__global__ void medianFilterDPPKernel(const DPPDetails details, + const ReadIOp input, + const WriteIOp output, + const SelectionIOp selection) { + MedianFilterDPP::exec( + details, input, output, selection); +} + +template +inline bool executeMedianFilter(const DPPDetails& details, + const ReadIOp& input, + const WriteIOp& output, + const SelectionIOp& selection, + Stream_& stream) { + if (!DPPDetails::valid(details)) return false; + const dim3 block(DPPDetails::TILE_WIDTH, + DPPDetails::TILE_HEIGHT, 1); + const dim3 grid((details.width + DPPDetails::TILE_WIDTH - 1) / + DPPDetails::TILE_WIDTH, + (details.height + DPPDetails::TILE_HEIGHT - 1) / + DPPDetails::TILE_HEIGHT, + 1); + medianFilterDPPKernel<<>>( + details, input, output, selection); + gpuErrchk(cudaGetLastError()); + return true; +} +#endif // defined(__NVCC__) + +} // namespace fk + +#endif // FK_MEDIAN_FILTER_DPP_H diff --git a/include/fused_kernel/algorithms/image_processing/neighborhood.h b/include/fused_kernel/algorithms/image_processing/neighborhood.h new file mode 100644 index 00000000..136bd525 --- /dev/null +++ b/include/fused_kernel/algorithms/image_processing/neighborhood.h @@ -0,0 +1,106 @@ +/* Copyright 2026 Oscar Amoros Huguet, Johnny Nunez + + 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. */ + +#ifndef FK_NEIGHBORHOOD_DPP_H +#define FK_NEIGHBORHOOD_DPP_H + +#include +#include + +namespace fk { + +template +struct NeighborhoodDPPPolicy { + using ValueType = T; + static constexpr int TILE_WIDTH = TILE_W; + static constexpr int TILE_HEIGHT = TILE_H; + static constexpr int MAX_WINDOW_WIDTH = MAX_WINDOW_W; + static constexpr int MAX_WINDOW_HEIGHT = MAX_WINDOW_H; + static constexpr int MAX_HALO_WIDTH = TILE_W + MAX_WINDOW_W - 1; + static constexpr int MAX_HALO_HEIGHT = TILE_H + MAX_WINDOW_H - 1; + + static_assert(TILE_W > 0 && TILE_H > 0, + "Neighbourhood tile dimensions must be positive"); + static_assert(MAX_WINDOW_W > 0 && MAX_WINDOW_H > 0, + "Neighbourhood max window dimensions must be positive"); + static_assert(TILE_W * TILE_H <= 1024, + "Neighbourhood tile exceeds CUDA block size"); +}; + +template +struct NeighborhoodDPPStage { +private: + using SelfType = NeighborhoodDPPStage; + +public: + using T = typename Policy::ValueType; + FK_STATIC_STRUCT(NeighborhoodDPPStage, SelfType) + + FK_HOST_DEVICE_FUSE bool valid(const int width, const int height, + const int windowWidth, + const int windowHeight, + const int anchorX, const int anchorY, + const bool requireOdd) { + return width > 0 && height > 0 && + windowWidth > 0 && windowHeight > 0 && + windowWidth <= Policy::MAX_WINDOW_WIDTH && + windowHeight <= Policy::MAX_WINDOW_HEIGHT && + anchorX >= 0 && anchorX < windowWidth && + anchorY >= 0 && anchorY < windowHeight && + (!requireOdd || ((windowWidth & 1) && (windowHeight & 1))); + } + + template + FK_HOST_DEVICE_STATIC T readReplicate(const int width, + const int height, + const ReadIOp& input, + int x, int y) { + x = x < 0 ? 0 : (x >= width ? width - 1 : x); + y = y < 0 ? 0 : (y >= height ? height - 1 : y); + return ReadIOp::Operation::exec(Point{x, y, 0}, input); + } + +#if defined(__NVCC__) + template + FK_DEVICE_STATIC void stageReplicate(const int width, + const int height, + const int windowWidth, + const int windowHeight, + const int anchorX, + const int anchorY, + const ReadIOp& input, + T* halo) { + const int haloWidth = Policy::TILE_WIDTH + windowWidth - 1; + const int haloHeight = Policy::TILE_HEIGHT + windowHeight - 1; + const int tileX = blockIdx.x * Policy::TILE_WIDTH; + const int tileY = blockIdx.y * Policy::TILE_HEIGHT; + const int originX = tileX - anchorX; + const int originY = tileY - anchorY; + const int tid = threadIdx.y * blockDim.x + threadIdx.x; + const int threads = blockDim.x * blockDim.y; + for (int index = tid; index < haloWidth * haloHeight; + index += threads) { + const int y = index / haloWidth; + const int x = index % haloWidth; + halo[index] = readReplicate( + width, height, input, originX + x, originY + y); + } + } +#endif +}; + +} // namespace fk + +#endif // FK_NEIGHBORHOOD_DPP_H diff --git a/tests/image_processing/test_median_filter_dpp.h b/tests/image_processing/test_median_filter_dpp.h new file mode 100644 index 00000000..405946f3 --- /dev/null +++ b/tests/image_processing/test_median_filter_dpp.h @@ -0,0 +1,174 @@ +/* Copyright 2026 Johnny Nunez + + 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. */ + +#include + +#include +#include +#include + +#include +#include +#include +#include + +using namespace fk; + +namespace { + +using Details = MedianFilterDPPDetails; +using Window = NeighborhoodWindow; +using MedianSelection = decltype(MedianWindowSelect::build()); +using MinSelection = decltype(MinimumWindowSelect::build()); + +enum class SelectionKind { MEDIAN, MINIMUM }; + +float inputValue(const int x, const int y) { + return static_cast((x * 11 + y * 7) % 29 - 14) * 0.125f; +} + +std::vector oracle(const std::vector& input, + const Details& details, + const SelectionKind kind, + const bool fused) { + std::vector output(details.width * details.height); + std::vector values; + values.reserve(details.windowWidth * details.windowHeight); + for (int oy = 0; oy < details.height; ++oy) { + for (int ox = 0; ox < details.width; ++ox) { + values.clear(); + for (int wy = 0; wy < details.windowHeight; ++wy) { + for (int wx = 0; wx < details.windowWidth; ++wx) { + int sx = ox + wx - details.anchorX; + int sy = oy + wy - details.anchorY; + sx = sx < 0 ? 0 : (sx >= details.width + ? details.width - 1 : sx); + sy = sy < 0 ? 0 : (sy >= details.height + ? details.height - 1 : sy); + float value = input[sy * details.width + sx]; + if (fused) value = value * 2.f + 1.f; + values.push_back(value); + } + } + std::sort(values.begin(), values.end()); + float selected = kind == SelectionKind::MEDIAN + ? values[values.size() / 2] : values.front(); + output[oy * details.width + ox] = + fused ? selected * 0.5f : selected; + } + } + return output; +} + +bool compare(const std::vector& output, + const std::vector& expected, + const char* label) { + for (size_t index = 0; index < output.size(); ++index) { + if (std::fabs(output[index] - expected[index]) > 1e-6f) { + std::printf("%s index=%zu got=%g expected=%g\n", + label, index, output[index], expected[index]); + return false; + } + } + return true; +} + +template +bool runCase(const Details& details, const Selection& selection, + const SelectionKind kind, const bool fused, + const char* label) { + std::vector input(details.width * details.height); + for (int y = 0; y < details.height; ++y) + for (int x = 0; x < details.width; ++x) + input[y * details.width + x] = inputValue(x, y); + const auto expected = oracle(input, details, kind, fused); + std::vector cpuOutput(input.size(), -999.f); + const RawPtr inputPtr{ + input.data(), PtrDims( + details.width, details.height, + details.width * sizeof(float))}; + const RawPtr outputPtr{ + cpuOutput.data(), PtrDims( + details.width, details.height, + details.width * sizeof(float))}; + const auto readBase = PerThreadRead::build(inputPtr); + const auto writeBase = PerThreadWrite::build(outputPtr); + if (fused) { + const auto read = readBase.then(Mul::build(2.f)) + .then(Add::build(1.f)); + const auto write = Mul::build(0.5f).then(writeBase); + MedianFilterDPP::exec( + details, read, write, selection); + } else { + MedianFilterDPP::exec( + details, readBase, writeBase, selection); + } + if (!compare(cpuOutput, expected, label)) return false; + +#if defined(__NVCC__) + Ptr2D gpuInput(details.width, details.height); + Ptr2D gpuOutput(details.width, details.height); + for (int y = 0; y < details.height; ++y) + for (int x = 0; x < details.width; ++x) + gpuInput.at(Point{x, y, 0}) = input[y * details.width + x]; + Stream stream; + gpuInput.upload(stream); + const auto gpuReadBase = + PerThreadRead::build(gpuInput); + const auto gpuWriteBase = + PerThreadWrite::build(gpuOutput); + bool launched = false; + if (fused) { + const auto read = gpuReadBase.then(Mul::build(2.f)) + .then(Add::build(1.f)); + const auto write = Mul::build(0.5f).then(gpuWriteBase); + launched = executeMedianFilter( + details, read, write, selection, stream); + } else { + launched = executeMedianFilter( + details, gpuReadBase, gpuWriteBase, selection, stream); + } + if (!launched) return false; + gpuOutput.download(stream); + stream.sync(); + std::vector result(input.size()); + for (int y = 0; y < details.height; ++y) + for (int x = 0; x < details.width; ++x) + result[y * details.width + x] = + gpuOutput.at(Point{x, y, 0}); + if (!compare(result, expected, label)) return false; +#endif + return true; +} + +} // namespace + +int launch() { + const auto median = MedianWindowSelect::build(); + const auto minimum = MinimumWindowSelect::build(); + bool ok = true; + ok = runCase(Details{37, 19, 3, 3, 1, 1}, median, + SelectionKind::MEDIAN, false, "median3") && ok; + ok = runCase(Details{35, 17, 5, 5, 2, 2}, median, + SelectionKind::MEDIAN, false, "median5") && ok; + ok = runCase(Details{33, 21, 7, 3, 5, 1}, median, + SelectionKind::MEDIAN, true, "median7x3-fused") && ok; + ok = runCase(Details{31, 15, 3, 5, 0, 3}, minimum, + SelectionKind::MINIMUM, false, "minimum-selection") && ok; + + const Details evenWindow{16, 16, 4, 3, 1, 1}; + if (Details::valid(evenWindow)) ok = false; + if (ok) std::printf("MedianFilterDPP contracts: PASS\n"); + return ok ? 0 : -1; +}