From f5bbcab5f22eb63327e1feb7a136768c59d363bf Mon Sep 17 00:00:00 2001 From: Tomas Farago Date: Mon, 24 Jun 2019 09:12:41 +0200 Subject: Fix stripe_filter kernel Modify only all horizontal frequencies and take interleaved storage into account. --- src/kernels/.filter.cl.swp | Bin 0 -> 12288 bytes src/kernels/filter.cl | 2 +- 2 files changed, 1 insertion(+), 1 deletion(-) create mode 100644 src/kernels/.filter.cl.swp diff --git a/src/kernels/.filter.cl.swp b/src/kernels/.filter.cl.swp new file mode 100644 index 0000000..92485a6 Binary files /dev/null and b/src/kernels/.filter.cl.swp differ diff --git a/src/kernels/filter.cl b/src/kernels/filter.cl index 445a735..ed68e16 100644 --- a/src/kernels/filter.cl +++ b/src/kernels/filter.cl @@ -38,7 +38,7 @@ stripe_filter (global float *input, const int height = get_global_size(1); const int index = idy * get_global_size(0) + idx; - if (((idy < 1) || (idy >= height - 1)) && (idx > 0) && (idx < width - 1)) + if (idy == 0 && idx > 1) output[index] = 0.0f; else output[index] = input[index]; -- cgit v1.2.1 From acd91fe68874f5e535a62a41da9fd4553ebfb01f Mon Sep 17 00:00:00 2001 From: Tomas Farago Date: Mon, 24 Jun 2019 12:02:53 +0200 Subject: Add Gaussian window to the stripe filter --- src/kernels/.filter.cl.swp | Bin 12288 -> 0 bytes src/kernels/filter.cl | 14 ++++++++--- src/ufo-filter-stripes-task.c | 57 ++++++++++++++++++++++++++++++++++++++++++ 3 files changed, 67 insertions(+), 4 deletions(-) delete mode 100644 src/kernels/.filter.cl.swp diff --git a/src/kernels/.filter.cl.swp b/src/kernels/.filter.cl.swp deleted file mode 100644 index 92485a6..0000000 Binary files a/src/kernels/.filter.cl.swp and /dev/null differ diff --git a/src/kernels/filter.cl b/src/kernels/filter.cl index ed68e16..ddc91f6 100644 --- a/src/kernels/filter.cl +++ b/src/kernels/filter.cl @@ -30,16 +30,22 @@ filter (global float *input, kernel void stripe_filter (global float *input, - global float *output) + global float *output, + const float sigma) { const int idx = get_global_id(0); const int idy = get_global_id(1); const int width = get_global_size(0); const int height = get_global_size(1); const int index = idy * get_global_size(0) + idx; + /* Swap frequencies for interleaved complex input */ + /* No need to negate the second half of frequencies for Gaussian */ + const float x = idx >= width >> 1 ? (width >> 1) - (idx >> 1) : idx >> 1; + const float weight = exp (- x * x / (2 * sigma * sigma)); - if (idy == 0 && idx > 1) - output[index] = 0.0f; - else + if (idy == 0) { + output[index] = input[index] * weight; + } else { output[index] = input[index]; + } } diff --git a/src/ufo-filter-stripes-task.c b/src/ufo-filter-stripes-task.c index 3b6aebd..98dd852 100644 --- a/src/ufo-filter-stripes-task.c +++ b/src/ufo-filter-stripes-task.c @@ -28,6 +28,7 @@ struct _UfoFilterStripesTaskPrivate { + gfloat sigma; cl_kernel kernel; }; @@ -41,9 +42,12 @@ G_DEFINE_TYPE_WITH_CODE (UfoFilterStripesTask, ufo_filter_stripes_task, UFO_TYPE enum { PROP_0, + PROP_SIGMA, N_PROPERTIES }; +static GParamSpec *properties[N_PROPERTIES] = { NULL, }; + UfoNode * ufo_filter_stripes_task_new (void) { @@ -71,9 +75,11 @@ ufo_filter_stripes_task_process (UfoTask *task, UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_mem)); UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &out_mem)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_float), &priv->sigma)); profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL); + g_message ("%lu %lu %g", requisition->dims[0], requisition->dims[1], priv->sigma); return TRUE; } @@ -148,13 +154,63 @@ ufo_task_interface_init (UfoTaskIface *iface) iface->process = ufo_filter_stripes_task_process; } +static void +ufo_filter_stripes_task_set_property (GObject *object, + guint property_id, + const GValue *value, + GParamSpec *pspec) +{ + UfoFilterStripesTaskPrivate *priv = UFO_FILTER_STRIPES_TASK_GET_PRIVATE (object); + + switch (property_id) { + case PROP_SIGMA: + priv->sigma = g_value_get_float (value); + break; + default: + G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec); + break; + } +} + +static void +ufo_filter_stripes_task_get_property (GObject *object, + guint property_id, + GValue *value, + GParamSpec *pspec) +{ + UfoFilterStripesTaskPrivate *priv = UFO_FILTER_STRIPES_TASK_GET_PRIVATE (object); + + switch (property_id) { + case PROP_SIGMA: + g_value_set_float (value, priv->sigma); + break; + default: + G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec); + break; + } +} + static void ufo_filter_stripes_task_class_init (UfoFilterStripesTaskClass *klass) { GObjectClass *oclass; oclass = G_OBJECT_CLASS (klass); + oclass->finalize = ufo_filter_stripes_task_finalize; + oclass->set_property = ufo_filter_stripes_task_set_property; + oclass->get_property = ufo_filter_stripes_task_get_property; + + properties[PROP_SIGMA] = + g_param_spec_float ("sigma", + "Sigma of the gaussian window", + "Sigma of the gaussian window", + 0.0, G_MAXFLOAT, 1e-7, + G_PARAM_READWRITE); + + for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) + g_object_class_install_property (oclass, i, properties[i]); + g_type_class_add_private(klass, sizeof(UfoFilterStripesTaskPrivate)); } @@ -164,4 +220,5 @@ ufo_filter_stripes_task_init (UfoFilterStripesTask *self) UfoFilterStripesTaskPrivate *priv; self->priv = priv = UFO_FILTER_STRIPES_TASK_GET_PRIVATE (self); priv->kernel = NULL; + priv->sigma = 1e-7; } -- cgit v1.2.1 From 225446c5f06febc9e37764bae416c4bb32b4765a Mon Sep 17 00:00:00 2001 From: Tomas Farago Date: Tue, 25 Jun 2019 08:57:12 +0200 Subject: Document filter-stripes --- docs/filters.rst | 20 ++++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/docs/filters.rst b/docs/filters.rst index 96a40ea..d962b31 100644 --- a/docs/filters.rst +++ b/docs/filters.rst @@ -801,6 +801,26 @@ Frequency filtering Theta parameter of Faris-Byer filter. +Stripe filtering +---------------- + +.. gobj:class:: filter-stripes + + Filter vertical stripes. The input and output are in 2D frequency domain. + The filter multiplies horizontal frequencies (for frequency ky=0) with a + Gaussian profile centered at 0 frequency. + + Example usage:: + + $ ufo-launch read path=sino.tif ! fft dimensions=2 ! filter-stripes sigma=1 ! ifft dimensions=2 ! write filename=sino-filtered.tif + + .. gobj:prop:: sigma:float + + Filter strength, which is the sigma of the gaussian. Small values, e.g. + 1e-7 cause only the zero frequency to remain in the signal, i.e. + stronger filtering. + + 1D stripe filtering ------------------- -- cgit v1.2.1 From 2346d9f843abefd1a483041bc1b3e8663a310e3a Mon Sep 17 00:00:00 2001 From: Tomas Farago Date: Tue, 25 Jun 2019 09:16:57 +0200 Subject: Remove print --- src/ufo-filter-stripes-task.c | 1 - 1 file changed, 1 deletion(-) diff --git a/src/ufo-filter-stripes-task.c b/src/ufo-filter-stripes-task.c index 98dd852..7b07c5d 100644 --- a/src/ufo-filter-stripes-task.c +++ b/src/ufo-filter-stripes-task.c @@ -79,7 +79,6 @@ ufo_filter_stripes_task_process (UfoTask *task, profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL); - g_message ("%lu %lu %g", requisition->dims[0], requisition->dims[1], priv->sigma); return TRUE; } -- cgit v1.2.1