Skip to content

Add MedianFilter (window median) #278

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
morousg merged 2 commits into from
Aug 1, 2026
Merged

Add MedianFilter (window median) #278

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d08c51c
feat(image): add median filter DPP
johnnynunez Jun 28, 2026
9279bcb
Merge branch 'main' into feat/median-filter
morousg Aug 1, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions include/fused_kernel/algorithms/image_processing/image_processing.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,9 @@
#include <fused_kernel/algorithms/image_processing/resize.h>
#include <fused_kernel/algorithms/image_processing/saturate.h>
#include <fused_kernel/algorithms/image_processing/warping.h>
#include <fused_kernel/algorithms/image_processing/neighborhood.h>
#include <fused_kernel/algorithms/image_processing/linear_filter.h>
#include <fused_kernel/algorithms/image_processing/median_filter.h>
#include <fused_kernel/algorithms/image_processing/morphology.h>

#endif // FK_IMAGE_PROCESSING
79 changes: 27 additions & 52 deletions include/fused_kernel/algorithms/image_processing/linear_filter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@
#define FK_LINEAR_FILTER_DPP_H

#include <fused_kernel/algorithms/basic_ops/arithmetic.h>
#include <fused_kernel/algorithms/image_processing/neighborhood.h>
#include <fused_kernel/core/data/point.h>
#include <fused_kernel/core/data/tuple.h>
#include <fused_kernel/core/execution_model/operation_model/operation_model.h>
Expand All @@ -31,19 +32,18 @@ template <typename T, int TILE_W = 16, int TILE_H = 8,
int MAX_KERNEL_W = 7, int MAX_KERNEL_H = 7>
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;
Expand All @@ -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<NeighborhoodPolicy>::valid(
details.width, details.height,
details.kernelWidth, details.kernelHeight,
details.anchorX, details.anchorY, false);
}
};

Expand Down Expand Up @@ -96,15 +92,8 @@ struct LinearFilterDPP<ParArch::CPU, DPPDetails> {
private:
using SelfType = LinearFilterDPP<ParArch::CPU, DPPDetails>;
using T = typename DPPDetails::ValueType;

template <typename ImageRead>
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)
Expand Down Expand Up @@ -133,8 +122,8 @@ struct LinearFilterDPP<ParArch::CPU, DPPDetails> {
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 =
Expand All @@ -159,15 +148,8 @@ struct LinearFilterDPP<ParArch::GPU_NVIDIA, DPPDetails> {
private:
using SelfType = LinearFilterDPP<ParArch::GPU_NVIDIA, DPPDetails>;
using T = typename DPPDetails::ValueType;

template <typename ImageRead>
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)
Expand Down Expand Up @@ -199,20 +181,13 @@ struct LinearFilterDPP<ParArch::GPU_NVIDIA, DPPDetails> {
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) {
Expand Down
Loading