Add RemoveStripes filter

src/CMakeLists.txt

@@ -4,6 +4,7 @@ cmake_minimum_required(VERSION 2.6)
 set(ufofilter_SRCS
     ufo-anka-backproject-task.c
     ufo-anka-padding-task.c
+    ufo-remove-stripes-task.c
     )
 
 file(GLOB ufofilter_KERNELS "kernels/*.cl")

src/kernels/complex.cl

@@ -0,0 +1,46 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+/**
+  * multiply_frequencies_with_real_sym:
+  * @frequencies: complex Fourier transform frequencies with interleaved
+  * real/imaginary values
+  * @output: multiplication result
+  * @coefficients: first half of symmetric coefficients for the multiplication
+  * (size = width / 2 + 1)
+  *
+  * Multiply every row of @frequencies with @coefficients which are half the *real*
+  * width + 1, i.e. width = global size / 2 because of the complex numbers. This
+  * kernel takes advantage of symmetry and expects @frequencies to be ordered as
+  * [0, 1, ..., width / 2 - 1, -width / 2, ..., -1]. After width / 2 the @coefficients
+  * are mirrored.
+  */
+kernel void
+multiply_frequencies_with_real_sym (global float *frequencies,
+                                    global float *output,
+                                    constant float *coefficients)
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int real_width = get_global_size (0) / 2;
+    const int index = idy * 2 * real_width + idx;
+    const int real_index = idx < real_width ? idx / 2 : real_width - idx / 2;
+
+    output[index] = frequencies[index] * coefficients[real_index];
+}

src/ufo-remove-stripes-task.c

@@ -0,0 +1,300 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+#include <math.h>
+
+#include "ufo-remove-stripes-task.h"
+
+/**
+ * SECTION:ufo-remove-stripes-task
+ * @Short_description: Remove stripes from 2D images, requires frequency domain
+ * input
+ * @Title: remove_stripes
+ *
+ */
+
+struct _UfoRemoveStripesTaskPrivate {
+    guint last_width;
+
+    /* OpenCL */
+    cl_context context;
+    cl_kernel kernel;
+    cl_mem filter_mem;
+
+    /* properties */
+    gfloat strength;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoRemoveStripesTask, ufo_remove_stripes_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_REMOVE_STRIPES_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_REMOVE_STRIPES_TASK, UfoRemoveStripesTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_STRENGTH,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+/**
+ * create_coefficients:
+ * @priv: UfoRemoveStripesTaskPrivate
+ * @width: image width (not the interleaved from fft, rather the real width)
+ *
+ * Compute symmetric filter coefficients.
+ */
+static void
+create_coefficients (UfoRemoveStripesTaskPrivate *priv, const guint width)
+{
+    gfloat *coefficients = g_malloc0 (sizeof (gfloat) * width);
+    gfloat sigma = priv->strength / (2.0 * sqrt (2.0 * log (2.0)));
+    /* Divided by 2 because of the symmetry and we need the +1 to correctly handle
+     * the frequency at width / 2 */
+    const guint real_width = width / 2 + 1;
+    cl_int cl_err;
+
+    g_message ("Creating for width: %u, last width: %u", width, priv->last_width);
+
+    if (!coefficients) {
+        g_warning ("Could not allocate memory for coeefficients");
+    }
+    if (width % 2) {
+        g_warning ("Width must be an even number");
+    }
+
+    for (gint x = 0; x < (gint) real_width; x++) {
+        coefficients[x] = 1.0f - exp (- x * x / (sigma * sigma * 2.0f));
+    }
+
+    if (priv->filter_mem) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->filter_mem));
+    }
+    priv->filter_mem = clCreateBuffer (priv->context,
+                                       CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                                       sizeof (cl_float) * real_width,
+                                       coefficients,
+                                       &cl_err);
+    UFO_RESOURCES_CHECK_CLERR (cl_err);
+    g_free (coefficients);
+    priv->last_width = width;
+}
+
+UfoNode *
+ufo_remove_stripes_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_REMOVE_STRIPES_TASK, NULL));
+}
+
+static void
+ufo_remove_stripes_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+    UfoRemoveStripesTaskPrivate *priv;
+
+    priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (task);
+    priv->context = ufo_resources_get_context (resources);
+    priv->kernel = ufo_resources_get_kernel (resources, "complex.cl", "multiply_frequencies_with_real_sym", error);
+    priv->filter_mem = NULL;
+    UFO_RESOURCES_CHECK_CLERR (clRetainContext (priv->context));
+    if (priv->kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+    }
+}
+
+static void
+ufo_remove_stripes_task_get_requisition (UfoTask *task,
+                                         UfoBuffer **inputs,
+                                         UfoRequisition *requisition)
+{
+    UfoRemoveStripesTaskPrivate *priv;
+
+    priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (task);
+    ufo_buffer_get_requisition (inputs[0], requisition);
+
+    if (!priv->filter_mem || requisition->dims[0] / 2 != priv->last_width) {
+        create_coefficients (priv, requisition->dims[0] / 2);
+    }
+}
+
+static guint
+ufo_remove_stripes_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_remove_stripes_task_get_num_dimensions (UfoTask *task,
+                                            guint input)
+{
+    g_return_val_if_fail (input == 0, 0);
+
+    return 2;
+}
+
+static UfoTaskMode
+ufo_remove_stripes_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+static gboolean
+ufo_remove_stripes_task_process (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoBuffer *output,
+                                 UfoRequisition *requisition)
+{
+    UfoRemoveStripesTaskPrivate *priv;
+    UfoGpuNode *node;
+    UfoProfiler *profiler;
+    cl_command_queue cmd_queue;
+    cl_mem in_mem;
+    cl_mem out_mem;
+
+    priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (task);
+    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+    out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+    in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
+
+    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_mem), &priv->filter_mem));
+
+    profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+    ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL);
+
+    return TRUE;
+}
+
+static void
+ufo_remove_stripes_task_set_property (GObject *object,
+                                      guint property_id,
+                                      const GValue *value,
+                                      GParamSpec *pspec)
+{
+    UfoRemoveStripesTaskPrivate *priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_STRENGTH:
+            priv->strength = g_value_get_float (value);
+            if (priv->last_width) {
+                /* Update only if we know how big the data is */
+                create_coefficients (priv, priv->last_width);
+            }
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_remove_stripes_task_get_property (GObject *object,
+                                      guint property_id,
+                                      GValue *value,
+                                      GParamSpec *pspec)
+{
+    UfoRemoveStripesTaskPrivate *priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_STRENGTH:
+            g_value_set_float (value, priv->strength);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_remove_stripes_task_finalize (GObject *object)
+{
+    UfoRemoveStripesTaskPrivate *priv;
+
+    priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE (object);
+
+    if (priv->filter_mem) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->filter_mem));
+        priv->filter_mem = NULL;
+    }
+    if (priv->kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel));
+        priv->kernel = NULL;
+    }
+    if (priv->context) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseContext (priv->context));
+        priv->context = NULL;
+    }
+
+    G_OBJECT_CLASS (ufo_remove_stripes_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_remove_stripes_task_setup;
+    iface->get_num_inputs = ufo_remove_stripes_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_remove_stripes_task_get_num_dimensions;
+    iface->get_mode = ufo_remove_stripes_task_get_mode;
+    iface->get_requisition = ufo_remove_stripes_task_get_requisition;
+    iface->process = ufo_remove_stripes_task_process;
+}
+
+static void
+ufo_remove_stripes_task_class_init (UfoRemoveStripesTaskClass *klass)
+{
+    GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+    oclass->set_property = ufo_remove_stripes_task_set_property;
+    oclass->get_property = ufo_remove_stripes_task_get_property;
+    oclass->finalize = ufo_remove_stripes_task_finalize;
+
+    properties[PROP_STRENGTH] =
+        g_param_spec_float ("strength",
+                            "Filter strength",
+                            "Filter strength, it is the full width at half maximum of a Gaussian "\
+                            "in the frequency domain. The real filter is then 1 - Gaussian",
+                            1e-3f, G_MAXFLOAT, 1.0f,
+                            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 (oclass, sizeof(UfoRemoveStripesTaskPrivate));
+}
+
+static void
+ufo_remove_stripes_task_init(UfoRemoveStripesTask *self)
+{
+    self->priv = UFO_REMOVE_STRIPES_TASK_GET_PRIVATE(self);
+
+    self->priv->strength = 1.0f;
+    self->priv->last_width = 0;
+}

src/ufo-remove-stripes-task.h

@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifndef __UFO_REMOVE_STRIPES_TASK_H
+#define __UFO_REMOVE_STRIPES_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_REMOVE_STRIPES_TASK             (ufo_remove_stripes_task_get_type())
+#define UFO_REMOVE_STRIPES_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_REMOVE_STRIPES_TASK, UfoRemoveStripesTask))
+#define UFO_IS_REMOVE_STRIPES_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_REMOVE_STRIPES_TASK))
+#define UFO_REMOVE_STRIPES_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_REMOVE_STRIPES_TASK, UfoRemoveStripesTaskClass))
+#define UFO_IS_REMOVE_STRIPES_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_REMOVE_STRIPES_TASK))
+#define UFO_REMOVE_STRIPES_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_REMOVE_STRIPES_TASK, UfoRemoveStripesTaskClass))
+
+typedef struct _UfoRemoveStripesTask           UfoRemoveStripesTask;
+typedef struct _UfoRemoveStripesTaskClass      UfoRemoveStripesTaskClass;
+typedef struct _UfoRemoveStripesTaskPrivate    UfoRemoveStripesTaskPrivate;
+
+/**
+ * UfoRemoveStripesTask:
+ *
+ * The contents of the #UfoRemoveStripesTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoRemoveStripesTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoRemoveStripesTaskPrivate *priv;
+};
+
+/**
+ * UfoRemoveStripesTaskClass:
+ *
+ * #UfoRemoveStripesTask class
+ */
+struct _UfoRemoveStripesTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_remove_stripes_task_new       (void);
+GType     ufo_remove_stripes_task_get_type  (void);
+
+G_END_DECLS
+
+#endif