Import uPIV related filters by Alexandre Lewkowicz

This commit adds relevant filters and kernels to compute rings for the PIV
application.

38 files changed, 5789 insertions, 2 deletions

diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 2811a1d..5f92044 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -7,28 +7,43 @@ set(ufofilter_SRCS
     ufo-buffer-task.c
     ufo-cut-sinogram-task.c
     ufo-center-of-rotation-task.c
+    ufo-concatenate-result-task.c
+    ufo-contrast-task.c
     ufo-dfi-sinc-task.c
+    ufo-denoise-task.c
     ufo-downsample-task.c
+    ufo-dump-ring-task.c
+    ufo-duplicate-task.c
     ufo-filter-task.c
     ufo-flatten-task.c
     ufo-flatten-inplace-task.c
     ufo-flat-field-correction-task.c
     ufo-fft-task.c
+    ufo-fftmult-task.c
+    ufo-filter-particle-task.c
     ufo-forwardproject-task.c
     ufo-gaussian-blur-task.c
     ufo-generate-task.c
+    ufo-get-dup-circ-task.c
     ufo-ifft-task.c
     ufo-interpolate-task.c
     ufo-sharpness-measure-task.c
     ufo-map-slice-task.c
     ufo-median-filter-task.c
     ufo-meta-balls-task.c
+    ufo-multi-search-task.c
     ufo-null-task.c
     ufo-opencl-task.c
+    ufo-ordfilt-task.c
     ufo-pad-task.c
     ufo-phase-retrieval-task.c
     ufo-polar-coordinates-task.c
     ufo-region-of-interest-task.c
+    ufo-reduce-task.c
+    ufo-remove-circle-task.c
+    ufo-ring-pattern-task.c
+    ufo-ringwriter-task.c
+    ufo-replicate-task.c
     ufo-sino-generator-task.c
     ufo-slice-task.c
     ufo-stack-task.c
@@ -40,6 +55,8 @@ set(ufofilter_SRCS
     ufo-zeropadding-task.c
     )
 
+set(ufoaux_SRCS ufo-priv.c)
+
 file(GLOB ufofilter_KERNELS "kernels/*.cl")
 #}}}
 #{{{ Variables
@@ -135,9 +152,9 @@ foreach(_src ${ufofilter_SRCS})
 
         # build single shared library per filter
         if (${CMAKE_SYSTEM_NAME} MATCHES "Darwin")
-            add_library(${target} MODULE ${_src} ${${_misc}})
+            add_library(${target} MODULE ${ufoaux_SRCS} ${_src} ${${_misc}})
         else()
-            add_library(${target} SHARED ${_src} ${${_misc}})
+            add_library(${target} SHARED ${ufoaux_SRCS} ${_src} ${${_misc}})
         endif()
 
         target_link_libraries(${target} ${ufofilter_LIBS})
diff --git a/src/kernels/denoise.cl b/src/kernels/denoise.cl
new file mode 100644
index 0000000..68d18c9
--- /dev/null
+++ b/src/kernels/denoise.cl
@@ -0,0 +1,82 @@
+#include "piv.cl"
+
+
+kernel void
+load_elements (global float *src,
+               global float *dst,
+               int dimension)
+{
+    const int2 imsize = (int2) (get_global_size(0), get_global_size(1));
+    const int x = get_global_id(0);
+    const int y = get_global_id(1);
+
+    unsigned short buffer[12];
+
+    int offset =  dimension * dimension * (x + y * get_global_size(0));
+    int2 center;
+    get_center (dimension, dimension, &center);
+    global float* tmpDst = dst + offset;
+
+    for(int j = 0; j < dimension; ++j) {
+        for(int i = 0; i < dimension; ++i) {
+            *tmpDst = src[get_pel_position (imsize, (int2)(i, j), (int2)(x, y), center)];
+            ++tmpDst;
+        }
+    }
+}
+
+kernel void
+sort_and_set (global const float *src, global float *dst, int num_elements,
+              int arrayLength, float threshold, local float *s_key)
+{
+    // ascendent order
+    unsigned dir = 1;
+    unsigned block_idx = get_global_id(0) / get_local_size(0);
+    //Offset to the beginning of subbatch and load data
+    global const float *d_SrcKey = src + block_idx * num_elements;
+
+    s_key[get_local_id(0)] = d_SrcKey[get_local_id(0)];
+    unsigned next_index = get_local_id(0) + arrayLength / 2;
+
+    s_key[next_index] = next_index < num_elements ? d_SrcKey[next_index] : 0xFFFF;
+
+    for (int size = 2; size < arrayLength; size <<= 1) {
+        //Bitonic merge.  ddd tells if ordering is ascendent or descendent
+        unsigned ddd = dir ^ ((get_local_id(0) & (size / 2)) != 0);
+
+        for (int stride = size / 2; stride > 0; stride >>= 1) {
+            barrier(CLK_LOCAL_MEM_FENCE);
+            unsigned pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            compare_and_exchange (&s_key[pos], &s_key[pos + stride], ddd);
+        }
+    }
+
+    {
+        for (int stride = arrayLength / 2; stride > 0; stride >>= 1) {
+            barrier(CLK_LOCAL_MEM_FENCE);
+            unsigned pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            compare_and_exchange (&s_key[pos], &s_key[pos + stride], dir);
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (get_local_id(0) == 0) {
+        dst[block_idx] = s_key[(int)(num_elements * threshold - 1)];
+    }
+}
+
+kernel void
+/* remove pixels darker than the background image */
+remove_background (global float *src, global float *dst)
+{
+    int2 pel = (int2)(get_global_id(0), get_global_id(1));
+    int idx = pel.x + pel.y * get_global_size(0);
+    /* foreground pixel */
+    float fg = src[idx];
+    /* background pixel */
+    float bk = dst[idx];
+    fg -= bk;
+
+    dst[idx] = fg < 0 ? 0 : fg;
+}
diff --git a/src/kernels/fftmult.cl b/src/kernels/fftmult.cl
new file mode 100644
index 0000000..7d4c578
--- /dev/null
+++ b/src/kernels/fftmult.cl
@@ -0,0 +1,17 @@
+kernel
+void mult(global float* a, global float* b, global float* res)
+{
+    int width = get_global_size(0) * 2;
+    int x = get_global_id(0) * 2;
+    int y = get_global_id(1);
+    int idx_r = x + y * width;
+    int idx_i = 1 + x + y * width;
+
+    float ra = a[idx_r];
+    float rb = b[idx_r];
+    float ia = a[idx_i];
+    float ib = b[idx_i];
+
+    res[idx_r] = ra * rb - ia * ib;
+    res[idx_i] = ra * ib + rb * ia;
+}
diff --git a/src/kernels/ordfilt.cl b/src/kernels/ordfilt.cl
new file mode 100644
index 0000000..70bf042
--- /dev/null
+++ b/src/kernels/ordfilt.cl
@@ -0,0 +1,85 @@
+#include "piv.cl"
+
+
+/* For each pixel, sorts neighboring elements, picks the low_threshold and
+ * high_threshold pixel, computes a value, and stores it in the dst image */
+__kernel void
+bitonic_ordfilt (global const float *src, global float *dst, int num_elements,
+                 int arrayLength, float low_threshold, float high_threshold,
+                 __local float *s_key, unsigned idx_offset)
+{
+    // ascendent order
+    unsigned dir = 1;
+    unsigned block_idx = get_global_id(0) / get_local_size(0);
+    //Offset to the beginning of subbatch and load data
+    global const float *d_SrcKey = src + block_idx * num_elements;
+
+    s_key[get_local_id(0)] = d_SrcKey[get_local_id(0)];
+    unsigned next_index = get_local_id(0) + arrayLength / 2;
+
+    s_key[next_index] = next_index < num_elements ? d_SrcKey[next_index] : 0xFFFF;
+
+    for (int size = 2; size < arrayLength; size <<= 1) {
+        //Bitonic merge.  ddd tells if ordering is ascendent or descendent
+        unsigned ddd = dir ^ ((get_local_id(0) & (size / 2)) != 0);
+
+        for (int stride = size / 2; stride > 0; stride >>= 1) {
+            barrier(CLK_LOCAL_MEM_FENCE);
+            unsigned pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            compare_and_exchange (&s_key[pos], &s_key[pos + stride], ddd);
+        }
+    }
+
+    {
+        for (int stride = arrayLength / 2; stride > 0; stride >>= 1) {
+            barrier(CLK_LOCAL_MEM_FENCE);
+            unsigned pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            compare_and_exchange (&s_key[pos], &s_key[pos + stride], dir);
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (get_local_id(0) == 0) {
+        float low_p = s_key[(int) (num_elements * low_threshold - 1)];
+        float high_p = s_key[(int) (num_elements * high_threshold - 1)];
+
+        unsigned global_idx = block_idx + idx_offset;
+        float intensity = (high_p + low_p) / 2.;
+        float contrast = high_p - low_p;
+        dst[global_idx] = intensity * (1 - contrast);
+    }
+}
+
+kernel void
+/* Load only elements that match the pattern */
+/* Dst is the size of image * number of ones found in pattern */
+load_elements_from_pattern (global float *src,
+                            global float *dst,
+                            global float *pattern,
+                            int dimension,
+                            int num_ones,
+                            unsigned height,
+                            unsigned y_offset)
+{
+    const int2 imsize = (int2) (get_global_size(0), height);
+    const int x = get_global_id(0);
+    const int y = get_global_id(1) + y_offset;
+    const int local_y = get_global_id(1);
+
+    unsigned short buffer[12];
+
+    int2 center;
+    get_center (dimension, dimension, &center);
+    /* Don't apply the offset for the dst buffer */
+    int offset =  num_ones * (x + local_y * imsize.x);
+    global float* tmpDst = dst + offset;
+    for(int j = 0; j < dimension; ++j) {
+        for(int i = 0; i < dimension; ++i) {
+            if (pattern[i + j * dimension]) {
+                *tmpDst = src[get_pel_position (imsize, (int2)(i,j), (int2)(x, y), center)];
+                ++tmpDst;
+            }
+        }
+    }
+}
diff --git a/src/kernels/piv.cl b/src/kernels/piv.cl
new file mode 100644
index 0000000..3a744aa
--- /dev/null
+++ b/src/kernels/piv.cl
@@ -0,0 +1,48 @@
+void
+compare_and_exchange (local float *keyA, local float *keyB, unsigned dir)
+{
+    float t;
+
+    if ((*keyA > *keyB) == dir) {
+        t = *keyA;
+        *keyA = *keyB;
+        *keyB = t;
+    }
+}
+
+int
+get_position (int position, int size)
+{
+    int newPos = 0;
+
+    if (position < 0) {
+        ++position;
+        position *= (-1);
+    }
+
+    int divisor = (position / size);
+    return divisor % 2 ? position % size : size - 1 - (position % size);
+}
+
+void
+get_center (const unsigned width, const unsigned height, int2 *center)
+{
+    center->x = width % 2 == 0 ? width / 2 - 1 : width / 2;
+    center->y = height % 2 == 0 ? height / 2 - 1  : height / 2;
+}
+
+int
+get_pel_position (int2 imageDim, int2 filterPos, int2 pel, int2 center)
+{
+    int2 posInImage = (int2)(filterPos.x - center.x + pel.x, filterPos.y - center.y + pel.y);
+
+    if (posInImage.x < 0 || posInImage.x >= imageDim.x) {
+        posInImage.x = get_position (posInImage.x, imageDim.x);
+    }
+
+    if (posInImage.y < 0 || posInImage.y >= imageDim.y) {
+        posInImage.y = get_position (posInImage.y, imageDim.y);
+    }
+
+    return (posInImage.x + posInImage.y * imageDim.x);
+}
diff --git a/src/ufo-concatenate-result-task.c b/src/ufo-concatenate-result-task.c
new file mode 100644
index 0000000..43cb17b
--- /dev/null
+++ b/src/ufo-concatenate-result-task.c
@@ -0,0 +1,381 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <stdlib.h> /* realloc */
+#include <string.h> /* memcpy */
+#include <assert.h>
+
+#include "ufo-concatenate-result-task.h"
+#include "ufo-ring-coordinates.h"
+
+
+struct _PivFileMetadata {
+    char *piv_file_name;
+    unsigned piv_file_idx;
+};
+
+typedef struct _PivFileMetadata PivFileMetadata;
+
+struct _UfoConcatenateResultTaskPrivate {
+    URCS **data;
+    PivFileMetadata *metadata;
+    unsigned alloc_size;
+    unsigned idx;
+    unsigned current_output_idx;
+    unsigned max_count;
+    unsigned ring_count;
+    unsigned img_alloc_size;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoConcatenateResultTask, ufo_concatenate_result_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_CONCATENATE_RESULT_TASK, UfoConcatenateResultTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_MAX_COUNT,
+    PROP_RING_COUNT,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_concatenate_result_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_CONCATENATE_RESULT_TASK, NULL));
+}
+
+static void
+ufo_concatenate_result_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_concatenate_result_task_get_requisition (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoRequisition *requisition)
+{
+    requisition->n_dims = 1;
+    /* Output size varies according to number of inputs to merge, and data
+     * whithin the data */
+    /* I don't know the output size, I'll allocate it myself once I have
+     * processed the data */
+    requisition->dims[0] = 0;
+}
+
+static guint
+ufo_concatenate_result_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_concatenate_result_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_concatenate_result_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_REDUCTOR | UFO_TASK_MODE_CPU;
+}
+
+static void
+ufo_concatenate_get_metadata(UfoBuffer *src, const char **piv_file_name,
+                             unsigned *piv_file_idx)
+{
+    GValue *value;
+
+    value = ufo_buffer_get_metadata(src, "piv_file_name");
+    *piv_file_name = g_value_get_string(value);
+
+    value = ufo_buffer_get_metadata(src, "piv_file_idx");
+    *piv_file_idx = g_value_get_uint(value);
+}
+
+static void
+initialize_field(UfoConcatenateResultTaskPrivate *priv, const char *piv_file_name,
+                 unsigned piv_file_idx)
+{
+    priv->idx = 0;
+    priv->alloc_size = 16;
+    priv->data[piv_file_idx] = g_malloc (16 * sizeof (UfoRingCoordinate) +
+                                         sizeof (float));
+    priv->data[piv_file_idx]->nb_elt = 0;
+    priv->metadata[piv_file_idx].piv_file_idx = piv_file_idx;
+    priv->metadata[piv_file_idx].piv_file_name = g_strdup (piv_file_name);
+}
+
+static void
+increase_buffer_size(UfoConcatenateResultTaskPrivate *priv)
+{
+    priv->img_alloc_size *= 2;
+    priv->data = g_realloc (priv->data, sizeof (URCS *) * priv->img_alloc_size);
+    priv->metadata = g_realloc (priv->metadata, sizeof (PivFileMetadata) * priv->img_alloc_size);
+
+    // Make sure that we inform that the newly allocated image buffer do not
+    // have any data yet.
+    for (unsigned i = priv->img_alloc_size / 2; i < priv->img_alloc_size; ++i) {
+        priv->data[i] = NULL;
+        priv->metadata[i].piv_file_idx = (unsigned) -1;
+        priv->metadata[i].piv_file_name = NULL;
+    }
+}
+
+static gboolean
+ufo_concatenate_result_task_process (UfoTask *task,
+                                     UfoBuffer **inputs,
+                                     UfoBuffer *output,
+                                     UfoRequisition *requisition)
+{
+    const char *piv_file_name;
+    unsigned piv_file_idx;
+    UfoConcatenateResultTaskPrivate *priv;
+
+    priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE (task);
+    ufo_concatenate_get_metadata (inputs[0], &piv_file_name, &piv_file_idx);
+    assert(priv->ring_count && "Ring count must be set\n");
+    // Check if buffer needs to be increased to be able to store more images.
+
+    if (piv_file_idx >= priv->img_alloc_size)
+        increase_buffer_size (priv);
+
+    if (priv->data[piv_file_idx] == NULL)
+        initialize_field (priv, piv_file_name, piv_file_idx);
+
+    float *input = ufo_buffer_get_host_array (inputs[0], NULL);
+    unsigned nb_coord = (unsigned) *input;
+    UfoRingCoordinate *coord = (UfoRingCoordinate *) (input + 1);
+
+    if (nb_coord > priv->max_count) {
+        g_print("Concatenate: max_count : Ignoring radius %f. %u rings found, "
+                "maximum is %u\n", coord[0].r, nb_coord, priv->max_count);
+        return TRUE;
+    }
+
+    if (priv->idx + nb_coord > priv->alloc_size) {
+        priv->alloc_size += nb_coord;
+        priv->data[piv_file_idx] = g_realloc (priv->data[piv_file_idx], sizeof (UfoRingCoordinate) * priv->alloc_size + sizeof (float));
+    }
+
+    for (unsigned i = 0; i < nb_coord; ++i)
+        priv->data[piv_file_idx]->coord[priv->idx++] = coord[i];
+
+    priv->data[piv_file_idx]->nb_elt += (float) nb_coord;
+
+    return TRUE;
+}
+
+static unsigned
+find_next_idx_for_output (UfoConcatenateResultTaskPrivate *priv)
+{
+    unsigned res = priv->current_output_idx;
+
+    while (res < priv->img_alloc_size && priv->data[res] == NULL)
+        ++res;
+
+    return res;
+}
+
+static void
+attach_output_metadata (UfoBuffer *output, PivFileMetadata *meta)
+{
+    GValue value_uint = G_VALUE_INIT;
+    GValue value_string = G_VALUE_INIT;
+
+    g_value_init (&value_uint, G_TYPE_UINT);
+    g_value_set_uint (&value_uint, meta->piv_file_idx);
+    ufo_buffer_set_metadata (output, "piv_file_idx", &value_uint);
+
+    g_value_init (&value_string, G_TYPE_STRING);
+    g_value_set_string (&value_string, meta->piv_file_name);
+    ufo_buffer_set_metadata(output, "piv_file_name", &value_string);
+}
+
+static gboolean
+ufo_concatenate_result_task_generate (UfoTask *task,
+                                      UfoBuffer *output,
+                                      UfoRequisition *requisition)
+{
+    UfoConcatenateResultTaskPrivate *priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE (task);
+    priv->current_output_idx = find_next_idx_for_output (priv);
+
+    if (priv->current_output_idx >= priv->img_alloc_size)
+        return FALSE;
+
+    unsigned ring_idx = priv->current_output_idx;
+    attach_output_metadata (output, &priv->metadata[ring_idx]);
+    URCS *rings = priv->data[ring_idx];
+    /* There will be at most idx elements */
+    UfoRequisition new_req = {
+        .dims[0] = 1 + (unsigned) rings->nb_elt * sizeof (UfoRingCoordinate) / sizeof (float),
+        .n_dims = 1,
+    };
+    ufo_buffer_resize (output, &new_req);
+    float *res = ufo_buffer_get_host_array (output, NULL);
+    memcpy (res, rings, (unsigned) (rings->nb_elt) * sizeof (UfoRingCoordinate) + sizeof (float));
+    ++priv->current_output_idx;
+    return TRUE;
+}
+
+static void
+ufo_concatenate_result_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoConcatenateResultTaskPrivate *priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MAX_COUNT:
+            priv->max_count = g_value_get_uint (value);
+            break;
+
+        case PROP_RING_COUNT:
+            if (priv->data != NULL)
+                g_free (priv->data);
+
+            if (priv->metadata != NULL)
+                g_free (priv->metadata);
+
+            priv->ring_count = g_value_get_uint (value);
+            priv->data = g_malloc (sizeof (URCS *) * priv->ring_count);
+            priv->metadata = g_malloc (sizeof (PivFileMetadata) * priv->ring_count);
+            priv->img_alloc_size = priv->ring_count;
+
+            for (unsigned i = 0; i < priv->ring_count; ++i) {
+                priv->data[i] = NULL;
+                priv->metadata[i].piv_file_idx = (unsigned) -1;
+                priv->metadata[i].piv_file_name = NULL;
+            }
+
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_concatenate_result_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoConcatenateResultTaskPrivate *priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MAX_COUNT:
+            g_value_set_uint (value, priv->max_count);
+            break;
+
+        case PROP_RING_COUNT:
+            g_value_set_uint (value, priv->ring_count);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_concatenate_result_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_concatenate_result_task_parent_class)->finalize (object);
+    UfoConcatenateResultTaskPrivate *priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE (object);
+
+    for (unsigned i = 0; i < priv->img_alloc_size; ++i) {
+        if (priv->data[i])
+            g_free (priv->data[i]);
+
+        if (priv->metadata[i].piv_file_name)
+            g_free (priv->metadata[i].piv_file_name);
+    }
+
+    g_free (priv->data);
+    g_free (priv->metadata);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_concatenate_result_task_setup;
+    iface->get_num_inputs = ufo_concatenate_result_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_concatenate_result_task_get_num_dimensions;
+    iface->get_mode = ufo_concatenate_result_task_get_mode;
+    iface->get_requisition = ufo_concatenate_result_task_get_requisition;
+    iface->process = ufo_concatenate_result_task_process;
+    iface->generate = ufo_concatenate_result_task_generate;
+}
+
+static void
+ufo_concatenate_result_task_class_init (UfoConcatenateResultTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_concatenate_result_task_set_property;
+    gobject_class->get_property = ufo_concatenate_result_task_get_property;
+    gobject_class->finalize = ufo_concatenate_result_task_finalize;
+
+    properties[PROP_MAX_COUNT] =
+        g_param_spec_uint ("max_count",
+            "The maximum number of rings desired per ring pattern",
+            "The maximum number of rings desired per ring pattern",
+            1, G_MAXUINT, 60,
+            G_PARAM_READWRITE);
+
+    properties[PROP_RING_COUNT] =
+        g_param_spec_uint ("ring_count",
+            "The number of ring pattern generated per image",
+            "The maximum number of rings desired per ring pattern",
+            0, G_MAXUINT, 0,
+            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoConcatenateResultTaskPrivate));
+}
+
+static void
+ufo_concatenate_result_task_init(UfoConcatenateResultTask *self)
+{
+    self->priv = UFO_CONCATENATE_RESULT_TASK_GET_PRIVATE(self);
+    self->priv->idx = 0;
+    self->priv->alloc_size = 0;
+    self->priv->data = NULL;
+    self->priv->metadata = NULL;
+    self->priv->max_count = 60;
+    self->priv->ring_count = 0;
+    self->priv->img_alloc_size = 0;
+    self->priv->current_output_idx = 0;
+}
diff --git a/src/ufo-concatenate-result-task.h b/src/ufo-concatenate-result-task.h
new file mode 100644
index 0000000..6774743
--- /dev/null
+++ b/src/ufo-concatenate-result-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_CONCATENATE_RESULT_TASK_H
+#define __UFO_CONCATENATE_RESULT_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_CONCATENATE_RESULT_TASK             (ufo_concatenate_result_task_get_type())
+#define UFO_CONCATENATE_RESULT_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_CONCATENATE_RESULT_TASK, UfoConcatenateResultTask))
+#define UFO_IS_CONCATENATE_RESULT_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_CONCATENATE_RESULT_TASK))
+#define UFO_CONCATENATE_RESULT_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_CONCATENATE_RESULT_TASK, UfoConcatenateResultTaskClass))
+#define UFO_IS_CONCATENATE_RESULT_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_CONCATENATE_RESULT_TASK))
+#define UFO_CONCATENATE_RESULT_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_CONCATENATE_RESULT_TASK, UfoConcatenateResultTaskClass))
+
+typedef struct _UfoConcatenateResultTask           UfoConcatenateResultTask;
+typedef struct _UfoConcatenateResultTaskClass      UfoConcatenateResultTaskClass;
+typedef struct _UfoConcatenateResultTaskPrivate    UfoConcatenateResultTaskPrivate;
+
+/**
+ * UfoConcatenateResultTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoConcatenateResultTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoConcatenateResultTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoConcatenateResultTaskPrivate *priv;
+};
+
+/**
+ * UfoConcatenateResultTaskClass:
+ *
+ * #UfoConcatenateResultTask class
+ */
+struct _UfoConcatenateResultTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_concatenate_result_task_new       (void);
+GType     ufo_concatenate_result_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-contrast-task.c b/src/ufo-contrast-task.c
new file mode 100644
index 0000000..b8f5b2b
--- /dev/null
+++ b/src/ufo-contrast-task.c
@@ -0,0 +1,311 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <math.h>
+#include <stdlib.h>
+
+#include "ufo-contrast-task.h"
+
+struct _UfoContrastTaskPrivate {
+    gboolean remove_high;
+};
+
+typedef struct _UfoHistogram {
+    float max;
+    float min;
+    double median;
+    unsigned num_bins;
+    /* value of pic in histogram */
+    double pic;
+    double step;
+    /* array of bins if num_bins size */
+    unsigned *bins;
+} UfoHistogram;
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoContrastTask, ufo_contrast_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_CONTRAST_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_CONTRAST_TASK, UfoContrastTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_REMOVE_HIGH,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_contrast_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_CONTRAST_TASK, NULL));
+}
+
+static void
+ufo_contrast_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_contrast_task_get_requisition (UfoTask *task,
+                                   UfoBuffer **inputs,
+                                   UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_contrast_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_contrast_task_get_num_dimensions (UfoTask *task,
+                                      guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_contrast_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static UfoHistogram *
+new_histogram(UfoBuffer *ufo_input)
+{
+    UfoRequisition input_req;
+    ufo_buffer_get_requisition(ufo_input, &input_req);
+    double num_elt = (double) input_req.dims[0] * (double) input_req.dims[1];
+    unsigned num_bins = (unsigned) sqrt(num_elt);
+    float *input = ufo_buffer_get_host_array(ufo_input, NULL);
+
+    UfoHistogram *res;
+
+    res = g_new0 (UfoHistogram, 1);
+    res->bins = g_new0 (unsigned, num_bins);
+    res->num_bins = num_bins;
+    res->min = input[0];
+    res->max = input[0];
+
+    /* Compute max and min element */
+    for (unsigned j = 0; j < input_req.dims[1]; ++j) {
+        for (unsigned i = 0; i < input_req.dims[0]; ++i) {
+            if (res->max < input[i + j * input_req.dims[0]])
+                res->max = input[i + j * input_req.dims[0]];
+
+            if (res->min > input[i + j * input_req.dims[0]])
+                res->min = input[i + j * input_req.dims[0]];
+        }
+    }
+
+    /* Compute histogram */
+    res->step = (double) (res->max - res->min) / (num_bins - 1);
+
+    for (unsigned j = 0; j < input_req.dims[1]; ++j) {
+        for (unsigned i = 0; i < input_req.dims[0]; ++i) {
+            unsigned idx = (unsigned) round((input[i + j * input_req.dims[0]] - res->min) / res->step);
+            ++res->bins[idx];
+        }
+    }
+
+    /* Compute median and pic value */
+    unsigned nb_elt_seen = 0;
+    unsigned max_elt_count_idx = 0;
+
+    for (unsigned i = 0; i < num_bins; ++i) {
+        nb_elt_seen += res->bins[i];
+
+        /* Once median value is set, don't set it again */
+        if (res->median == 0 &&
+            nb_elt_seen >= (input_req.dims[0] * input_req.dims[1]) / 2)
+            res->median = i * res->step + res->min;
+
+        if (res->bins[max_elt_count_idx] < res->bins[i])
+            max_elt_count_idx = i;
+    }
+
+    res->pic = max_elt_count_idx * res->step + res->min;
+    return res;
+}
+
+/* search [start, end[ interval, return index with highest value */
+static double
+histogram_get_pic (UfoHistogram *histogram, unsigned start, unsigned end)
+{
+    unsigned res = start;
+
+    for (unsigned i = start + 1; i < end; ++i) {
+        if (histogram->bins[i] > histogram->bins[res])
+            res = i;
+    }
+
+    return res * histogram->step + histogram->min;
+}
+
+/* Rescale image from low -> high to 0 -> 1 values and enhance constrast by
+ * gamma. Gamma == 1 makes a linear mapping. gamma < 1 produces brighter image.
+ * gamma > 1 produces darker image */
+static void
+imadjust (UfoBuffer *ufo_src, UfoBuffer *ufo_dst, double low, double high,
+          double gamma, float new_high)
+{
+    UfoRequisition req;
+    ufo_buffer_get_requisition (ufo_src, &req);
+    float *src = ufo_buffer_get_host_array(ufo_src, NULL);
+    float *dst = ufo_buffer_get_host_array(ufo_dst, NULL);
+
+    for (unsigned j = 0; j < req.dims[1]; ++j) {
+        for (unsigned i = 0; i < req.dims[0]; ++i) {
+            if (src[i + j * req.dims[0]] >= high)
+                dst[i + j * req.dims[0]] = new_high;
+            else if (src[i + j * req.dims[0]] <= low)
+                dst[i + j * req.dims[0]] = 0.0;
+            else {
+                double normalized = (src[i + j * req.dims[0]] - low) / (double) (high - low);
+                double val = pow(normalized, gamma);
+                dst[i + j * req.dims[0]] = (float) val;
+            }
+        }
+    }
+}
+
+static gboolean
+ufo_contrast_task_process (UfoTask *task,
+                           UfoBuffer **inputs,
+                           UfoBuffer *output,
+                           UfoRequisition *requisition)
+{
+    UfoContrastTaskPrivate *priv = UFO_CONTRAST_TASK_GET_PRIVATE (task);
+    UfoRequisition input_req;
+    ufo_buffer_get_requisition (inputs[0], &input_req);
+    UfoHistogram *histogram = new_histogram (inputs[0]);
+    /* gamma < 1 to make image more bright and enhance contrast */
+    double gamma = 0.3;
+
+    /* Remove values under pic, enhance contrast and normalize image */
+    double pic = histogram_get_pic(histogram, 1, histogram->num_bins - 1);
+
+    /* Remove high pixels of output */
+    if (priv->remove_high) {
+        double crop_max = histogram->max - (histogram->max - pic) / 2;
+        imadjust (inputs[0], output, pic,  crop_max, gamma, 0.0);
+    }
+    else {
+        /* transpose image from [pic, 1] to [0, 1] */
+        imadjust (inputs[0], output, pic,  histogram->max, gamma, 1.0);
+    }
+
+    g_free (histogram->bins);
+    g_free (histogram);
+    return TRUE;
+}
+
+static void
+ufo_contrast_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoContrastTaskPrivate *priv = UFO_CONTRAST_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_REMOVE_HIGH:
+            priv->remove_high = g_value_get_boolean(value);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_contrast_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoContrastTaskPrivate *priv = UFO_CONTRAST_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_REMOVE_HIGH:
+            g_value_set_boolean(value, priv->remove_high);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_contrast_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_contrast_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_contrast_task_setup;
+    iface->get_num_inputs = ufo_contrast_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_contrast_task_get_num_dimensions;
+    iface->get_mode = ufo_contrast_task_get_mode;
+    iface->get_requisition = ufo_contrast_task_get_requisition;
+    iface->process = ufo_contrast_task_process;
+}
+
+static void
+ufo_contrast_task_class_init (UfoContrastTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_contrast_task_set_property;
+    gobject_class->get_property = ufo_contrast_task_get_property;
+    gobject_class->finalize = ufo_contrast_task_finalize;
+
+    properties[PROP_REMOVE_HIGH] =
+        g_param_spec_boolean ("remove_high",
+            "Tells whether or not to set high intensity pixels to 0",
+            "Tells whether or not to set high intensity pixels to 0",
+            0,
+            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoContrastTaskPrivate));
+}
+
+static void
+ufo_contrast_task_init(UfoContrastTask *self)
+{
+    self->priv = UFO_CONTRAST_TASK_GET_PRIVATE(self);
+    self->priv->remove_high = 0;
+}
diff --git a/src/ufo-contrast-task.h b/src/ufo-contrast-task.h
new file mode 100644
index 0000000..f03907d
--- /dev/null
+++ b/src/ufo-contrast-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_CONTRAST_TASK_H
+#define __UFO_CONTRAST_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_CONTRAST_TASK             (ufo_contrast_task_get_type())
+#define UFO_CONTRAST_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_CONTRAST_TASK, UfoContrastTask))
+#define UFO_IS_CONTRAST_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_CONTRAST_TASK))
+#define UFO_CONTRAST_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_CONTRAST_TASK, UfoContrastTaskClass))
+#define UFO_IS_CONTRAST_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_CONTRAST_TASK))
+#define UFO_CONTRAST_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_CONTRAST_TASK, UfoContrastTaskClass))
+
+typedef struct _UfoContrastTask           UfoContrastTask;
+typedef struct _UfoContrastTaskClass      UfoContrastTaskClass;
+typedef struct _UfoContrastTaskPrivate    UfoContrastTaskPrivate;
+
+/**
+ * UfoContrastTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoContrastTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoContrastTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoContrastTaskPrivate *priv;
+};
+
+/**
+ * UfoContrastTaskClass:
+ *
+ * #UfoContrastTask class
+ */
+struct _UfoContrastTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_contrast_task_new       (void);
+GType     ufo_contrast_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-denoise-task.c b/src/ufo-denoise-task.c
new file mode 100644
index 0000000..1157d14
--- /dev/null
+++ b/src/ufo-denoise-task.c
@@ -0,0 +1,356 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <math.h>
+#include "ufo-denoise-task.h"
+#include "ufo-priv.h"
+
+
+struct _UfoDenoiseTaskPrivate {
+    cl_kernel k_sort_and_set;
+    cl_kernel k_load_elements;
+    cl_kernel k_remove_background;
+    unsigned matrix_size;
+    gpointer context;
+    UfoResources *resources;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoDenoiseTask, ufo_denoise_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_DENOISE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_DENOISE_TASK, UfoDenoiseTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_MATRIX_SIZE,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_denoise_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_DENOISE_TASK, NULL));
+}
+
+static void
+ufo_denoise_task_setup (UfoTask *task,
+                   UfoResources *resources,
+                   GError **error)
+{
+    UfoDenoiseTaskPrivate *priv;
+
+    priv = UFO_DENOISE_TASK_GET_PRIVATE (task);
+    priv->context = ufo_resources_get_context (resources);
+    priv->resources = resources;
+
+    priv->k_sort_and_set = ufo_resources_get_kernel (resources, "denoise.cl", "sort_and_set", error);
+
+    if (priv->k_sort_and_set != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_sort_and_set));
+
+    priv->k_load_elements = ufo_resources_get_kernel (resources, "denoise.cl", "load_elements", error);
+
+    if (priv->k_load_elements != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_load_elements));
+
+    priv->k_remove_background = ufo_resources_get_kernel (resources, "denoise.cl", "remove_background", error);
+
+    if (priv->k_remove_background != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_remove_background));
+}
+
+static void
+ufo_denoise_task_get_requisition (UfoTask *task,
+                                  UfoBuffer **inputs,
+                                  UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition (inputs[0], requisition);
+}
+
+static guint
+ufo_denoise_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_denoise_task_get_num_dimensions (UfoTask *task,
+                                     guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_denoise_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+static void
+launch_kernel_1D(cl_kernel kernel, UfoBuffer *ufo_src, UfoBuffer *ufo_dst,
+                 cl_command_queue cmd_queue, size_t dimension)
+{
+    cl_mem dst;
+    cl_mem src;
+    UfoRequisition requisition;
+    size_t global_work_size[1];
+    size_t local_work_size[1];
+
+    dst = ufo_buffer_get_device_array(ufo_dst, cmd_queue);
+    src = ufo_buffer_get_device_array(ufo_src, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof (cl_mem),
+                                               &src));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof (cl_mem),
+                                               &dst));
+    size_t num_elements = dimension * dimension;
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof (cl_int),
+                                               &num_elements));
+    /* Power of 2 above dimension * dimension */
+    size_t array_length = (size_t) ceil_power_of_two (dimension * dimension);
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 3, sizeof (cl_int),
+                                               &array_length));
+    float threshold = (float) 0.3;
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 4, sizeof (cl_float),
+                                               &threshold));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 5, sizeof (cl_float) *
+                                               array_length, NULL));
+
+    /* Launch the kernel over 1D grid */
+    ufo_buffer_get_requisition(ufo_src, &requisition);
+    global_work_size[0] = requisition.dims[0] * requisition.dims[1] *
+                          (array_length / 2);
+    /* Power of two below number of elements to sort */
+    local_work_size[0] = array_length / 2;
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue,
+                                                       kernel,
+                                                       1, NULL, global_work_size,
+                                                       local_work_size,
+                                                       0, NULL, NULL));
+}
+
+static void
+get_max_work_group_size (UfoResources *resources, size_t *x_worker_count,
+                         size_t * y_worker_count)
+{
+    *x_worker_count = ULLONG_MAX;
+    GList *devices = ufo_resources_get_devices (resources);
+    GList *it;
+    g_list_for (devices, it) {
+        cl_device_id device = (cl_device_id) it->data;
+        size_t byte_count = 0;
+        size_t max_group_size = 0;
+        clGetDeviceInfo (device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (size_t),
+                         &max_group_size, &byte_count);
+        g_assert (sizeof (size_t) == byte_count);
+        if (max_group_size < *x_worker_count)
+            *x_worker_count = max_group_size;
+    }
+    *x_worker_count = (unsigned) sqrtf((float)*x_worker_count);
+    *y_worker_count = *x_worker_count;
+}
+
+static void
+launch_kernel_2D(UfoDenoiseTaskPrivate *priv,
+                 cl_kernel kernel, UfoBuffer *ufo_src, UfoBuffer *ufo_dst,
+                 cl_command_queue cmd_queue, size_t dimension)
+{
+    cl_mem dst;
+    cl_mem src;
+    UfoRequisition requisition;
+    size_t global_work_size[2];
+    size_t local_work_size[2];
+    dst = ufo_buffer_get_device_array(ufo_dst, cmd_queue);
+    src = ufo_buffer_get_device_array(ufo_src, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof (cl_mem),
+                                               &src));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof (cl_mem),
+                                               &dst));
+    if (dimension)
+        UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof (cl_int),
+                                                   &dimension));
+
+    ufo_buffer_get_requisition(ufo_src, &requisition);
+    global_work_size[0] = requisition.dims[0];
+    global_work_size[1] = requisition.dims[1];
+    size_t y_worker_count, x_worker_count;
+    get_max_work_group_size(priv->resources, &x_worker_count, &y_worker_count);
+
+    while (global_work_size[1] % y_worker_count)
+        --y_worker_count;
+
+    while (global_work_size[0] % x_worker_count)
+        --x_worker_count;
+
+    local_work_size[0] = x_worker_count; /* Multiple of image_width=1080 */
+    local_work_size[1] = y_worker_count; /* Multiple of image_height=1280 */
+    /* Launch the kernel over 1D grid */
+    local_work_size[0] = x_worker_count; /* Multiple of image_width=1080 */
+    local_work_size[1] = y_worker_count; /* Multiple of image_height=1280 */
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue,
+                                                       kernel,
+                                                       2, NULL, global_work_size,
+                                                       local_work_size,
+                                                       0, NULL, NULL));
+}
+
+static void
+get_background_image (UfoDenoiseTaskPrivate *priv, UfoBuffer *src, UfoBuffer *dst,
+                      cl_command_queue cmd_queue)
+{
+    size_t dimension = priv->matrix_size;
+    UfoRequisition image_requisition;
+    ufo_buffer_get_requisition(src, &image_requisition);
+    UfoRequisition requisition = {
+        .n_dims = 3,
+        .dims[0] = image_requisition.dims[0],
+        .dims[1] = image_requisition.dims[1],
+        .dims[2] = dimension * dimension,
+    };
+    UfoBuffer *ufo_buffer = ufo_buffer_new(&requisition, priv->context);
+
+    /* loads surrounding dimension * dimension pixels of each pixel in image.
+     * Result in buffer */
+    launch_kernel_2D (priv, priv->k_load_elements, src, ufo_buffer, cmd_queue, dimension);
+
+    /* Create background image using sorted neigbhooring pixels */
+    launch_kernel_1D (priv->k_sort_and_set, ufo_buffer, dst, cmd_queue, dimension);
+
+    g_object_unref(ufo_buffer);
+}
+
+static gboolean
+ufo_denoise_task_process (UfoTask *task,
+                          UfoBuffer **inputs,
+                          UfoBuffer *output,
+                          UfoRequisition *requisition)
+{
+    UfoDenoiseTaskPrivate *priv;
+    UfoGpuNode *node;
+    cl_command_queue cmd_queue;
+
+    priv = UFO_DENOISE_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);
+
+    get_background_image (priv, inputs[0], output, cmd_queue);
+    launch_kernel_2D (priv, priv->k_remove_background, inputs[0], output, cmd_queue, 0);
+    return TRUE;
+}
+
+static void
+ufo_denoise_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoDenoiseTaskPrivate *priv = UFO_DENOISE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MATRIX_SIZE:
+            priv->matrix_size = g_value_get_uint(value);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_denoise_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoDenoiseTaskPrivate *priv = UFO_DENOISE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MATRIX_SIZE:
+            g_value_set_uint (value, priv->matrix_size);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_denoise_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_denoise_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_denoise_task_setup;
+    iface->get_num_inputs = ufo_denoise_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_denoise_task_get_num_dimensions;
+    iface->get_mode = ufo_denoise_task_get_mode;
+    iface->get_requisition = ufo_denoise_task_get_requisition;
+    iface->process = ufo_denoise_task_process;
+}
+
+static void
+ufo_denoise_task_class_init (UfoDenoiseTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_denoise_task_set_property;
+    gobject_class->get_property = ufo_denoise_task_get_property;
+    gobject_class->finalize = ufo_denoise_task_finalize;
+
+    properties[PROP_MATRIX_SIZE] =
+        g_param_spec_uint ("matrix_size",
+            "determines the number of surrounding pixels to be compared with",
+            "determines the number of surrounding pixels to be compared with",
+            1, G_MAXUINT, 13,
+            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoDenoiseTaskPrivate));
+}
+
+static void
+ufo_denoise_task_init(UfoDenoiseTask *self)
+{
+    self->priv = UFO_DENOISE_TASK_GET_PRIVATE(self);
+    self->priv->context = NULL;
+    self->priv->matrix_size = 13;
+    self->priv->k_sort_and_set = NULL;
+    self->priv->k_load_elements = NULL;
+    self->priv->k_remove_background = NULL;
+}
diff --git a/src/ufo-denoise-task.h b/src/ufo-denoise-task.h
new file mode 100644
index 0000000..6410599
--- /dev/null
+++ b/src/ufo-denoise-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_DENOISE_TASK_H
+#define __UFO_DENOISE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_DENOISE_TASK             (ufo_denoise_task_get_type())
+#define UFO_DENOISE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_DENOISE_TASK, UfoDenoiseTask))
+#define UFO_IS_DENOISE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_DENOISE_TASK))
+#define UFO_DENOISE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_DENOISE_TASK, UfoDenoiseTaskClass))
+#define UFO_IS_DENOISE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_DENOISE_TASK))
+#define UFO_DENOISE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_DENOISE_TASK, UfoDenoiseTaskClass))
+
+typedef struct _UfoDenoiseTask           UfoDenoiseTask;
+typedef struct _UfoDenoiseTaskClass      UfoDenoiseTaskClass;
+typedef struct _UfoDenoiseTaskPrivate    UfoDenoiseTaskPrivate;
+
+/**
+ * UfoDenoiseTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoDenoiseTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoDenoiseTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoDenoiseTaskPrivate *priv;
+};
+
+/**
+ * UfoDenoiseTaskClass:
+ *
+ * #UfoDenoiseTask class
+ */
+struct _UfoDenoiseTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_denoise_task_new       (void);
+GType     ufo_denoise_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
diff --git a/src/ufo-dump-ring-task.c b/src/ufo-dump-ring-task.c
new file mode 100644
index 0000000..582a783
--- /dev/null
+++ b/src/ufo-dump-ring-task.c
@@ -0,0 +1,239 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <math.h>
+
+#include "ufo-dump-ring-task.h"
+#include "ufo-ring-coordinates.h"
+#include "ufo-priv.h"
+
+struct _UfoDumpRingTaskPrivate {
+    unsigned scale;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoDumpRingTask, ufo_dump_ring_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_DUMP_RING_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_DUMP_RING_TASK, UfoDumpRingTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_SCALE,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_dump_ring_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_DUMP_RING_TASK, NULL));
+}
+
+static void
+ufo_dump_ring_task_setup (UfoTask *task,
+                          UfoResources *resources,
+                          GError **error)
+{
+}
+
+static void
+ufo_dump_ring_task_get_requisition (UfoTask *task,
+                                    UfoBuffer **inputs,
+                                    UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_dump_ring_task_get_num_inputs (UfoTask *task)
+{
+    return 2;
+}
+
+static guint
+ufo_dump_ring_task_get_num_dimensions (UfoTask *task,
+                                       guint input)
+{
+    if (input == 0)
+        /* First input is source image */
+        return 2;
+    else
+        /* Second input is coordinates array */
+        return 1;
+}
+
+static UfoTaskMode
+ufo_dump_ring_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static void
+copy_image_and_get_max (UfoBuffer *ufo_img, UfoBuffer *ufo_dst, float *max)
+{
+    float *img = ufo_buffer_get_host_array (ufo_img, NULL);
+    float *dst = ufo_buffer_get_host_array (ufo_dst, NULL);
+    *max = img[0];
+    UfoRequisition req;
+    ufo_buffer_get_requisition (ufo_img, &req);
+
+    for (size_t y = 0; y < req.dims[1]; ++y) {
+        for (size_t x = 0; x < req.dims[0]; ++x) {
+            float value = img[y * req.dims[0] + x];
+            dst[y * req.dims[0] + x] = value;
+            *max = *max < value ? value : *max;
+        }
+    }
+}
+
+static void
+dump_circles (UfoDumpRingTaskPrivate *priv, UfoBuffer *ufo_rings,
+              UfoBuffer *ufo_dst, float max)
+{
+    if (max > 100)
+        max += 100;
+    else
+        max *= 2;
+
+    float *rings = ufo_buffer_get_host_array (ufo_rings, NULL);
+    float *dst = ufo_buffer_get_host_array (ufo_dst, NULL);
+    unsigned nb_rings = (unsigned) *rings;
+    UfoRequisition req;
+    ufo_buffer_get_requisition (ufo_dst, &req);
+    UfoRingCoordinate *coord = (UfoRingCoordinate *) (rings + 1);
+
+    for (unsigned i = 0; i < nb_rings; ++i) {
+        float x = roundf (coord[i].x * (float) priv->scale);
+        float y = roundf (coord[i].y * (float) priv->scale);
+        float r = roundf (coord[i].r * (float) priv->scale);
+
+        for (float t = 0; t < 2 * G_PI; t += (float) 0.005) {
+            int x2 = (int) round (r * cos (t) + x);
+            int y2 = (int) round (r * sin (t) + y);
+            if (x2 >= 0 && x2 < (int) req.dims[0] && y2 >=0 && y2 < (int) req.dims[1]) {
+                unsigned idx = (unsigned) ((unsigned) x2 + (unsigned) y2 * req.dims[0]);
+                dst[idx] = max;
+            }
+        }
+    }
+}
+
+static gboolean
+ufo_dump_ring_task_process (UfoTask *task,
+                            UfoBuffer **inputs,
+                            UfoBuffer *output,
+                            UfoRequisition *requisition)
+{
+    float max;
+    copy_image_and_get_max (inputs[0], output, &max);
+    UfoDumpRingTaskPrivate *priv = UFO_DUMP_RING_TASK_GET_PRIVATE (task);
+    /* use max_value for drawing circles */
+    dump_circles (priv, inputs[1], output, max);
+    return TRUE;
+}
+
+static void
+ufo_dump_ring_task_set_property (GObject *object,
+                                 guint property_id,
+                                 const GValue *value,
+                                 GParamSpec *pspec)
+{
+    UfoDumpRingTaskPrivate *priv = UFO_DUMP_RING_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            priv->scale = g_value_get_uint(value);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_dump_ring_task_get_property (GObject *object,
+                                 guint property_id,
+                                 GValue *value,
+                                 GParamSpec *pspec)
+{
+    UfoDumpRingTaskPrivate *priv = UFO_DUMP_RING_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            g_value_set_uint (value, priv->scale);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_dump_ring_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_dump_ring_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_dump_ring_task_setup;
+    iface->get_num_inputs = ufo_dump_ring_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_dump_ring_task_get_num_dimensions;
+    iface->get_mode = ufo_dump_ring_task_get_mode;
+    iface->get_requisition = ufo_dump_ring_task_get_requisition;
+    iface->process = ufo_dump_ring_task_process;
+}
+
+static void
+ufo_dump_ring_task_class_init (UfoDumpRingTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_dump_ring_task_set_property;
+    gobject_class->get_property = ufo_dump_ring_task_get_property;
+    gobject_class->finalize = ufo_dump_ring_task_finalize;
+
+    properties[PROP_SCALE] =
+        g_param_spec_uint ("scale",
+                           "Says by how much rings should be increased",
+                           "Says by how much rings should be increased",
+                           1, G_MAXUINT, 1,
+                           G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoDumpRingTaskPrivate));
+}
+
+static void
+ufo_dump_ring_task_init(UfoDumpRingTask *self)
+{
+    self->priv = UFO_DUMP_RING_TASK_GET_PRIVATE(self);
+    self->priv->scale = 1;
+}
diff --git a/src/ufo-dump-ring-task.h b/src/ufo-dump-ring-task.h
new file mode 100644
index 0000000..e1c203d
--- /dev/null
+++ b/src/ufo-dump-ring-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_DUMP_RING_TASK_H
+#define __UFO_DUMP_RING_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_DUMP_RING_TASK             (ufo_dump_ring_task_get_type())
+#define UFO_DUMP_RING_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_DUMP_RING_TASK, UfoDumpRingTask))
+#define UFO_IS_DUMP_RING_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_DUMP_RING_TASK))
+#define UFO_DUMP_RING_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_DUMP_RING_TASK, UfoDumpRingTaskClass))
+#define UFO_IS_DUMP_RING_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_DUMP_RING_TASK))
+#define UFO_DUMP_RING_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_DUMP_RING_TASK, UfoDumpRingTaskClass))
+
+typedef struct _UfoDumpRingTask           UfoDumpRingTask;
+typedef struct _UfoDumpRingTaskClass      UfoDumpRingTaskClass;
+typedef struct _UfoDumpRingTaskPrivate    UfoDumpRingTaskPrivate;
+
+/**
+ * UfoDumpRingTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoDumpRingTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoDumpRingTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoDumpRingTaskPrivate *priv;
+};
+
+/**
+ * UfoDumpRingTaskClass:
+ *
+ * #UfoDumpRingTask class
+ */
+struct _UfoDumpRingTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_dump_ring_task_new       (void);
+GType     ufo_dump_ring_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-duplicate-task.c b/src/ufo-duplicate-task.c
new file mode 100644
index 0000000..3541a90
--- /dev/null
+++ b/src/ufo-duplicate-task.c
@@ -0,0 +1,130 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <stdlib.h>
+#include "ufo-duplicate-task.h"
+
+
+struct _UfoDuplicateTaskPrivate {
+    unsigned alloc_size;
+    unsigned idx;
+    UfoBuffer *data;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoDuplicateTask, ufo_duplicate_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_DUPLICATE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_DUPLICATE_TASK, UfoDuplicateTaskPrivate))
+
+UfoNode *
+ufo_duplicate_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_DUPLICATE_TASK, NULL));
+}
+
+static void
+ufo_duplicate_task_setup (UfoTask *task,
+                          UfoResources *resources,
+                          GError **error)
+{
+}
+
+static void
+ufo_duplicate_task_get_requisition (UfoTask *task,
+                                    UfoBuffer **inputs,
+                                    UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_duplicate_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_duplicate_task_get_num_dimensions (UfoTask *task,
+                                       guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_duplicate_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_REDUCTOR;
+}
+
+static gboolean
+ufo_duplicate_task_process (UfoTask *task,
+                            UfoBuffer **inputs,
+                            UfoBuffer *output,
+                            UfoRequisition *requisition)
+{
+    UfoDuplicateTaskPrivate *priv = UFO_DUPLICATE_TASK_GET_PRIVATE (task);
+
+    if (priv->idx > priv->alloc_size) {
+        priv->alloc_size *= 2;
+        size_t nb_bytes = priv->alloc_size * sizeof (UfoBuffer *);
+        priv->data = realloc (priv->data, nb_bytes);
+    }
+    return TRUE;
+}
+
+static void
+ufo_duplicate_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_duplicate_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_duplicate_task_setup;
+    iface->get_num_inputs = ufo_duplicate_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_duplicate_task_get_num_dimensions;
+    iface->get_mode = ufo_duplicate_task_get_mode;
+    iface->get_requisition = ufo_duplicate_task_get_requisition;
+    iface->process = ufo_duplicate_task_process;
+}
+
+static void
+ufo_duplicate_task_class_init (UfoDuplicateTaskClass *klass)
+{
+    GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+    oclass->finalize = ufo_duplicate_task_finalize;
+
+    g_type_class_add_private (oclass, sizeof(UfoDuplicateTaskPrivate));
+}
+
+static void
+ufo_duplicate_task_init(UfoDuplicateTask *self)
+{
+    self->priv = UFO_DUPLICATE_TASK_GET_PRIVATE(self);
+    self->priv->alloc_size = 256;
+    self->priv->data = malloc (sizeof (UfoBuffer *) * self->priv->alloc_size);
+    self->priv->idx = 0;
+}
diff --git a/src/ufo-duplicate-task.h b/src/ufo-duplicate-task.h
new file mode 100644
index 0000000..f40b5fa
--- /dev/null
+++ b/src/ufo-duplicate-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_DUPLICATE_TASK_H
+#define __UFO_DUPLICATE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_DUPLICATE_TASK             (ufo_duplicate_task_get_type())
+#define UFO_DUPLICATE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_DUPLICATE_TASK, UfoDuplicateTask))
+#define UFO_IS_DUPLICATE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_DUPLICATE_TASK))
+#define UFO_DUPLICATE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_DUPLICATE_TASK, UfoDuplicateTaskClass))
+#define UFO_IS_DUPLICATE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_DUPLICATE_TASK))
+#define UFO_DUPLICATE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_DUPLICATE_TASK, UfoDuplicateTaskClass))
+
+typedef struct _UfoDuplicateTask           UfoDuplicateTask;
+typedef struct _UfoDuplicateTaskClass      UfoDuplicateTaskClass;
+typedef struct _UfoDuplicateTaskPrivate    UfoDuplicateTaskPrivate;
+
+/**
+ * UfoDuplicateTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoDuplicateTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoDuplicateTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoDuplicateTaskPrivate *priv;
+};
+
+/**
+ * UfoDuplicateTaskClass:
+ *
+ * #UfoDuplicateTask class
+ */
+struct _UfoDuplicateTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_duplicate_task_new       (void);
+GType     ufo_duplicate_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-fftmult-task.c b/src/ufo-fftmult-task.c
new file mode 100644
index 0000000..68d1bf2
--- /dev/null
+++ b/src/ufo-fftmult-task.c
@@ -0,0 +1,224 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#ifdef __APPLE__
+# include <OpenCL/cl.h>
+#else
+# include <CL/cl.h>
+#endif
+
+#include <math.h>
+#include "ufo-fftmult-task.h"
+#include "ufo-priv.h"
+
+struct _UfoFftmultTaskPrivate {
+    cl_kernel k_fftmult;
+    UfoResources *resources;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoFftmultTask, ufo_fftmult_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_FFTMULT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FFTMULT_TASK, UfoFftmultTaskPrivate))
+
+UfoNode *
+ufo_fftmult_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_FFTMULT_TASK, NULL));
+}
+
+static void
+ufo_fftmult_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+    UfoFftmultTaskPrivate *priv;
+
+    priv = UFO_FFTMULT_TASK_GET_PRIVATE (task);
+    priv->resources = resources;
+
+    priv->k_fftmult = ufo_resources_get_kernel (resources, "fftmult.cl", "mult", error);
+
+    if (priv->k_fftmult != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_fftmult));
+}
+
+static void
+ufo_fftmult_task_get_requisition (UfoTask *task,
+                                  UfoBuffer **inputs,
+                                  UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[1], requisition);
+}
+
+static guint
+ufo_fftmult_task_get_num_inputs (UfoTask *task)
+{
+    return 2;
+}
+
+static guint
+ufo_fftmult_task_get_num_dimensions (UfoTask *task,
+                                     guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_fftmult_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+static void
+get_ring_metadata(UfoBuffer *src, unsigned *number_ones, unsigned *radius)
+{
+    GValue *value;
+    value = ufo_buffer_get_metadata(src, "radius");
+    *radius = g_value_get_uint(value);
+    value = ufo_buffer_get_metadata(src, "number_ones");
+    *number_ones = g_value_get_uint(value);
+}
+
+static void
+get_max_work_group_size (UfoResources *resources, size_t *x_worker_count,
+                         size_t * y_worker_count)
+{
+    *x_worker_count = ULLONG_MAX;
+    GList *devices = ufo_resources_get_devices (resources);
+    GList *it;
+
+    g_list_for (devices, it) {
+        cl_device_id device = (cl_device_id) it->data;
+        size_t byte_count = 0;
+        size_t max_group_size = 0;
+        clGetDeviceInfo (device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (size_t),
+                         &max_group_size, &byte_count);
+        g_assert (sizeof (size_t) == byte_count);
+        if (max_group_size < *x_worker_count)
+            *x_worker_count = max_group_size;
+    }
+
+    *x_worker_count = (unsigned) sqrtf((float)*x_worker_count);
+    *y_worker_count = *x_worker_count;
+}
+
+static void
+launch_kernel_2D(UfoFftmultTaskPrivate *priv,
+                 UfoBuffer *ufo_a, UfoBuffer *ufo_b,
+                 UfoBuffer *ufo_dst, cl_command_queue cmd_queue)
+{
+    cl_kernel kernel = priv->k_fftmult;
+    cl_mem a, b, dst;
+    UfoRequisition requisition;
+    size_t global_work_size[2];
+    size_t local_work_size[2];
+    dst = ufo_buffer_get_device_array(ufo_dst, cmd_queue);
+    a = ufo_buffer_get_device_array(ufo_a, cmd_queue);
+    b = ufo_buffer_get_device_array(ufo_b, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof (cl_mem), &a));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof (cl_mem), &b));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof (cl_mem), &dst));
+
+    /* Launch the kernel over 2D grid, using dst requisition which reperesents a
+     * crop of the image */
+    ufo_buffer_get_requisition (ufo_dst, &requisition);
+    global_work_size[0] = requisition.dims[0] / 2;
+    /* Buffer may have mod extra rows, don't take them into account */
+    g_assert(requisition.dims[0] % 2 == 0 && "FFT images are multiples of 2\n");
+    global_work_size[1] = requisition.dims[1];
+    size_t y_worker_count, x_worker_count;
+    get_max_work_group_size(priv->resources, &x_worker_count, &y_worker_count);
+
+    while (global_work_size[1] % y_worker_count)
+        --y_worker_count;
+
+    while (global_work_size[0] % x_worker_count)
+        --x_worker_count;
+
+    local_work_size[0] = x_worker_count; /* Multiple of image_width=1080 */
+    local_work_size[1] = y_worker_count; /* Multiple of image_height=1280 */
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue,
+                                                       kernel,
+                                                       2, NULL, global_work_size,
+                                                       local_work_size,
+                                                       0, NULL, NULL));
+}
+
+static gboolean
+ufo_fftmult_task_process (UfoTask *task,
+                          UfoBuffer **inputs,
+                          UfoBuffer *output,
+                          UfoRequisition *requisition)
+{
+    UfoFftmultTaskPrivate *priv;
+    priv = UFO_FFTMULT_TASK_GET_PRIVATE (task);
+    /* Forwarding ring radius metada to next plugin */
+    unsigned radius, number_ones;
+    get_ring_metadata(inputs[1], &number_ones, &radius);
+
+    UfoGpuNode *node;
+    cl_command_queue cmd_queue;
+
+    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+
+    priv = UFO_FFTMULT_TASK_GET_PRIVATE (task);
+    launch_kernel_2D (priv, inputs[0], inputs[1], output, cmd_queue);
+    return TRUE;
+}
+
+static void
+ufo_fftmult_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_fftmult_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_fftmult_task_setup;
+    iface->get_num_inputs = ufo_fftmult_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_fftmult_task_get_num_dimensions;
+    iface->get_mode = ufo_fftmult_task_get_mode;
+    iface->get_requisition = ufo_fftmult_task_get_requisition;
+    iface->process = ufo_fftmult_task_process;
+}
+
+static void
+ufo_fftmult_task_class_init (UfoFftmultTaskClass *klass)
+{
+    GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+    oclass->finalize = ufo_fftmult_task_finalize;
+
+    g_type_class_add_private (oclass, sizeof(UfoFftmultTaskPrivate));
+}
+
+static void
+ufo_fftmult_task_init(UfoFftmultTask *self)
+{
+    self->priv = UFO_FFTMULT_TASK_GET_PRIVATE(self);
+}
diff --git a/src/ufo-fftmult-task.h b/src/ufo-fftmult-task.h
new file mode 100644
index 0000000..9d4416d
--- /dev/null
+++ b/src/ufo-fftmult-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_FFTMULT_TASK_H
+#define __UFO_FFTMULT_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_FFTMULT_TASK             (ufo_fftmult_task_get_type())
+#define UFO_FFTMULT_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FFTMULT_TASK, UfoFftmultTask))
+#define UFO_IS_FFTMULT_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FFTMULT_TASK))
+#define UFO_FFTMULT_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FFTMULT_TASK, UfoFftmultTaskClass))
+#define UFO_IS_FFTMULT_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FFTMULT_TASK))
+#define UFO_FFTMULT_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FFTMULT_TASK, UfoFftmultTaskClass))
+
+typedef struct _UfoFftmultTask           UfoFftmultTask;
+typedef struct _UfoFftmultTaskClass      UfoFftmultTaskClass;
+typedef struct _UfoFftmultTaskPrivate    UfoFftmultTaskPrivate;
+
+/**
+ * UfoFftmultTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoFftmultTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoFftmultTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoFftmultTaskPrivate *priv;
+};
+
+/**
+ * UfoFftmultTaskClass:
+ *
+ * #UfoFftmultTask class
+ */
+struct _UfoFftmultTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_fftmult_task_new       (void);
+GType     ufo_fftmult_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-filter-particle-task.c b/src/ufo-filter-particle-task.c
new file mode 100644
index 0000000..e054fc5
--- /dev/null
+++ b/src/ufo-filter-particle-task.c
@@ -0,0 +1,602 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <math.h>
+
+#include "ufo-filter-particle-task.h"
+#include "ufo-ring-coordinates.h"
+
+struct _UfoFilterParticleTaskPrivate {
+    unsigned *img;
+    size_t num_pixels;
+    float threshold;
+    float min;
+};
+
+/*Maximal number of particles that can be tagged*/
+#define nVectMax 100000
+
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoFilterParticleTask, ufo_filter_particle_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_FILTER_PARTICLE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_PARTICLE_TASK, UfoFilterParticleTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_MIN,
+    PROP_THRESHOLD,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_filter_particle_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_FILTER_PARTICLE_TASK, NULL));
+}
+
+static void
+ufo_filter_particle_task_setup (UfoTask *task,
+                                UfoResources *resources,
+                                GError **error)
+{
+}
+
+static void
+ufo_filter_particle_task_get_requisition (UfoTask *task,
+                                          UfoBuffer **inputs,
+                                          UfoRequisition *requisition)
+{
+    UfoFilterParticleTaskPrivate *priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE (task);
+    requisition->n_dims = 1;
+    ufo_buffer_get_requisition(inputs[0], requisition);
+    size_t num_pixels = requisition->dims[0] * requisition->dims[1];
+
+    /* Allocate internal buffer to hold copy of image */
+    /* Reuse buffer when size is the smae */
+    if (priv->num_pixels != num_pixels) {
+        priv->num_pixels = num_pixels;
+        priv->img = realloc (priv->img, sizeof (unsigned) * num_pixels);
+    }
+
+    /* Output size varies according to  data whithin the input */
+    /* I don't know the output size, I'll allocate it myself once I have
+     * processed the data */
+    requisition->dims[0] = 0;
+    requisition->n_dims = 1;
+}
+
+static guint
+ufo_filter_particle_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_filter_particle_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_filter_particle_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static void
+part_cen_radius(float *img, unsigned *bin, unsigned nc, unsigned nr,
+                unsigned pc, double *cx, double *cy, double *rr)
+{
+    /* nr == number of rows */
+    /* nc == number of columns */
+    /* pc == pixel count, (i-e nb of detected rings) */
+    /* cx, cy, rr, (x, y, r) coordinates of each ring */
+    unsigned i,j,k, itop, ibot, jleft, jright;
+    unsigned b;
+    unsigned *tops,*bots,*lefts,*rights;
+    float cenx, radx, ceny, rady, wei;
+
+    tops = malloc((pc+1)*sizeof(unsigned) * 100);
+    bots = malloc((pc+1)*sizeof(unsigned) * 100);
+    lefts = malloc((pc+1)*sizeof(unsigned) * 100);
+    rights = malloc((pc+1)*sizeof(unsigned) * 100);
+
+    for (k=0;k<=pc;k++) {
+        tops[k]= (unsigned)nr;  bots[k]=0;  lefts[k]=(unsigned)nc; rights[k]=0;
+    }
+
+    /* Find extremities of each clustered pixels */
+    for (i=0;i<nr;i++) {
+        for(j=0;j<nc;j++) {
+            b = bin[i * nc + j];
+            if ( tops[b]>i ) tops[b]=i;
+            if ( bots[b]<i ) bots[b]=i;
+            if ( lefts[b]>j ) lefts[b]=j;
+            if ( rights[b]<j ) rights[b]=j;
+        }
+    }
+
+    /* calculate center of each clustered pixel using img to put weight on each
+     * candidate pixel, at the same time calculate radius of the cluster */
+    for (k = 0; k < pc; k++) {
+        itop=tops[k+1];     ibot=bots[k+1];
+        jleft=lefts[k+1];   jright=rights[k+1];
+
+        cenx=0.0;    radx=0.0;
+        ceny=0.0;    rady=0.0;
+        wei=0.0;
+
+        for (i=itop;i<=ibot;i++) {
+            for (j=jleft;j<=jright;j++)
+            {
+                if (bin[i * nc + j]==k+1) {
+                    float value = img[i * nc + j];
+                    cenx += (float) j * value;
+                    radx = radx + (float) (j * j) * value;
+                    ceny = ceny + (float) i * value;
+                    rady = rady + (float) (i * i) * value;
+                    wei = wei + value;
+                }
+            }
+        }
+        cenx=cenx/wei;
+        ceny=ceny/wei;
+        radx=sqrtf (fabsf ((radx/wei) - cenx*cenx)) + (float) 0.25;
+        rady=sqrtf (fabsf ((rady/wei) - ceny*ceny)) + (float) 0.25;
+
+        cx[k] = cenx+1;
+        cy[k] = ceny+1;
+        rr[k] = sqrt(radx*rady);
+    }
+
+    free(tops);  free(bots);  free(lefts);  free(rights);
+
+}
+
+static void
+pfind (float *input, unsigned *out, unsigned nc, unsigned nr, unsigned *pcount,
+       double threshold)
+{
+    unsigned pc;
+    unsigned ndup;
+    unsigned i,j,k;
+    int ncase;
+    char iodd;
+    unsigned ncc;
+    unsigned  *dupx,*dupy;
+    unsigned *dupindex, *binloc;
+    unsigned k1,k2,temp;
+ 
+    unsigned b;
+ 
+    pc=0;
+    ndup=0;
+ 
+    dupx = malloc (nVectMax * sizeof (unsigned));
+    dupy = malloc (nVectMax * sizeof (unsigned));
+
+    /* 0, 0 is top left of image */
+    for(j=0;j<nc;j++) {
+        for(i=0;i<nr;i++) {
+            /* Bit mask of surrounding pixels */
+            ncase=0;
+            /* Set if bottom left pixel is set */
+            iodd=0;
+            if(input[i * nc +  j] > threshold) {
+                /* Bottom left pixel */
+                if (i < (nr - 1) && j > 0 && input[(i + 1) * nc + (j - 1)] > threshold)
+                { ncase += 1; iodd = 1; }
+                /* Bottom right pixel */
+                if (i < (nr - 1) && j < (nc - 1) && input[(i + 1) * nc + (j + 1)] > threshold)
+                    ncase += 2;
+                /* Top right pixel */
+                if (i > 0 && j < (nc - 1) && input[(i - 1) * nc + (j + 1)] > threshold)
+                    ncase += 4;
+                /* Top left pixel */
+                if (i > 0 && j > 0 && input[(i - 1) * nc + (j - 1)] > threshold)
+                    ncase += 8;
+                /* Bottom pixel */
+                if (i < (nr - 1) && input[(i + 1) * nc + j] > threshold)
+                    ncase += 16;
+                /* Right pixel */
+                if (j < (nc - 1) && input[i * nc + (j + 1)] > threshold)
+                    ncase += 32;
+                /* Top pixel */
+                if (i > 0 && input[(i - 1) * nc + j] > threshold)
+                    ncase += 64;
+                /* Left pixel */
+                if (j > 0 && input[i * nc + (j - 1)] > threshold)
+                    ncase += 128;
+
+                /* Pixel has no : Bottom, Right, Top, Left */
+                if (ncase <= 15) {
+                    //  if (NO1 == 1)
+                    //    bin[i][j] = ++pc;
+                    //  else
+                    out[i * nc + j] = 0;
+                }
+                /* Pixel has no : Right, Top, Left */
+                /* and Pixel has : Bottom */
+                else if (ncase <= 31) {
+                    if(iodd == 1)
+                        out[i * nc + j] = out[(i + 1) * nc + (j - 1)]; /* Clone id */
+                    else
+                        out[i * nc + j] = ++pc;
+                }
+                /* Pixel has no : Top, Left, Bottom */
+                /* Pixel has : Right */
+                else if (ncase <= 47)
+                    out[i * nc + j] = ++pc;
+                /* Pixel has no : Top, Left */
+                /* Pixel has : Right, Bottom */
+                else if (ncase <= 63) {
+                    if (iodd == 1)
+                        out[i * nc + j] = out[(i + 1) * nc + (j - 1)];
+                    else
+                        out[i * nc + j] = ++pc;
+                }
+                /* Pixel has no : Left */
+                /* Pixel has : Top */
+                else if (ncase <= 127)
+                    out[i * nc + j] = out[(i - 1) * nc + j];
+                /* Pixel has : Left */
+                else
+                    out[i * nc + j] = out[i * nc + (j - 1)];
+
+
+                /* Bottom Left && Bottom && Top && not(Left) */
+                if ((ncase & 1) && (ncase & 16) && (ncase & 64) && !(ncase & 128)) {
+                    /* Find duplicates, and locate them by x, and y list */
+                    /* These duplicates shall be joined later */
+                    k1 = out[i * nc + j];
+                    /* 11
+                     *  1
+                     * 22
+                     * in this case, bottom left was processed later, which
+                     * explains the comparison */
+                    k2 = out[(i + 1) * nc + (j - 1)];
+                    dupx[ndup] = (k1 > k2 ? k2 : k1);
+                    dupy[ndup] = (k1 > k2 ? k1 : k2);
+                    ndup++;
+                }
+            }
+            else
+                out[i * nc + j] = 0;
+        }
+    }
+
+    dupindex = malloc ((unsigned) (pc + 1) * sizeof (unsigned));
+    /* binloc stores the 1D position (on row basis, like matlab) of each
+     * candidate block of pixel (Most left-top of block)
+     * or MAX_POS_VAL otherwise */
+    binloc = malloc ((unsigned) (pc + 1) * sizeof(unsigned));
+    for(k=0;k<=pc;k++) {
+        binloc[k]= nr * nc;
+        dupindex[k] = k;
+    }
+
+    for(j=0;j<nc;j++) {
+        for(i=0;i<nr;i++) {
+            b = out[i * nc + j];
+            if(b > 0 && binloc[b] > (j*nr+i) )
+                binloc[b]=j*nr+i;
+        }
+    }
+
+    for(k=0; k < ndup; k++) {
+        /* 11
+         *  1
+         * 22 in this case*/
+        if( binloc[dupx[k]] > binloc[dupy[k]] ) {
+            /* Dupx has always been afected the smaller id
+             * ==> with row wise basis, locations dupx > dupy is always false */
+            temp = dupx[k];   dupx[k]=dupy[k];    dupy[k]=temp;
+        }
+    }
+
+    /* Sort dupx to be in crescent order */
+    /* i-e
+     * 22
+     *  2
+     * 33
+     *  3
+     *111
+     * Here dupx = { 2, 1 } ==> Not in crescent order */
+    if (ndup != 0) {
+        for (k=0;k<ndup-1;k++) {
+            for(j=k+1;j<ndup;j++) {
+                if(binloc[dupx[k]]>binloc[dupx[j]]) {
+                    temp=dupx[k];
+                    dupx[k]=dupx[j];
+                    dupx[j]=temp;
+                    temp=dupy[k];
+                    dupy[k]=dupy[j];
+                    dupy[j]=temp;
+                }
+                else if( dupx[k]== dupx[j] && binloc[dupy[k]] > binloc[dupy[j]]) {
+                    temp=dupy[k];
+                    dupy[k]=dupy[j];
+                    dupy[j]=temp;
+                }
+            }
+        }
+    }
+
+    /* dupindex gives the new index for each index, when two index are duplicates
+     * dupindex will affect those two indexes a same value */
+    for(k=0;k<ndup;k++) {
+        k1=dupx[k];     k2=dupy[k];
+        for(j=k;j<ndup;j++) {
+            if(dupy[j]== dupy[k]||dupx[j]== dupy[k]) {
+                dupindex[ dupx[j] ]=dupindex[k1];
+                dupindex[ dupy[j] ]=dupindex[k1];
+            }
+        }
+    }
+
+    /* copy dupindex in dupx */
+    for (k=0;k<=pc;k++) {
+        dupx[k]= dupindex[k];
+        dupy[k]=0;
+    }
+    /* order dupx in crescent order, is it not already ordered though ?!? */
+    for (k=0;k<=pc;k++) {
+        for(j=k+1;j<=pc;j++) {
+            if(dupx[k]>dupx[j]) {
+                temp=dupx[k];
+                dupx[k]=dupx[j];
+                dupx[j]=temp;
+            }
+        }
+    }
+
+    ncc=0;
+    dupy[0]=0;
+    /* map dupindex values to values from 1 to number_of_detected_ring, this
+     * allows to have a constant step of one, eg: converting {1,3,6} -> {1,2,3} */
+    for(k=1;k<=pc;k++) {
+        if(dupx[k]>dupx[k-1])
+            dupy[ dupx[k] ] = ++ncc;
+    }
+
+    *pcount=ncc;
+    /* set dupindex to its newly mapped value, (making dupindex have a step of
+     * one */
+    for (k=0;k<=pc;k++)
+        dupindex[k]= dupy[ dupindex[k] ];
+
+    /* Make all duplicates have the same ID */
+    for(i=0;i<nr;i++) {
+        for(j=0;j<nc;j++)
+            out[i * nc + j]= dupindex[ out[i * nc + j] ];
+    }
+
+    free(dupx);
+    free(dupy);
+    free(dupindex);
+    free(binloc);
+
+    return;
+}
+
+static void
+get_radius_metadata(UfoBuffer *src, unsigned *radius)
+{
+    GValue *value;
+
+    value = ufo_buffer_get_metadata(src, "radius");
+    *radius = g_value_get_uint(value);
+}
+
+static double
+get_max_value(float* img, size_t width, size_t height)
+{
+    double res = 0;
+
+    for (size_t y = 0; y < height; ++y) {
+        for (size_t x = 0; x < width; ++x) {
+            if (res < img[y * width + x])
+                res = img[y * width + x];
+        }
+    }
+
+    return res;
+}
+
+static unsigned
+min(unsigned a, unsigned b)
+{
+    return a > b ? b : a;
+}
+
+static gboolean
+ufo_filter_particle_task_process (UfoTask *task,
+                                  UfoBuffer **inputs,
+                                  UfoBuffer *output,
+                                  UfoRequisition *requisition)
+{
+    (void) requisition;
+    UfoFilterParticleTaskPrivate *priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE (task);
+    UfoRequisition req;
+    ufo_buffer_get_requisition(inputs[0], &req);
+    unsigned pcount = 0;
+    float *input = ufo_buffer_get_host_array(inputs[0], NULL);
+    unsigned radius;
+    get_radius_metadata(inputs[0], &radius);
+
+    /* Give unique ID to each cluster of pixels */
+    double max_value = get_max_value(input, req.dims[0], req.dims[1]);
+    double threshold = priv->threshold * max_value;
+
+    // When max_value is too low, it is unlikely for there to be any rings
+    if (max_value < priv->min) {
+        g_print("Filter_particle: min : Ignoring radius %u. Inputs max value is"
+                " %f, at least %f expected\n", radius, max_value, priv->min);
+        UfoRequisition out_req = {
+            .n_dims = 1,
+            .dims[0] = 1
+        };
+        ufo_buffer_resize (output, &out_req);
+        float *out = ufo_buffer_get_host_array(output, NULL);
+        *out = 0;
+        return TRUE;
+    }
+
+    pfind (input, priv->img, (unsigned) req.dims[0], (unsigned) req.dims[1], &pcount, threshold);
+
+    /* Compute center of each cluster and store coordinates in cx, cy */
+    double *cx = malloc (sizeof (double) * pcount);
+    double *cy = malloc (sizeof (double) * pcount);
+    double *rr = malloc (sizeof (double) * pcount);
+    part_cen_radius(input, priv->img, (unsigned) req.dims[0],
+                    (unsigned) req.dims[1], pcount, cx, cy, rr);
+
+    /* a list of pcount UfoRingCoordinate are generated in output */
+    /* first element is the number of rings detected */
+    UfoRequisition out_req = {
+        .n_dims = 1,
+        .dims[0] = 1 + pcount * sizeof (UfoRingCoordinate) / sizeof (float),
+    };
+    ufo_buffer_resize (output, &out_req);
+    float *out = ufo_buffer_get_host_array(output, NULL);
+    UfoRingCoordinate *res = (UfoRingCoordinate *) (out + 1);
+    *out = (float) pcount;
+
+    for (unsigned i = 0; i < pcount; ++i) {
+        res[i].x = (float) cx[i];
+        res[i].y = (float) cy[i];
+        res[i].r = (float) radius;
+        unsigned x =  min((unsigned)round(cx[i]), (unsigned) (req.dims[0] - 1));
+        unsigned y =  min((unsigned)round(cy[i]), (unsigned) (req.dims[1] - 1));
+        res[i].intensity = input[y * req.dims[0] + x];
+    }
+
+    return TRUE;
+}
+
+static void
+ufo_filter_particle_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoFilterParticleTaskPrivate *priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MIN:
+            priv->min = g_value_get_float(value);
+            break;
+        case PROP_THRESHOLD:
+            priv->threshold = g_value_get_float(value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_filter_particle_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoFilterParticleTaskPrivate *priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_MIN:
+            g_value_set_float(value, priv->min);
+            break;
+        case PROP_THRESHOLD:
+            g_value_set_float(value, priv->threshold);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_filter_particle_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_filter_particle_task_parent_class)->finalize (object);
+    UfoFilterParticleTaskPrivate *priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE (object);
+    free (priv->img);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_filter_particle_task_setup;
+    iface->get_num_inputs = ufo_filter_particle_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_filter_particle_task_get_num_dimensions;
+    iface->get_mode = ufo_filter_particle_task_get_mode;
+    iface->get_requisition = ufo_filter_particle_task_get_requisition;
+    iface->process = ufo_filter_particle_task_process;
+}
+
+static void
+ufo_filter_particle_task_class_init (UfoFilterParticleTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_filter_particle_task_set_property;
+    gobject_class->get_property = ufo_filter_particle_task_get_property;
+    gobject_class->finalize = ufo_filter_particle_task_finalize;
+
+    properties[PROP_MIN] =
+        g_param_spec_float ("min",
+            "Ignore image when it's maximal value is less than min",
+            "Ignore image when it's maximal value is less than min",
+            0, 1, 0.125f,
+            G_PARAM_READWRITE);
+    properties[PROP_THRESHOLD] =
+        g_param_spec_float ("threshold",
+            "Set threshold value which is relative to the image maximum value",
+            "Set threshold value which is relative to the image maximum value",
+            0, 1, 0.8f,
+            G_PARAM_READWRITE);
+
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoFilterParticleTaskPrivate));
+}
+
+static void
+ufo_filter_particle_task_init(UfoFilterParticleTask *self)
+{
+    self->priv = UFO_FILTER_PARTICLE_TASK_GET_PRIVATE(self);
+    self->priv->img = NULL;
+    self->priv->num_pixels = 0;
+    self->priv->min = 0.125f;
+    self->priv->threshold = 0.8f;
+}
diff --git a/src/ufo-filter-particle-task.h b/src/ufo-filter-particle-task.h
new file mode 100644
index 0000000..6220a0d
--- /dev/null
+++ b/src/ufo-filter-particle-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_FILTER_PARTICLE_TASK_H
+#define __UFO_FILTER_PARTICLE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_FILTER_PARTICLE_TASK             (ufo_filter_particle_task_get_type())
+#define UFO_FILTER_PARTICLE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_PARTICLE_TASK, UfoFilterParticleTask))
+#define UFO_IS_FILTER_PARTICLE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_PARTICLE_TASK))
+#define UFO_FILTER_PARTICLE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_PARTICLE_TASK, UfoFilterParticleTaskClass))
+#define UFO_IS_FILTER_PARTICLE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_PARTICLE_TASK))
+#define UFO_FILTER_PARTICLE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_PARTICLE_TASK, UfoFilterParticleTaskClass))
+
+typedef struct _UfoFilterParticleTask           UfoFilterParticleTask;
+typedef struct _UfoFilterParticleTaskClass      UfoFilterParticleTaskClass;
+typedef struct _UfoFilterParticleTaskPrivate    UfoFilterParticleTaskPrivate;
+
+/**
+ * UfoFilterParticleTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoFilterParticleTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoFilterParticleTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoFilterParticleTaskPrivate *priv;
+};
+
+/**
+ * UfoFilterParticleTaskClass:
+ *
+ * #UfoFilterParticleTask class
+ */
+struct _UfoFilterParticleTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_filter_particle_task_new       (void);
+GType     ufo_filter_particle_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-get-dup-circ-task.c b/src/ufo-get-dup-circ-task.c
new file mode 100644
index 0000000..e8a45e5
--- /dev/null
+++ b/src/ufo-get-dup-circ-task.c
@@ -0,0 +1,234 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <math.h> /* sqrt */
+#include <string.h> /* memcpy */
+
+#include "ufo-get-dup-circ-task.h"
+#include "ufo-ring-coordinates.h"
+
+struct _UfoGetDupCircTaskPrivate {
+    float threshold;
+    UfoRingCoordinate *coord;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoGetDupCircTask, ufo_get_dup_circ_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_GET_DUP_CIRC_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_GET_DUP_CIRC_TASK, UfoGetDupCircTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_THRESHOLD,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_get_dup_circ_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_GET_DUP_CIRC_TASK, NULL));
+}
+
+static void
+ufo_get_dup_circ_task_setup (UfoTask *task,
+                             UfoResources *resources,
+                             GError **error)
+{
+}
+
+static void
+ufo_get_dup_circ_task_get_requisition (UfoTask *task,
+                                       UfoBuffer **inputs,
+                                       UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition (inputs[0], requisition);
+}
+
+static guint
+ufo_get_dup_circ_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_get_dup_circ_task_get_num_dimensions (UfoTask *task,
+                                          guint input)
+{
+    return 1;
+}
+
+static UfoTaskMode
+ufo_get_dup_circ_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static void
+merge_rings (UfoBuffer *output, UfoGetDupCircTaskPrivate *priv, unsigned nb_elt)
+{
+    float *res = ufo_buffer_get_host_array (output, NULL);
+    UfoRingCoordinate *coord = (UfoRingCoordinate *) (res + 1);
+    unsigned nb_coord = 0;
+    for (unsigned i = 0; i < nb_elt; ++i) {
+        /* We set rings to a radius of 0 when they have already been merged */
+        if (priv->coord[i].r == 0)
+            continue;
+        coord[nb_coord] = priv->coord[i];
+        float nb_merged_elt = 1;
+        for (unsigned j = i + 1; j < nb_elt; ++j) {
+            if (priv->coord[j].r == 0)
+                continue;
+            float distance = sqrtf ((priv->coord[j].x - priv->coord[i].x) *
+                                    (priv->coord[j].x - priv->coord[i].x) +
+                                    (priv->coord[j].y - priv->coord[i].y) *
+                                    (priv->coord[j].y - priv->coord[i].y));
+            float radius_diff = priv->coord[j].r - priv->coord[i].r;
+            radius_diff = radius_diff < 0 ? -radius_diff : radius_diff;
+            if (distance < priv->threshold && radius_diff < priv->threshold) {
+                coord[nb_coord].x += priv->coord[j].x;
+                coord[nb_coord].y += priv->coord[j].y;
+                coord[nb_coord].r += priv->coord[j].r;
+                /* Say that the j-th ring has been merged */
+                priv->coord[j].r = 0;
+                ++nb_merged_elt;
+            }
+        }
+        coord[nb_coord].x /= nb_merged_elt;
+        coord[nb_coord].y /= nb_merged_elt;
+        coord[nb_coord].r /= nb_merged_elt;
+        /* When merging occurs, search again for new candidates */
+        if (nb_merged_elt > 1) {
+            priv->coord[i] = coord[nb_coord];
+            --i;
+        }
+        else
+            ++nb_coord;
+    }
+    *res = (float) nb_coord;
+}
+
+static gboolean
+ufo_get_dup_circ_task_process (UfoTask *task,
+                               UfoBuffer **inputs,
+                               UfoBuffer *output,
+                               UfoRequisition *requisition)
+{
+    UfoGetDupCircTaskPrivate *priv = UFO_GET_DUP_CIRC_TASK_GET_PRIVATE (task);
+    float *input = ufo_buffer_get_host_array (inputs[0], NULL);
+    unsigned nb_elt = (unsigned) *input;
+    UfoRingCoordinate *coord = (UfoRingCoordinate *) (input + 1);
+    priv->coord = g_malloc (sizeof (UfoRingCoordinate) * nb_elt);
+    memcpy (priv->coord, coord, nb_elt * sizeof (UfoRingCoordinate));
+    merge_rings (output, priv, nb_elt);
+    return TRUE;
+}
+
+static void
+ufo_get_dup_circ_task_set_property (GObject *object,
+                                    guint property_id,
+                                    const GValue *value,
+                                    GParamSpec *pspec)
+{
+    UfoGetDupCircTaskPrivate *priv = UFO_GET_DUP_CIRC_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_THRESHOLD:
+            priv->threshold = g_value_get_float (value);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_get_dup_circ_task_get_property (GObject *object,
+                                    guint property_id,
+                                    GValue *value,
+                                    GParamSpec *pspec)
+{
+    UfoGetDupCircTaskPrivate *priv = UFO_GET_DUP_CIRC_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_THRESHOLD:
+            g_value_set_float (value, priv->threshold);
+            break;
+
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_get_dup_circ_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_get_dup_circ_task_parent_class)->finalize (object);
+    UfoGetDupCircTaskPrivate *priv = UFO_GET_DUP_CIRC_TASK_GET_PRIVATE (object);
+    g_free (priv->coord);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_get_dup_circ_task_setup;
+    iface->get_num_inputs = ufo_get_dup_circ_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_get_dup_circ_task_get_num_dimensions;
+    iface->get_mode = ufo_get_dup_circ_task_get_mode;
+    iface->get_requisition = ufo_get_dup_circ_task_get_requisition;
+    iface->process = ufo_get_dup_circ_task_process;
+}
+
+static void
+ufo_get_dup_circ_task_class_init (UfoGetDupCircTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_get_dup_circ_task_set_property;
+    gobject_class->get_property = ufo_get_dup_circ_task_get_property;
+    gobject_class->finalize = ufo_get_dup_circ_task_finalize;
+
+    properties[PROP_THRESHOLD] =
+        g_param_spec_float ("threshold",
+                            "Give maximum ring distance, and radius difference",
+                            "Give maximum ring distance, and radius difference",
+                            1, G_MAXFLOAT, 10,
+                            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoGetDupCircTaskPrivate));
+}
+
+static void
+ufo_get_dup_circ_task_init(UfoGetDupCircTask *self)
+{
+    self->priv = UFO_GET_DUP_CIRC_TASK_GET_PRIVATE(self);
+    self->priv->threshold = 10;
+    self->priv->coord = NULL;
+}
diff --git a/src/ufo-get-dup-circ-task.h b/src/ufo-get-dup-circ-task.h
new file mode 100644
index 0000000..138fec5
--- /dev/null
+++ b/src/ufo-get-dup-circ-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_GET_DUP_CIRC_TASK_H
+#define __UFO_GET_DUP_CIRC_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_GET_DUP_CIRC_TASK             (ufo_get_dup_circ_task_get_type())
+#define UFO_GET_DUP_CIRC_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_GET_DUP_CIRC_TASK, UfoGetDupCircTask))
+#define UFO_IS_GET_DUP_CIRC_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_GET_DUP_CIRC_TASK))
+#define UFO_GET_DUP_CIRC_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_GET_DUP_CIRC_TASK, UfoGetDupCircTaskClass))
+#define UFO_IS_GET_DUP_CIRC_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_GET_DUP_CIRC_TASK))
+#define UFO_GET_DUP_CIRC_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_GET_DUP_CIRC_TASK, UfoGetDupCircTaskClass))
+
+typedef struct _UfoGetDupCircTask           UfoGetDupCircTask;
+typedef struct _UfoGetDupCircTaskClass      UfoGetDupCircTaskClass;
+typedef struct _UfoGetDupCircTaskPrivate    UfoGetDupCircTaskPrivate;
+
+/**
+ * UfoGetDupCircTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoGetDupCircTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoGetDupCircTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoGetDupCircTaskPrivate *priv;
+};
+
+/**
+ * UfoGetDupCircTaskClass:
+ *
+ * #UfoGetDupCircTask class
+ */
+struct _UfoGetDupCircTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_get_dup_circ_task_new       (void);
+GType     ufo_get_dup_circ_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-multi-search-task.c b/src/ufo-multi-search-task.c
new file mode 100644
index 0000000..1996f19
--- /dev/null
+++ b/src/ufo-multi-search-task.c
@@ -0,0 +1,513 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <float.h>
+#include <math.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include "ufo-multi-search-task.h"
+#include "ufo-ring-coordinates.h"
+
+
+struct _UfoMultiSearchTaskPrivate {
+    /* number of elements desired in computations of polynomial */
+    unsigned radii_range;
+    float threshold;
+    unsigned displacement;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoMultiSearchTask, ufo_multi_search_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_MULTI_SEARCH_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_MULTI_SEARCH_TASK, UfoMultiSearchTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_RADII_RANGE,
+    PROP_THRESHOLD,
+    PROP_DISPLACEMENT,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_multi_search_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_MULTI_SEARCH_TASK, NULL));
+}
+
+static void
+ufo_multi_search_task_setup (UfoTask *task,
+                             UfoResources *resources,
+                             GError **error)
+{
+}
+
+static void
+ufo_multi_search_task_get_requisition (UfoTask *task,
+                                       UfoBuffer **inputs,
+                                       UfoRequisition *requisition)
+{
+    /* input[0] : contrasted image */
+    /* input[1] : coordinate list */
+    ufo_buffer_get_requisition (inputs[1], requisition);
+}
+
+static guint
+ufo_multi_search_task_get_num_inputs (UfoTask *task)
+{
+    return 2;
+}
+
+static guint
+ufo_multi_search_task_get_num_dimensions (UfoTask *task,
+                                          guint input)
+{
+    return 1;
+}
+
+static UfoTaskMode
+ufo_multi_search_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static int
+min (int l, int r)
+{
+    if (l > r)
+        return r;
+
+    return l;
+}
+
+static int
+max (int l, int r)
+{
+    if (l > r)
+        return l;
+    return r;
+}
+
+static void
+get_coords(int *left, int *right, int *top, int *bot, int rad,
+           UfoRequisition *req, UfoRingCoordinate *center)
+{
+    int l = (int) roundf (center->x - (float) rad);
+    int r = (int) roundf (center->x + (float) rad);
+    int t = (int) roundf (center->y - (float) rad);
+    int b = (int) roundf (center->y + (float) rad);
+    *left = max (l, 0);
+    *right = min (r, (int) (req->dims[0]) - 1);
+    // Bottom most point is req->dims[1]
+    *top = max (t, 0);
+    // Top most point is 0
+    *bot = min (b, (int) (req->dims[1]) - 1);
+
+}
+
+/* Sum each pixel of ring from center with a radius of r */
+static float
+compute_intensity (UfoBuffer *ufo_image, UfoRingCoordinate *center, int r)
+{
+    float intensity = 0;
+    UfoRequisition req;
+    ufo_buffer_get_requisition(ufo_image, &req);
+    float *image = ufo_buffer_get_host_array(ufo_image, NULL);
+    int x_center = (int) roundf (center->x);
+    int y_center = (int) roundf (center->y);
+    int left, right, top, bot;
+    get_coords(&left, &right, &top, &bot, r, &req, center);
+    unsigned counter = 0;
+    for (int y = top; y <= bot; ++y) {
+        for (int x = left; x <= right; ++x) {
+            int xx = (x - x_center) * (x - x_center);
+            int yy = (y - y_center) * (y - y_center);
+            /* Check if point is on ring */
+            if (fabs (sqrtf ((float) (xx + yy)) - (float) r) < 0.5) {
+                intensity += image[x + y * (int) req.dims[0]];
+                ++counter;
+            }
+        }
+    }
+    return intensity / (float) counter;
+}
+
+/*  X_1^order ..... X^0
+ *  |
+ *  |
+ *  X_i^order ..... X^0
+ */
+static float *
+vandermonde_new (unsigned x, unsigned nb_elt, unsigned order)
+{
+    float *vandermonde = malloc (sizeof (float) * nb_elt * (order + 1));
+    for (unsigned j = 0; j <= order; ++j) {
+        for (unsigned i = 0; i < nb_elt; ++i) {
+            vandermonde[i * (order + 1) + j] = powf ((float) (x + i),
+                                                     (float) (order - j));
+        }
+    }
+    return vandermonde;
+}
+
+/* Computes the projectios of vector A(:, j) over e */
+static void
+compute_projection (float *e, float *A, unsigned j, unsigned row,
+                    unsigned column, float *dst)
+{
+    float sc1 = 0;
+    float sc2 = 0;
+    /* Scalar product */
+    for (unsigned i  = 0; i < row; ++i) {
+        sc1 += e[i] * A[i * column + j];
+        sc2 += e[i] * e[i];
+    }
+
+    /* Compute projection */
+    for (unsigned i  = 0; i < row; ++i) {
+        dst[i] = e[i] * sc1 / sc2;
+    }
+}
+
+static float *
+Gram_Schmidt_U(float *A, unsigned row, unsigned column)
+{
+    float *U = malloc (sizeof (float) * row * column);
+    float *norms = malloc (sizeof (float) * column);
+    float *proj_sum = malloc (sizeof (float) * row);
+    float *e = malloc (sizeof (float) * row);
+    float *proj = malloc (sizeof (float) * row);
+    for (unsigned j = 0; j < column; ++j) {
+
+        /* Compute the sum of projections of column j over each vector in the
+         * orthonrmal basis */
+        memset (proj_sum, 0, sizeof (float) * row);
+        for (int k = 0; k < (int) j; ++k) {
+            /* k represents a column */
+            for (unsigned i = 0; i < row; ++i) {
+                e[i] = U[i * column + (unsigned) k] / norms[k];
+            }
+            compute_projection (e, A, j, row, column, proj);
+            for (unsigned i = 0; i < row; ++i) {
+                proj_sum[i] += proj[i];
+            }
+        }
+
+        /* compute column j of matrix U */
+        for (unsigned i = 0; i < row; ++i) {
+            U[i * column + j] = A[i * column + j] - proj_sum[i];
+        }
+
+        /* compute new norm of column j */
+        norms[j] = 0;
+        for (unsigned i = 0; i < row; ++i) {
+            norms[j] += U[i * column + j] * U[i * column + j];
+        }
+        norms[j] = sqrtf (norms[j]);
+    }
+
+    free (norms);
+    free (proj_sum);
+    free (e);
+    free (proj);
+
+    return U;
+}
+
+static float *
+Gram_Schmidt_Q (float *A, unsigned row, unsigned column)
+{
+    float *Q = malloc (sizeof (float) * row * column);
+    float *U = Gram_Schmidt_U(A, row, column);
+    for (unsigned j = 0; j < column; ++j) {
+        float norm = 0;
+        for (unsigned i = 0; i < row; ++i) {
+            norm += U[i * column + j] * U[i * column + j];
+        }
+        norm = sqrtf (norm);
+        for (unsigned i = 0; i < row; ++i) {
+            Q[i * column + j] = U[i * column + j] / norm;
+        }
+    }
+    free (U);
+    return Q;
+}
+/* Transpose first matrix, and multiply to second matrix */
+/* column_Q is the number of column in Q befor being transposed */
+static float *
+matrix_transpose_mul2(float *Q, float *A, unsigned column_Q,
+                            unsigned row, unsigned column_A)
+{
+    float *res = calloc (1, sizeof (float) * column_Q * column_A);
+    for (unsigned i = 0; i < column_Q; ++i) {
+        for (unsigned j = 0; j < column_A; ++j) {
+            for (unsigned k = 0; k < row; ++k) {
+                res[i * column_A + j] += Q[k * column_Q + i] * A[k * column_A + j];
+            }
+        }
+    }
+    return res;
+}
+
+/* Transpose first matrix, and multiply to second matrix */
+static float *
+matrix_transpose_mul(float *Q, float *A, unsigned row, unsigned column)
+{
+    float *res = calloc (1, sizeof (float) * column * column);
+    for (unsigned i = 0; i < column; ++i) {
+        for (unsigned j = 0; j < column; ++j) {
+            for (unsigned k = 0; k < row; ++k) {
+                res[i * column + j] += Q[k * column + i] * A[k * column + j];
+            }
+        }
+    }
+    return res;
+}
+/* P(r_min) = values[0] */
+/* P(r_min + 1) = values[1] */
+/* P(r_min + nb_elt - 1) = values[nb_elt - 1] */
+/* The step is always of 1 */
+static void
+polyfit (float *values, unsigned nb_elt, unsigned r_min,
+         float *a, float *b, float *c)
+{
+    /* We have as many column as the order + 1 of the polynomial */
+    unsigned column = 3;
+    /* Compute vandermonde matrix */
+    float *vandermonde = vandermonde_new(r_min, nb_elt, 2);
+
+    /* Compute V = QR using Gram-Schmidt method */
+    float *Q = Gram_Schmidt_Q(vandermonde, nb_elt, column);
+    float *R = matrix_transpose_mul (Q, vandermonde, nb_elt, column);
+    /* Q' x Values */
+    float *Qtxy = matrix_transpose_mul2 (Q, values, column, nb_elt, 1);
+
+    *c = Qtxy[2] / R[2 * column + 2];
+    *b = (Qtxy[1] - R[1 * column + 2] * *c) / R[1 * column + 1];
+    *a = (Qtxy[0] - R[2] * *c - R[1] * *b) / R[0];
+
+    free (vandermonde);
+    free (Q);
+    free(R);
+    free(Qtxy);
+}
+
+/* From a given image, vary the radius size and compare the variation of the
+ * intensity for each radii.  From these intensities compute a polynomial P(r)
+ * where r represents a radius */
+static void
+create_profile_advanced (UfoMultiSearchTaskPrivate *priv, UfoBuffer *image,
+                         UfoRingCoordinate *center, float *a, float *b, float *c)
+{
+    unsigned min_rad = 1;
+    if (center->r > priv->radii_range)
+        min_rad = (unsigned) center->r - priv->radii_range;
+    unsigned max_rad = (unsigned) center->r + priv->radii_range;
+    /* The value associated to each radius */
+    float values[max_rad - min_rad + 1];
+
+    unsigned pic_idx = 0;
+    values[pic_idx] = 0;
+    float max_a = 0;
+    int displacement = (int) priv->displacement;
+    UfoRingCoordinate copy = *center;
+    for (int x = -displacement; x <= displacement; ++x) {
+        for (int y = -displacement; y <= displacement; ++y) {
+            UfoRingCoordinate urc =
+            { copy.x + (float) x, copy.y + (float) y, copy.r, 0.0f, 0.0f };
+            for (unsigned r = 0; r <= max_rad - min_rad; ++r) {
+                values[r] = compute_intensity (image, &urc, (int) (r + min_rad));
+                if (values[r] > values[pic_idx]) {
+                    pic_idx = r;
+                }
+            }
+
+            polyfit (values, max_rad - min_rad + 1, min_rad, a, b, c);
+            if (*a <= max_a) {
+                center->x = urc.x;
+                center->y = urc.y;
+                center->r = -*b / (2.0f * *a);
+                max_a = *a;
+            }
+        }
+    }
+    *a = max_a;
+}
+
+static char
+center_search (UfoMultiSearchTaskPrivate *priv, UfoBuffer *image,
+               UfoRingCoordinate *src, UfoRingCoordinate *dst)
+{
+    float a, b, c;
+    /* Compute polynomial aX^2 + bX + c */
+    create_profile_advanced(priv, image, src, &a, &b, &c);
+    /* When contrast is too low, we delete ring */
+    /* A represents the steepness of the polynomial, the more steep i-e the more
+     * negative a is, the more contrast we have.  The ideal function is a
+     * dirac */
+    if (a <= -priv->threshold) {
+        *dst = *src;
+        dst->contrast = a;
+        return 1;
+    }
+
+    return 0;
+}
+
+static gboolean
+ufo_multi_search_task_process (UfoTask *task,
+                               UfoBuffer **inputs,
+                               UfoBuffer *output,
+                               UfoRequisition *requisition)
+{
+    (void) requisition;
+    URCS *src = (URCS *) ufo_buffer_get_host_array(inputs[1], NULL);
+    URCS *dst = (URCS *) ufo_buffer_get_host_array(output, NULL);
+    unsigned nb_elt = (unsigned) src->nb_elt;
+    UfoMultiSearchTaskPrivate *priv = UFO_MULTI_SEARCH_TASK_GET_PRIVATE (task);
+
+    unsigned del_count = 0;
+    /* Check rings contrast, if it's too low, delete it */
+    for (unsigned idx = 0; idx < nb_elt; ++idx) {
+        if (!center_search(priv, inputs[0], &src->coord[idx],
+                           &dst->coord[idx - del_count]))
+            del_count++;
+    }
+    dst->nb_elt = src->nb_elt - (float) del_count;
+    return TRUE;
+}
+
+static void
+ufo_multi_search_task_set_property (GObject *object,
+                                    guint property_id,
+                                    const GValue *value,
+                                    GParamSpec *pspec)
+{
+    UfoMultiSearchTaskPrivate *priv = UFO_MULTI_SEARCH_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_RADII_RANGE:
+            priv->radii_range = g_value_get_uint (value);
+            break;
+        case PROP_THRESHOLD:
+            priv->threshold = g_value_get_float (value);
+            break;
+        case PROP_DISPLACEMENT:
+            priv->displacement = g_value_get_uint (value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_multi_search_task_get_property (GObject *object,
+                                    guint property_id,
+                                    GValue *value,
+                                    GParamSpec *pspec)
+{
+    UfoMultiSearchTaskPrivate *priv = UFO_MULTI_SEARCH_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_RADII_RANGE:
+            g_value_set_uint (value, priv->radii_range);
+            break;
+        case PROP_THRESHOLD:
+            g_value_set_float (value, priv->threshold);
+            break;
+        case PROP_DISPLACEMENT:
+            g_value_set_uint (value, priv->displacement);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_multi_search_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_multi_search_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_multi_search_task_setup;
+    iface->get_num_inputs = ufo_multi_search_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_multi_search_task_get_num_dimensions;
+    iface->get_mode = ufo_multi_search_task_get_mode;
+    iface->get_requisition = ufo_multi_search_task_get_requisition;
+    iface->process = ufo_multi_search_task_process;
+}
+
+static void
+ufo_multi_search_task_class_init (UfoMultiSearchTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_multi_search_task_set_property;
+    gobject_class->get_property = ufo_multi_search_task_get_property;
+    gobject_class->finalize = ufo_multi_search_task_finalize;
+
+    properties[PROP_RADII_RANGE] =
+        g_param_spec_uint ("radii_range",
+                          "Gives the radius scanning range",
+                          "Gives the radius scanning range",
+                          0, G_MAXUINT, 3,
+                          G_PARAM_READWRITE);
+
+    properties[PROP_THRESHOLD] =
+        g_param_spec_float ("threshold",
+                            "Give minimum contrast a ring should have",
+                            "Give minimum contrast a ring should have",
+                            0, G_MAXFLOAT, 0.01f,
+                            G_PARAM_READWRITE);
+
+    properties[PROP_DISPLACEMENT] =
+        g_param_spec_uint ("displacement",
+                            "How much rings center can be displaced",
+                            "How much rings center can be displaced",
+                            0, G_MAXUINT, 2,
+                            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoMultiSearchTaskPrivate));
+}
+
+static void
+ufo_multi_search_task_init(UfoMultiSearchTask *self)
+{
+    self->priv = UFO_MULTI_SEARCH_TASK_GET_PRIVATE(self);
+    self->priv->radii_range = 3;
+    self->priv->threshold = 0.01f;
+    self->priv->displacement = 2;
+}
diff --git a/src/ufo-multi-search-task.h b/src/ufo-multi-search-task.h
new file mode 100644
index 0000000..3b043a2
--- /dev/null
+++ b/src/ufo-multi-search-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_MULTI_SEARCH_TASK_H
+#define __UFO_MULTI_SEARCH_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_MULTI_SEARCH_TASK             (ufo_multi_search_task_get_type())
+#define UFO_MULTI_SEARCH_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_MULTI_SEARCH_TASK, UfoMultiSearchTask))
+#define UFO_IS_MULTI_SEARCH_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_MULTI_SEARCH_TASK))
+#define UFO_MULTI_SEARCH_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_MULTI_SEARCH_TASK, UfoMultiSearchTaskClass))
+#define UFO_IS_MULTI_SEARCH_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_MULTI_SEARCH_TASK))
+#define UFO_MULTI_SEARCH_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_MULTI_SEARCH_TASK, UfoMultiSearchTaskClass))
+
+typedef struct _UfoMultiSearchTask           UfoMultiSearchTask;
+typedef struct _UfoMultiSearchTaskClass      UfoMultiSearchTaskClass;
+typedef struct _UfoMultiSearchTaskPrivate    UfoMultiSearchTaskPrivate;
+
+/**
+ * UfoMultiSearchTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoMultiSearchTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoMultiSearchTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoMultiSearchTaskPrivate *priv;
+};
+
+/**
+ * UfoMultiSearchTaskClass:
+ *
+ * #UfoMultiSearchTask class
+ */
+struct _UfoMultiSearchTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_multi_search_task_new       (void);
+GType     ufo_multi_search_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-ordfilt-task.c b/src/ufo-ordfilt-task.c
new file mode 100644
index 0000000..96a2426
--- /dev/null
+++ b/src/ufo-ordfilt-task.c
@@ -0,0 +1,329 @@
+/*
+ * Copyright (C) 2011-2013 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 <limits.h>
+#include <glib.h>
+
+#include "ufo-ordfilt-task.h"
+#include "ufo-priv.h"
+
+struct _UfoOrdfiltTaskPrivate {
+    cl_kernel k_bitonic_ordfilt;
+    cl_kernel k_load_elements_from_patern;
+    size_t max_alloc_size;
+    gpointer context;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoOrdfiltTask, ufo_ordfilt_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_ORDFILT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_ORDFILT_TASK, UfoOrdfiltTaskPrivate))
+
+UfoNode *
+ufo_ordfilt_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_ORDFILT_TASK, NULL));
+}
+
+static void
+get_max_alloc_size (UfoResources *resources, UfoOrdfiltTaskPrivate *priv)
+{
+    priv->max_alloc_size = ULLONG_MAX;
+    GList *devices = ufo_resources_get_devices (resources);
+    GList *it;
+
+    g_list_for (devices, it) {
+        cl_device_id device = (cl_device_id) it->data;
+        size_t byte_count = 0;
+        size_t max_alloc_size = 0;
+        clGetDeviceInfo (device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (size_t), &max_alloc_size, &byte_count);
+        g_assert (sizeof (size_t) == byte_count);
+
+        if (max_alloc_size < priv->max_alloc_size)
+            priv->max_alloc_size = max_alloc_size;
+    }
+}
+
+static void
+ufo_ordfilt_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+    UfoOrdfiltTaskPrivate *priv;
+
+    priv = UFO_ORDFILT_TASK_GET_PRIVATE (task);
+    priv->context = ufo_resources_get_context (resources);
+    get_max_alloc_size (resources, priv);
+
+    priv->k_bitonic_ordfilt = ufo_resources_get_kernel (resources, "ordfilt.cl", "bitonic_ordfilt", error);
+
+    if (priv->k_bitonic_ordfilt != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_bitonic_ordfilt));
+
+    priv->k_load_elements_from_patern = ufo_resources_get_kernel (resources, "ordfilt.cl", "load_elements_from_pattern", error);
+
+    if (priv->k_load_elements_from_patern != NULL)
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->k_load_elements_from_patern));
+}
+
+static void
+ufo_ordfilt_task_get_requisition (UfoTask *task,
+                                  UfoBuffer **inputs,
+                                  UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition (inputs[0], requisition);
+}
+
+static guint
+ufo_ordfilt_task_get_num_inputs (UfoTask *task)
+{
+    return 2;
+}
+
+static guint
+ufo_ordfilt_task_get_num_dimensions (UfoTask *task,
+                                     guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_ordfilt_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+/* K is ordfilt kernel */
+static void
+launch_kernel_1D (cl_kernel k, UfoBuffer *ufo_src, UfoBuffer *ufo_dst,
+                  cl_command_queue cmd_queue, size_t num_elements,
+                  unsigned idx_offset, unsigned mod)
+{
+    cl_mem dst;
+    cl_mem src;
+    UfoRequisition requisition;
+    size_t global_work_size[1];
+    size_t local_work_size[1];
+
+    dst = ufo_buffer_get_device_array (ufo_dst, cmd_queue);
+    src = ufo_buffer_get_device_array (ufo_src, cmd_queue);
+
+    /* Power of 2 above num_elements */
+    size_t array_length = (size_t) ceil_power_of_two (num_elements);
+    cl_float low_threshold = 0.25;
+    cl_float high_threshold = 0.50;
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 0, sizeof (cl_mem), &src));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 1, sizeof (cl_mem), &dst));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 2, sizeof (cl_int), &num_elements));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 3, sizeof (cl_int), &array_length));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 4, sizeof (cl_float), &low_threshold));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 5, sizeof (cl_float), &high_threshold));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 6, sizeof (cl_float) * array_length, NULL));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (k, 7, sizeof (cl_uint), &idx_offset));
+
+    /* Launch the kernel over 1D grid */
+    ufo_buffer_get_requisition(ufo_src, &requisition);
+
+    /* buffer may have mod extra rows -> so dont do extra work */
+    global_work_size[0] = requisition.dims[0] * (requisition.dims[1] - mod) *
+                          (array_length / 2);
+    /* Power of two below number of elements to sort */
+    local_work_size[0] = array_length / 2;
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue,
+                                                       k,
+                                                       1, NULL, global_work_size,
+                                                       local_work_size,
+                                                       0, NULL, NULL));
+}
+
+static void
+launch_kernel_2D(cl_kernel kernel, UfoBuffer *ufo_src, UfoBuffer *ufo_pattern,
+                 UfoBuffer *ufo_dst, cl_command_queue cmd_queue, size_t dimension,
+                 size_t num_ones, unsigned height, unsigned y_offset, unsigned mod)
+{
+    cl_mem dst;
+    cl_mem src;
+    cl_mem pattern;
+    UfoRequisition requisition;
+    size_t global_work_size[2];
+    size_t local_work_size[2];
+    dst = ufo_buffer_get_device_array(ufo_dst, cmd_queue);
+    src = ufo_buffer_get_device_array(ufo_src, cmd_queue);
+    pattern = ufo_buffer_get_device_array(ufo_pattern, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof (cl_mem), &src));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof (cl_mem), &dst));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof (cl_mem), &pattern));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 3, sizeof (cl_int), &dimension));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 4, sizeof (cl_int), &num_ones));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 5, sizeof (cl_uint), &height));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 6, sizeof (cl_uint), &y_offset));
+
+    /* Launch the kernel over 2D grid, using dst requisition which reperesents a
+     * crop of the image */
+    ufo_buffer_get_requisition (ufo_dst, &requisition);
+    global_work_size[0] = requisition.dims[0];
+    /* Buffer may have mod extra rows, don't take them into account */
+    global_work_size[1] = requisition.dims[1] - mod;
+    unsigned y_worker_count = 32;
+    unsigned x_worker_count = 32;
+
+    while (global_work_size[1] % y_worker_count)
+        --y_worker_count;
+
+    while (global_work_size[0] % x_worker_count)
+        --x_worker_count;
+
+    local_work_size[0] = x_worker_count; /* Multiple of image_width=1080 */
+    local_work_size[1] = y_worker_count; /* Multiple of image_height=1280 */
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue,
+                                                       kernel,
+                                                       2, NULL, global_work_size,
+                                                       local_work_size,
+                                                       0, NULL, NULL));
+}
+
+static void
+get_ring_metadata(UfoBuffer *pattern, unsigned *number_ones, unsigned *radius)
+{
+    GValue *value;
+    value = ufo_buffer_get_metadata(pattern, "number_ones");
+    *number_ones = g_value_get_uint(value);
+
+    value = ufo_buffer_get_metadata(pattern, "radius");
+    *radius = g_value_get_uint(value);
+}
+
+static void
+compute_ordfilt (UfoOrdfiltTaskPrivate *priv, UfoBuffer *src, UfoBuffer *pattern,
+                 UfoBuffer *dst, cl_command_queue cmd_queue)
+{
+
+    UfoRequisition image_requisition;
+    UfoRequisition pattern_requisition;
+    ufo_buffer_get_requisition(src, &image_requisition);
+    ufo_buffer_get_requisition(pattern, &pattern_requisition);
+
+    unsigned number_ones;
+    unsigned radius;
+    get_ring_metadata(pattern, &number_ones, &radius);
+
+    unsigned height = (unsigned) image_requisition.dims[1];
+    unsigned width = (unsigned) image_requisition.dims[0];
+    /* Tells us in how many chunks to chop image */
+    unsigned iter_count = (unsigned) (1 + sizeof (float) * height * width * number_ones / (priv->max_alloc_size + 1));
+    unsigned y_offset = 0;
+
+    /* On first iteration process mod rows more, this is needed when the height
+     * of image is not divisible by iter_count */
+    unsigned mod = (unsigned) (image_requisition.dims[1] % iter_count);
+
+    UfoRequisition requisition = {
+        .n_dims = 3,
+        .dims[0] = image_requisition.dims[0],
+        .dims[1] = image_requisition.dims[1] / iter_count + mod,
+        .dims[2] = number_ones,
+    };
+    UfoBuffer *ufo_buffer = ufo_buffer_new(&requisition, priv->context);
+
+    /* loads surrounding number_ones pixels of each pixel in image.
+     * Result in buffer */
+    launch_kernel_2D (priv->k_load_elements_from_patern, src, pattern,
+                      ufo_buffer, cmd_queue, pattern_requisition.dims[0],
+                      number_ones, height, y_offset, 0);
+    /* Create image of threshold telling how likely a pixel is a center of a
+     * ring */
+    launch_kernel_1D (priv->k_bitonic_ordfilt, ufo_buffer, dst, cmd_queue,
+                      number_ones, width * y_offset, 0);
+
+    for (unsigned iter = 0; iter < iter_count; ++iter) {
+        /* start at mod offset, since first iteration manipulated
+         * iter * height + mod rows */
+        y_offset = mod + iter * (height - mod) / iter_count;
+        launch_kernel_2D (priv->k_load_elements_from_patern, src, pattern,
+                          ufo_buffer, cmd_queue, pattern_requisition.dims[0],
+                          number_ones, height, y_offset, mod);
+        launch_kernel_1D (priv->k_bitonic_ordfilt, ufo_buffer, dst, cmd_queue,
+                          number_ones, width * y_offset, mod);
+    }
+
+    g_object_unref(ufo_buffer);
+}
+
+static gboolean
+ufo_ordfilt_task_process (UfoTask *task,
+                          UfoBuffer **inputs,
+                          UfoBuffer *output,
+                          UfoRequisition *requisition)
+{
+    UfoOrdfiltTaskPrivate *priv;
+    UfoGpuNode *node;
+    cl_command_queue cmd_queue;
+
+    priv = UFO_ORDFILT_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);
+
+    priv = UFO_ORDFILT_TASK_GET_PRIVATE (task);
+    compute_ordfilt (priv, inputs[0], inputs[1], output, cmd_queue);
+
+    return TRUE;
+}
+
+static void
+ufo_ordfilt_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_ordfilt_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_ordfilt_task_setup;
+    iface->get_num_inputs = ufo_ordfilt_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_ordfilt_task_get_num_dimensions;
+    iface->get_mode = ufo_ordfilt_task_get_mode;
+    iface->get_requisition = ufo_ordfilt_task_get_requisition;
+    iface->process = ufo_ordfilt_task_process;
+}
+
+static void
+ufo_ordfilt_task_class_init (UfoOrdfiltTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->finalize = ufo_ordfilt_task_finalize;
+
+    g_type_class_add_private (gobject_class, sizeof(UfoOrdfiltTaskPrivate));
+}
+
+static void
+ufo_ordfilt_task_init(UfoOrdfiltTask *self)
+{
+    self->priv = UFO_ORDFILT_TASK_GET_PRIVATE(self);
+}
diff --git a/src/ufo-ordfilt-task.h b/src/ufo-ordfilt-task.h
new file mode 100644
index 0000000..3b3d61f
--- /dev/null
+++ b/src/ufo-ordfilt-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_ORDFILT_TASK_H
+#define __UFO_ORDFILT_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_ORDFILT_TASK             (ufo_ordfilt_task_get_type())
+#define UFO_ORDFILT_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_ORDFILT_TASK, UfoOrdfiltTask))
+#define UFO_IS_ORDFILT_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_ORDFILT_TASK))
+#define UFO_ORDFILT_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_ORDFILT_TASK, UfoOrdfiltTaskClass))
+#define UFO_IS_ORDFILT_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_ORDFILT_TASK))
+#define UFO_ORDFILT_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_ORDFILT_TASK, UfoOrdfiltTaskClass))
+
+typedef struct _UfoOrdfiltTask           UfoOrdfiltTask;
+typedef struct _UfoOrdfiltTaskClass      UfoOrdfiltTaskClass;
+typedef struct _UfoOrdfiltTaskPrivate    UfoOrdfiltTaskPrivate;
+
+/**
+ * UfoOrdfiltTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoOrdfiltTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoOrdfiltTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoOrdfiltTaskPrivate *priv;
+};
+
+/**
+ * UfoOrdfiltTaskClass:
+ *
+ * #UfoOrdfiltTask class
+ */
+struct _UfoOrdfiltTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_ordfilt_task_new       (void);
+GType     ufo_ordfilt_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-priv.c b/src/ufo-priv.c
new file mode 100644
index 0000000..b37ef26
--- /dev/null
+++ b/src/ufo-priv.c
@@ -0,0 +1,12 @@
+#include "ufo-priv.h"
+
+guint
+ceil_power_of_two (guint x)
+{
+    guint res = 1;
+
+    while (res < x)
+        res *= 2;
+
+    return res;
+}
diff --git a/src/ufo-priv.h b/src/ufo-priv.h
new file mode 100644
index 0000000..98f08d7
--- /dev/null
+++ b/src/ufo-priv.h
@@ -0,0 +1,17 @@
+#ifndef UFO_PRIV_H
+#define UFO_PRIV_H
+
+#include <glib.h>
+
+G_BEGIN_DECLS
+
+#define g_list_for(list, it) \
+        for (it = g_list_first (list); \
+             it != NULL; \
+             it = g_list_next (it))
+
+guint ceil_power_of_two (guint x);
+
+G_END_DECLS
+
+#endif
diff --git a/src/ufo-reduce-task.c b/src/ufo-reduce-task.c
new file mode 100644
index 0000000..7fb7384
--- /dev/null
+++ b/src/ufo-reduce-task.c
@@ -0,0 +1,124 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "ufo-reduce-task.h"
+
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoReduceTask, ufo_reduce_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_REDUCE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_REDUCE_TASK, UfoReduceTaskPrivate))
+
+
+UfoNode *
+ufo_reduce_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_REDUCE_TASK, NULL));
+}
+
+static void
+ufo_reduce_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_reduce_task_get_requisition (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[0], requisition);
+    requisition->dims[0] /= 2;
+    requisition->dims[1] /= 2;
+}
+
+static guint
+ufo_reduce_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_reduce_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_reduce_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static gboolean
+ufo_reduce_task_process (UfoTask *task,
+                         UfoBuffer **inputs,
+                         UfoBuffer *output,
+                         UfoRequisition *requisition)
+{
+    UfoRequisition input_req;
+    ufo_buffer_get_requisition(inputs[0], &input_req);
+    float *src = ufo_buffer_get_host_array(inputs[0], NULL);
+    float *out = ufo_buffer_get_host_array(output, NULL);
+
+    for (unsigned i = 0; i < requisition->dims[0]; ++i) {
+        for (unsigned j = 0; j < requisition->dims[1]; ++j) {
+            out[i + j * requisition->dims[0]] = src[2 * i + 2 * j * input_req.dims[0]];
+            out[i + j * requisition->dims[0]] += src[2 * i + 1 + 2 * j * input_req.dims[0]];
+            out[i + j * requisition->dims[0]] += src[2 * i + j * 2 * input_req.dims[0] + input_req.dims[0]];
+            out[i + j * requisition->dims[0]] += src[2 * i + 1 + j * 2 * input_req.dims[0] + input_req.dims[0]];
+        }
+    }
+
+    return TRUE;
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_reduce_task_setup;
+    iface->get_num_inputs = ufo_reduce_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_reduce_task_get_num_dimensions;
+    iface->get_mode = ufo_reduce_task_get_mode;
+    iface->get_requisition = ufo_reduce_task_get_requisition;
+    iface->process = ufo_reduce_task_process;
+}
+
+static void
+ufo_reduce_task_class_init (UfoReduceTaskClass *klass)
+{
+}
+
+static void
+ufo_reduce_task_init(UfoReduceTask *self)
+{
+}
diff --git a/src/ufo-reduce-task.h b/src/ufo-reduce-task.h
new file mode 100644
index 0000000..bb55f6b
--- /dev/null
+++ b/src/ufo-reduce-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_REDUCE_TASK_H
+#define __UFO_REDUCE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_REDUCE_TASK             (ufo_reduce_task_get_type())
+#define UFO_REDUCE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_REDUCE_TASK, UfoReduceTask))
+#define UFO_IS_REDUCE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_REDUCE_TASK))
+#define UFO_REDUCE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_REDUCE_TASK, UfoReduceTaskClass))
+#define UFO_IS_REDUCE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_REDUCE_TASK))
+#define UFO_REDUCE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_REDUCE_TASK, UfoReduceTaskClass))
+
+typedef struct _UfoReduceTask           UfoReduceTask;
+typedef struct _UfoReduceTaskClass      UfoReduceTaskClass;
+typedef struct _UfoReduceTaskPrivate    UfoReduceTaskPrivate;
+
+/**
+ * UfoReduceTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoReduceTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoReduceTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoReduceTaskPrivate *priv;
+};
+
+/**
+ * UfoReduceTaskClass:
+ *
+ * #UfoReduceTask class
+ */
+struct _UfoReduceTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_reduce_task_new       (void);
+GType     ufo_reduce_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-remove-circle-task.c b/src/ufo-remove-circle-task.c
new file mode 100644
index 0000000..c1b2e9d
--- /dev/null
+++ b/src/ufo-remove-circle-task.c
@@ -0,0 +1,330 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <math.h>
+#include "ufo-remove-circle-task.h"
+#include "ufo-ring-coordinates.h"
+
+
+struct _UfoRemoveCircleTaskPrivate {
+    float threshold;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoRemoveCircleTask, ufo_remove_circle_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_REMOVE_CIRCLE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_REMOVE_CIRCLE_TASK, UfoRemoveCircleTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_THRESHOLD,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_remove_circle_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_REMOVE_CIRCLE_TASK, NULL));
+}
+
+static void
+ufo_remove_circle_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_remove_circle_task_get_requisition (UfoTask *task,
+                                        UfoBuffer **inputs,
+                                        UfoRequisition *requisition)
+{
+    /* At most we will have same sized buffer as input since we are only
+     * removing circles  */
+    ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_remove_circle_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_remove_circle_task_get_num_dimensions (UfoTask *task,
+                                           guint input)
+{
+    return 1;
+}
+
+static UfoTaskMode
+ufo_remove_circle_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static unsigned
+is_intersection(UfoRemoveCircleTaskPrivate *priv, UfoRingCoordinate *left,
+                UfoRingCoordinate *center, UfoRingCoordinate *right)
+{
+    float threshold = priv->threshold;
+    float lr_dx = left->x - right->x;
+    float lr_dy = left->y - right->y;
+    float lr_dist = sqrtf(lr_dx * lr_dx + lr_dy * lr_dy);
+
+    /* We want left and right to intersect */
+    if (left->r + right->r <= lr_dist)
+        return 0;
+
+    UfoRingCoordinate lr_mid = {
+        .x = (left->x + right->x) / 2,
+        .y = (left->y + right->y) / 2,
+        .r = (left->r + right->r) / 2
+    };
+    float r_inner = (left->r + right->r - lr_dist) / 2;
+    float r_outer = (left->r + right->r + lr_dist) / 2;
+
+    float dr_in = fabsf(center->r - r_inner);
+    float dr_out = fabsf(center->r - r_outer);
+    /* Center radi must have a raddi approximately equal to the inner or outer
+     * ring radius */
+    if (dr_out >= threshold && dr_in >= threshold)
+        return 0;
+
+    /* Compute distance from center to the midpoint location of left and right */
+    float cm_dx = center->x - lr_mid.x;
+    float cm_dy = center->y - lr_mid.y;
+    float cm_dist = sqrtf(cm_dx * cm_dx + cm_dy * cm_dy);
+    /* Center has to be close enough to the optimal intersection center point */
+    if (cm_dist >= lr_dist / 2)
+        return 0;
+
+    return 1;
+}
+
+static void remove_circle(UfoRemoveCircleTaskPrivate *priv, URCS *src, URCS *dst)
+{
+    unsigned nb_elt = (unsigned) src->nb_elt;
+    /* Number of elements to keep */
+    unsigned res_counter = 0;
+    /* Check if any of the elements is actually the intersection of two others
+     * rings */
+    for (unsigned center = 0; center < nb_elt; ++center) {
+        /* Flag to check if center is an intersecting ring */
+        unsigned is_intersecting_ring = 0;
+        /* Check if there exists a  pair (not containing center) of rings such
+         * as center rings is indeed the intersection of that pair of rings */
+        for (unsigned left = 0; left < nb_elt; ++left) {
+            if (left == center)
+                continue;
+
+            for (unsigned right = left + 1; right < nb_elt; ++right) {
+                if (right == center)
+                    continue;
+
+                if (is_intersection(priv, &src->coord[left], &src->coord[center], &src->coord[right])) {
+                    is_intersecting_ring = 1;
+                    break;
+                }
+            }
+
+            /* Stop searching if we know that centers needs removing */
+            if (is_intersecting_ring)
+                break;
+        }
+
+        if (!is_intersecting_ring) {
+            dst->coord[res_counter] = src->coord[center];
+            ++res_counter;
+        }
+    }
+    dst->nb_elt = (float) res_counter;
+}
+
+static
+int get_min_contrast(URCS *src)
+{
+    int min_i = 0;
+    char found = 0;
+    for (int i = 0; i < src->nb_elt; ++i){
+        if (src->coord[i].contrast) {
+            found = 1;
+            if (src->coord[min_i].contrast * src->coord[min_i].intensity >
+                src->coord[i].contrast * src->coord[i].intensity ||
+                src->coord[min_i].contrast == 0)
+                min_i = i;
+        }
+    }
+    if (!found)
+        return -1;
+    return min_i;
+}
+
+static
+void remove_inner_circle(URCS *src, URCS *dst)
+{
+    int min_i = get_min_contrast(src);
+    unsigned counter = 0;
+
+    while (min_i != -1 && counter < src->nb_elt) {
+        for (int i = 0; i < src->nb_elt; ++i) {
+            if (!src->coord[i].contrast || i == min_i)
+                continue;
+
+            float dx = src->coord[i].x - src->coord[min_i].x;
+            float dy = src->coord[i].y - src->coord[min_i].y;
+            float dist = sqrtf(dx * dx + dy * dy);
+            /* Check if other ring is touching current ring */
+            float big_rad = src->coord[min_i].r;
+            float small_rad = src->coord[i].r;
+
+            if (small_rad > big_rad) {
+                small_rad = big_rad;
+                big_rad = src->coord[i].r;
+            }
+
+            if (fabs(dist + small_rad - big_rad) < 6) {
+                /* Check if other ring contrast is less to current ring */
+                /* If so remove it, as it likely is a false positive, i-e it was
+                 * using part of current ring */
+                if ((src->coord[i].contrast * src->coord[i].intensity) /
+                    (src->coord[min_i].contrast * src->coord[min_i].intensity) < 0.75)
+                    src->coord[i].contrast = 0;
+            }
+        }
+        dst->coord[counter] = src->coord[min_i];
+        ++counter;
+        /* Notify that we have check this ring to not process it again */
+        src->coord[min_i].contrast = 0;
+        min_i = get_min_contrast(src);
+    }
+    dst->nb_elt = (float) counter;
+    if (counter >= src->nb_elt && min_i != -1)
+        g_print("Remove Circle : Some memory corruption occured.  This is a bug "
+                " in the software that needs fixing\n");
+}
+
+static gboolean
+ufo_remove_circle_task_process (UfoTask *task,
+                                UfoBuffer **inputs,
+                                UfoBuffer *output,
+                                UfoRequisition *requisition)
+{
+    UfoBuffer* dup = ufo_buffer_dup(inputs[0]);
+    ufo_buffer_copy(inputs[0], dup);
+    URCS *coord = (URCS *) ufo_buffer_get_host_array(dup, NULL);
+    URCS *out_coord = (URCS *) ufo_buffer_get_host_array(output, NULL);
+    UfoRemoveCircleTaskPrivate *priv = UFO_REMOVE_CIRCLE_TASK_GET_PRIVATE(task);
+    remove_inner_circle(coord, out_coord);
+
+    ufo_buffer_copy(output, dup);
+    remove_circle(priv, coord, out_coord);
+
+    g_object_unref(dup);
+    return TRUE;
+}
+
+static void
+ufo_remove_circle_task_set_property (GObject *object,
+                                     guint property_id,
+                                     const GValue *value,
+                                     GParamSpec *pspec)
+{
+    UfoRemoveCircleTaskPrivate *priv = UFO_REMOVE_CIRCLE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_THRESHOLD:
+            priv->threshold = g_value_get_float (value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_remove_circle_task_get_property (GObject *object,
+                                     guint property_id,
+                                     GValue *value,
+                                     GParamSpec *pspec)
+{
+    UfoRemoveCircleTaskPrivate *priv = UFO_REMOVE_CIRCLE_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_THRESHOLD:
+            g_value_set_float (value, priv->threshold);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_remove_circle_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_remove_circle_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_remove_circle_task_setup;
+    iface->get_num_inputs = ufo_remove_circle_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_remove_circle_task_get_num_dimensions;
+    iface->get_mode = ufo_remove_circle_task_get_mode;
+    iface->get_requisition = ufo_remove_circle_task_get_requisition;
+    iface->process = ufo_remove_circle_task_process;
+}
+
+static void
+ufo_remove_circle_task_class_init (UfoRemoveCircleTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_remove_circle_task_set_property;
+    gobject_class->get_property = ufo_remove_circle_task_get_property;
+    gobject_class->finalize = ufo_remove_circle_task_finalize;
+
+    properties[PROP_THRESHOLD] =
+        g_param_spec_float ("threshold",
+                            "Set maximum Inner and outer ring radii size difference",
+                            "Set maximum inner and outer ring radii size difference",
+                            1, G_MAXFLOAT, 5,
+                            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoRemoveCircleTaskPrivate));
+}
+
+static void
+ufo_remove_circle_task_init(UfoRemoveCircleTask *self)
+{
+    self->priv = UFO_REMOVE_CIRCLE_TASK_GET_PRIVATE(self);
+    self->priv->threshold = 5;
+}
diff --git a/src/ufo-remove-circle-task.h b/src/ufo-remove-circle-task.h
new file mode 100644
index 0000000..c7e5dcf
--- /dev/null
+++ b/src/ufo-remove-circle-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_CIRCLE_TASK_H
+#define __UFO_REMOVE_CIRCLE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_REMOVE_CIRCLE_TASK             (ufo_remove_circle_task_get_type())
+#define UFO_REMOVE_CIRCLE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_REMOVE_CIRCLE_TASK, UfoRemoveCircleTask))
+#define UFO_IS_REMOVE_CIRCLE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_REMOVE_CIRCLE_TASK))
+#define UFO_REMOVE_CIRCLE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_REMOVE_CIRCLE_TASK, UfoRemoveCircleTaskClass))
+#define UFO_IS_REMOVE_CIRCLE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_REMOVE_CIRCLE_TASK))
+#define UFO_REMOVE_CIRCLE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_REMOVE_CIRCLE_TASK, UfoRemoveCircleTaskClass))
+
+typedef struct _UfoRemoveCircleTask           UfoRemoveCircleTask;
+typedef struct _UfoRemoveCircleTaskClass      UfoRemoveCircleTaskClass;
+typedef struct _UfoRemoveCircleTaskPrivate    UfoRemoveCircleTaskPrivate;
+
+/**
+ * UfoRemoveCircleTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoRemoveCircleTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoRemoveCircleTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoRemoveCircleTaskPrivate *priv;
+};
+
+/**
+ * UfoRemoveCircleTaskClass:
+ *
+ * #UfoRemoveCircleTask class
+ */
+struct _UfoRemoveCircleTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_remove_circle_task_new       (void);
+GType     ufo_remove_circle_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-replicate-task.c b/src/ufo-replicate-task.c
new file mode 100644
index 0000000..98cb5cc
--- /dev/null
+++ b/src/ufo-replicate-task.c
@@ -0,0 +1,134 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <stdlib.h>
+#include "ufo-replicate-task.h"
+
+struct _UfoReplicateTaskPrivate {
+    unsigned alloc_size;
+    unsigned idx;
+    UfoBuffer *data;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoReplicateTask, ufo_replicate_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_REPLICATE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_REPLICATE_TASK, UfoReplicateTaskPrivate))
+
+UfoNode *
+ufo_replicate_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_REPLICATE_TASK, NULL));
+}
+
+static void
+ufo_replicate_task_setup (UfoTask *task,
+                          UfoResources *resources,
+                          GError **error)
+{
+}
+
+static void
+ufo_replicate_task_get_requisition (UfoTask *task,
+                                    UfoBuffer **inputs,
+                                    UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_replicate_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_replicate_task_get_num_dimensions (UfoTask *task,
+                                       guint input)
+{
+    return 2;
+}
+
+static UfoTaskMode
+ufo_replicate_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_REDUCTOR;
+}
+
+static gboolean
+ufo_replicate_task_process (UfoTask *task,
+                            UfoBuffer **inputs,
+                            UfoBuffer *output,
+                            UfoRequisition *requisition)
+{
+    UfoReplicateTaskPrivate *priv = UFO_REPLICATE_TASK_GET_PRIVATE (task);
+
+    if (priv->idx > priv->alloc_size) {
+        priv->alloc_size *= 2;
+        size_t nb_bytes = priv->alloc_size * sizeof (UfoBuffer *);
+        priv->data = realloc (priv->data, nb_bytes);
+    }
+
+    return TRUE;
+}
+
+static void
+ufo_replicate_task_finalize (GObject *object)
+{
+    UfoReplicateTaskPrivate *priv;
+
+    priv = UFO_REPLICATE_TASK_GET_PRIVATE (object);
+    g_free (priv->data);
+    G_OBJECT_CLASS (ufo_replicate_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_replicate_task_setup;
+    iface->get_num_inputs = ufo_replicate_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_replicate_task_get_num_dimensions;
+    iface->get_mode = ufo_replicate_task_get_mode;
+    iface->get_requisition = ufo_replicate_task_get_requisition;
+    iface->process = ufo_replicate_task_process;
+}
+
+static void
+ufo_replicate_task_class_init (UfoReplicateTaskClass *klass)
+{
+    GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+    oclass->finalize = ufo_replicate_task_finalize;
+
+    g_type_class_add_private (oclass, sizeof(UfoReplicateTaskPrivate));
+}
+
+static void
+ufo_replicate_task_init(UfoReplicateTask *self)
+{
+    self->priv = UFO_REPLICATE_TASK_GET_PRIVATE(self);
+    self->priv->alloc_size = 256;
+    self->priv->data = g_malloc (sizeof (UfoBuffer *) * self->priv->alloc_size);
+    self->priv->idx = 0;
+}
diff --git a/src/ufo-replicate-task.h b/src/ufo-replicate-task.h
new file mode 100644
index 0000000..af39d88
--- /dev/null
+++ b/src/ufo-replicate-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_REPLICATE_TASK_H
+#define __UFO_REPLICATE_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_REPLICATE_TASK             (ufo_replicate_task_get_type())
+#define UFO_REPLICATE_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_REPLICATE_TASK, UfoReplicateTask))
+#define UFO_IS_REPLICATE_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_REPLICATE_TASK))
+#define UFO_REPLICATE_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_REPLICATE_TASK, UfoReplicateTaskClass))
+#define UFO_IS_REPLICATE_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_REPLICATE_TASK))
+#define UFO_REPLICATE_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_REPLICATE_TASK, UfoReplicateTaskClass))
+
+typedef struct _UfoReplicateTask           UfoReplicateTask;
+typedef struct _UfoReplicateTaskClass      UfoReplicateTaskClass;
+typedef struct _UfoReplicateTaskPrivate    UfoReplicateTaskPrivate;
+
+/**
+ * UfoReplicateTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoReplicateTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoReplicateTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoReplicateTaskPrivate *priv;
+};
+
+/**
+ * UfoReplicateTaskClass:
+ *
+ * #UfoReplicateTask class
+ */
+struct _UfoReplicateTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_replicate_task_new       (void);
+GType     ufo_replicate_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-ring-coordinates.h b/src/ufo-ring-coordinates.h
new file mode 100644
index 0000000..e52c800
--- /dev/null
+++ b/src/ufo-ring-coordinates.h
@@ -0,0 +1,20 @@
+#ifndef UFO_RING_COORDINATES_H
+#define UFO_RING_COORDINATES_H
+
+struct _UfoRingCoordinate {
+    float x;
+    float y;
+    float r;
+    float contrast;
+    float intensity;
+};
+
+struct _UfoRingCoordinatesStream {
+    float nb_elt;
+    struct _UfoRingCoordinate *coord;
+};
+
+typedef struct _UfoRingCoordinate UfoRingCoordinate;
+typedef struct _UfoRingCoordinatesStream URCS;
+
+#endif
diff --git a/src/ufo-ring-pattern-task.c b/src/ufo-ring-pattern-task.c
new file mode 100644
index 0000000..510764e
--- /dev/null
+++ b/src/ufo-ring-pattern-task.c
@@ -0,0 +1,314 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <math.h>
+#include <stdlib.h>
+
+#include "ufo-ring-pattern-task.h"
+
+
+struct _UfoRingPatternTaskPrivate {
+    unsigned ring_thickness;
+    unsigned ring_end;
+    unsigned ring_start;
+    unsigned ring_current;
+    unsigned ring_step;
+    unsigned width;
+    unsigned height;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoRingPatternTask, ufo_ring_pattern_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_RING_PATTERN_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_RING_PATTERN_TASK, UfoRingPatternTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_RING_START,
+    PROP_RING_STEP,
+    PROP_RING_END,
+    PROP_RING_THICKNESS,
+    PROP_WIDTH,
+    PROP_HEIGHT,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_ring_pattern_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_RING_PATTERN_TASK, NULL));
+}
+
+static void
+ufo_ring_pattern_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_ring_pattern_task_get_requisition (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoRequisition *requisition)
+{
+    UfoRingPatternTaskPrivate *priv = UFO_RING_PATTERN_TASK_GET_PRIVATE (task);
+    requisition->dims[0] = priv->width;
+    requisition->dims[1] = priv->height;
+    requisition->n_dims = 2;
+}
+
+static guint
+ufo_ring_pattern_task_get_num_inputs (UfoTask *task)
+{
+    return 0;
+}
+
+static guint
+ufo_ring_pattern_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    return 0;
+}
+
+static UfoTaskMode
+ufo_ring_pattern_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_GENERATOR | UFO_TASK_MODE_CPU;
+}
+
+static void add_ring_metadata(UfoBuffer *output, unsigned number_ones,
+                              unsigned radius)
+{
+    GValue value = G_VALUE_INIT;
+    g_value_init(&value, G_TYPE_UINT);
+
+    g_value_set_uint(&value, number_ones);
+    ufo_buffer_set_metadata (output, "number_ones", &value);
+
+    g_value_set_uint(&value, radius);
+    ufo_buffer_set_metadata (output, "radius", &value);
+}
+
+static gboolean
+ufo_ring_pattern_task_generate (UfoTask *task,
+                                UfoBuffer *output,
+                                UfoRequisition *requisition)
+{
+    UfoRingPatternTaskPrivate *priv = UFO_RING_PATTERN_TASK_GET_PRIVATE (task);
+
+    if (priv->ring_current > priv->ring_end)
+        return FALSE;
+
+    float *out = ufo_buffer_get_host_array (output, NULL);
+    int dimx = (int) priv->width;
+    int dimy = (int) priv->height;
+    unsigned counter = 0;
+
+    for (int y = -(dimy / 2); y < dimy / 2; ++y) {
+        for (int x = -(dimx / 2); x < dimx / 2; ++x) {
+            double dist = sqrt (x * x + y * y) - priv->ring_current;
+            dist = dist < 0 ? -dist : dist;
+            if (dist < (double) priv->ring_thickness / 2) {
+                out[(x + (dimx)) % dimx + ((y + (dimy)) % dimy) * dimx] = 1;
+                ++counter;
+            }
+            else
+                out[(x + (dimx)) % dimx + ((y + (dimy)) % dimy) * dimx] = 0;
+        }
+    }
+
+    add_ring_metadata(output, counter, priv->ring_current);
+
+    for (int y = -(dimy / 2); y < dimy / 2; ++y) {
+        for (int x = -(dimx / 2); x < dimx / 2; ++x) {
+            double dist = sqrt (x * x + y * y) - priv->ring_current;
+            dist = dist < 0 ? -dist : dist;
+            if (dist < (double) priv->ring_thickness / 2) {
+                out[(x + (dimx)) % dimx + ((y + (dimy)) % dimy) * dimx] /=
+                    (float) counter;
+            }
+            else
+                out[(x + (dimx)) % dimx + ((y + (dimy)) % dimy) * dimx] = 0;
+        }
+    }
+
+    priv->ring_current += priv->ring_step;
+    return TRUE;
+}
+
+static void
+ufo_ring_pattern_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoRingPatternTaskPrivate *priv = UFO_RING_PATTERN_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_RING_START:
+            priv->ring_start = g_value_get_uint(value);
+            priv->ring_current = g_value_get_uint(value);
+            break;
+        case PROP_RING_STEP:
+            priv->ring_step = g_value_get_uint(value);
+            break;
+        case PROP_RING_END:
+            priv->ring_end = g_value_get_uint(value);
+            break;
+        case PROP_RING_THICKNESS:
+            priv->ring_thickness = g_value_get_uint(value);
+            break;
+        case PROP_WIDTH:
+            priv->width = ceil_power_of_two (g_value_get_uint (value));
+            break;
+        case PROP_HEIGHT:
+            priv->height = ceil_power_of_two (g_value_get_uint (value));
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_ring_pattern_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoRingPatternTaskPrivate *priv = UFO_RING_PATTERN_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_RING_START:
+            g_value_set_uint (value, priv->ring_start);
+            break;
+        case PROP_RING_STEP:
+            g_value_set_uint (value, priv->ring_step);
+            break;
+        case PROP_RING_END:
+            g_value_set_uint (value, priv->ring_end);
+            break;
+        case PROP_RING_THICKNESS:
+            g_value_set_uint (value, priv->ring_thickness);
+            break;
+        case PROP_WIDTH:
+            g_value_set_uint (value, priv->width);
+            break;
+        case PROP_HEIGHT:
+            g_value_set_uint (value, priv->height);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_ring_pattern_task_finalize (GObject *object)
+{
+    G_OBJECT_CLASS (ufo_ring_pattern_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_ring_pattern_task_setup;
+    iface->get_num_inputs = ufo_ring_pattern_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_ring_pattern_task_get_num_dimensions;
+    iface->get_mode = ufo_ring_pattern_task_get_mode;
+    iface->get_requisition = ufo_ring_pattern_task_get_requisition;
+    iface->generate = ufo_ring_pattern_task_generate;
+}
+
+static void
+ufo_ring_pattern_task_class_init (UfoRingPatternTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_ring_pattern_task_set_property;
+    gobject_class->get_property = ufo_ring_pattern_task_get_property;
+    gobject_class->finalize = ufo_ring_pattern_task_finalize;
+
+    properties[PROP_RING_START] =
+        g_param_spec_uint ("ring_start",
+                           "give starting radius size",
+                           "give starting radius size",
+                           1, G_MAXUINT, 5,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_RING_STEP] =
+        g_param_spec_uint ("ring_step",
+                           "Gives ring step",
+                           "Gives ring step",
+                           1, G_MAXUINT, 2,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_RING_END] =
+        g_param_spec_uint ("ring_end",
+                           "give ending radius size",
+                           "give ending radius size",
+                           1, G_MAXUINT, 5,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_RING_THICKNESS] =
+        g_param_spec_uint ("ring_thickness",
+                           "give desired ring thickness",
+                           "give desired ring thickness",
+                           1, G_MAXUINT, 13,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_WIDTH] =
+        g_param_spec_uint ("width",
+                           "Give x size of output image",
+                           "Give x size of output image",
+                           1, G_MAXUINT, 1024,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_HEIGHT] =
+        g_param_spec_uint ("height",
+                           "Give y size of output image",
+                           "Give y size of output image",
+                           1, G_MAXUINT, 1024,
+                           G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (gobject_class, i, properties[i]);
+
+    g_type_class_add_private (gobject_class, sizeof(UfoRingPatternTaskPrivate));
+}
+
+static void
+ufo_ring_pattern_task_init(UfoRingPatternTask *self)
+{
+    self->priv = UFO_RING_PATTERN_TASK_GET_PRIVATE(self);
+    self->priv->ring_start = 5;
+    self->priv->ring_end = 5;
+    self->priv->ring_thickness = 13;
+    self->priv->ring_current = self->priv->ring_start;
+    self->priv->width = 1024;
+    self->priv->height = 1024;
+    self->priv->ring_step = 2;
+}
diff --git a/src/ufo-ring-pattern-task.h b/src/ufo-ring-pattern-task.h
new file mode 100644
index 0000000..131fbb6
--- /dev/null
+++ b/src/ufo-ring-pattern-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_RING_PATTERN_TASK_H
+#define __UFO_RING_PATTERN_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_RING_PATTERN_TASK             (ufo_ring_pattern_task_get_type())
+#define UFO_RING_PATTERN_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_RING_PATTERN_TASK, UfoRingPatternTask))
+#define UFO_IS_RING_PATTERN_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_RING_PATTERN_TASK))
+#define UFO_RING_PATTERN_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_RING_PATTERN_TASK, UfoRingPatternTaskClass))
+#define UFO_IS_RING_PATTERN_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_RING_PATTERN_TASK))
+#define UFO_RING_PATTERN_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_RING_PATTERN_TASK, UfoRingPatternTaskClass))
+
+typedef struct _UfoRingPatternTask           UfoRingPatternTask;
+typedef struct _UfoRingPatternTaskClass      UfoRingPatternTaskClass;
+typedef struct _UfoRingPatternTaskPrivate    UfoRingPatternTaskPrivate;
+
+/**
+ * UfoRingPatternTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoRingPatternTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoRingPatternTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoRingPatternTaskPrivate *priv;
+};
+
+/**
+ * UfoRingPatternTaskClass:
+ *
+ * #UfoRingPatternTask class
+ */
+struct _UfoRingPatternTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_ring_pattern_task_new       (void);
+GType     ufo_ring_pattern_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file
diff --git a/src/ufo-ringwriter-task.c b/src/ufo-ringwriter-task.c
new file mode 100644
index 0000000..dcbfeb1
--- /dev/null
+++ b/src/ufo-ringwriter-task.c
@@ -0,0 +1,278 @@
+/*
+ * Copyright (C) 2011-2015 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/>.
+ *
+ * Authored by: Alexandre Lewkowicz (lewkow_a@epita.fr)
+ */
+
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+#include "ufo-ringwriter-task.h"
+#include "ufo-ring-coordinates.h"
+
+struct _UfoRingwriterTaskPrivate {
+    char *filename;
+    unsigned scale;
+    FILE *file;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoRingwriterTask, ufo_ringwriter_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_RINGWRITER_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_RINGWRITER_TASK, UfoRingwriterTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_PATH,
+    PROP_SCALE,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_ringwriter_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_RINGWRITER_TASK, NULL));
+}
+
+static void
+ufo_ringwriter_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+}
+
+static void
+ufo_ringwriter_task_get_requisition (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoRequisition *requisition)
+{
+    requisition->n_dims = 0;
+}
+
+static guint
+ufo_ringwriter_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_ringwriter_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    return 1;
+}
+
+static UfoTaskMode
+ufo_ringwriter_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_CPU;
+}
+
+static void
+get_file_metadata (UfoBuffer *src, const char **piv_file_name, unsigned *piv_file_idx)
+{
+    GValue *value;
+
+    value = ufo_buffer_get_metadata(src, "piv_file_name");
+    *piv_file_name = g_value_get_string(value);
+
+    value = ufo_buffer_get_metadata(src, "piv_file_idx");
+    *piv_file_idx = g_value_get_uint(value);
+}
+
+static void
+ufo_ringwriter_write_metadata (UfoBuffer *src, FILE *file)
+{
+    const char *piv_file_name;
+    unsigned piv_file_idx;
+    char str[256];
+    int count;
+
+    get_file_metadata (src, &piv_file_name, &piv_file_idx);
+    count = sprintf(str, "filename %s\n", piv_file_name);
+
+    if (count < 0)
+        g_print ("Unable to convert data to output text file\n");
+
+    fwrite (str, sizeof (char), (size_t) count, file);
+    count = sprintf(str, "index %u\n", piv_file_idx);
+
+    if (count < 0)
+        g_print ("Unable to convert data to output text file\n");
+
+    fwrite (str, sizeof (char), (size_t) count, file);
+}
+
+static gboolean
+ufo_ringwriter_task_process (UfoTask *task,
+                             UfoBuffer **inputs,
+                             UfoBuffer *output,
+                             UfoRequisition *requisition)
+{
+    UfoRingwriterTaskPrivate *priv = UFO_RINGWRITER_TASK_GET_PRIVATE (task);
+
+    if (priv->file == NULL) {
+        static int file_count = 0;
+        char filename[256];
+        sprintf(filename, "%s%i.txt", priv->filename, file_count);
+        priv->file = fopen (filename, "w");
+        ++file_count;
+    }
+
+    URCS *ring_stream = (URCS *) ufo_buffer_get_host_array (inputs[0], NULL);
+
+    FILE *file = priv->file;
+    ufo_ringwriter_write_metadata(inputs[0], file);
+    char str[256];
+    int count = sprintf(str, "ring_count %u\n", (unsigned) ring_stream->nb_elt);
+
+    if (count < 0)
+        g_print ("Unable to convert data to output text file\n");
+
+    fwrite (str, sizeof (char), (size_t) count, file);
+
+    for (unsigned i = 0; i < ring_stream->nb_elt; ++i) {
+        int x = (int) roundf (ring_stream->coord[i].x * (float) priv->scale);
+        int y = (int) roundf (ring_stream->coord[i].y * (float) priv->scale);
+        float r =  ring_stream->coord[i].r * (float) priv->scale;
+        count = sprintf(str, "ring_coord %i %i %f\n", x, y, r);
+        if (count < 0)
+            g_print ("Unable to convert data to output text file\n");
+        size_t write_count = 0;
+        char *tmp = str;
+
+        while (write_count != (size_t) count) {
+            size_t c = fwrite (tmp, sizeof (char), (size_t) count, file);
+            if (c == 0)
+                g_print ("Unable to write data to output text file\n");
+            tmp += c;
+            write_count += c;
+        }
+    }
+    return TRUE;
+}
+
+static void
+ufo_ringwriter_task_set_property (GObject *object,
+                                  guint property_id,
+                                  const GValue *value,
+                                  GParamSpec *pspec)
+{
+    UfoRingwriterTaskPrivate *priv = UFO_RINGWRITER_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            priv->scale = g_value_get_uint(value);
+            break;
+        case PROP_PATH:
+            priv->filename = g_strdup (g_value_get_string (value));
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_ringwriter_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoRingwriterTaskPrivate *priv = UFO_RINGWRITER_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            g_value_set_uint (value, priv->scale);
+            break;
+        case PROP_PATH:
+            g_value_set_string (value, priv->filename);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_ringwriter_task_finalize (GObject *object)
+{
+    UfoRingwriterTaskPrivate *priv = UFO_RINGWRITER_TASK_GET_PRIVATE (object);
+
+    if (priv->file)
+        fclose (priv->file);
+
+    if (priv->filename)
+        g_free (priv->filename);
+
+    G_OBJECT_CLASS (ufo_ringwriter_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_ringwriter_task_setup;
+    iface->get_num_inputs = ufo_ringwriter_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_ringwriter_task_get_num_dimensions;
+    iface->get_mode = ufo_ringwriter_task_get_mode;
+    iface->get_requisition = ufo_ringwriter_task_get_requisition;
+    iface->process = ufo_ringwriter_task_process;
+}
+
+static void
+ufo_ringwriter_task_class_init (UfoRingwriterTaskClass *klass)
+{
+    GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+    oclass->set_property = ufo_ringwriter_task_set_property;
+    oclass->get_property = ufo_ringwriter_task_get_property;
+    oclass->finalize = ufo_ringwriter_task_finalize;
+
+    properties[PROP_SCALE] =
+        g_param_spec_uint ("scale",
+                           "Says by how much rings should be increased",
+                           "Says by how much rings should be increased",
+                           1, G_MAXUINT, 1,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_PATH] =
+        g_param_spec_string ("filename",
+                             "Path for the output file to write data",
+                             "Path for the output file to write data",
+                             "results",
+                             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(UfoRingwriterTaskPrivate));
+}
+
+static void
+ufo_ringwriter_task_init(UfoRingwriterTask *self)
+{
+    self->priv = UFO_RINGWRITER_TASK_GET_PRIVATE(self);
+    self->priv->filename = g_strdup ("results");
+    self->priv->file = NULL;
+    self->priv->scale = 1;
+}
diff --git a/src/ufo-ringwriter-task.h b/src/ufo-ringwriter-task.h
new file mode 100644
index 0000000..bacf3ef
--- /dev/null
+++ b/src/ufo-ringwriter-task.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 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_RINGWRITER_TASK_H
+#define __UFO_RINGWRITER_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_RINGWRITER_TASK             (ufo_ringwriter_task_get_type())
+#define UFO_RINGWRITER_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_RINGWRITER_TASK, UfoRingwriterTask))
+#define UFO_IS_RINGWRITER_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_RINGWRITER_TASK))
+#define UFO_RINGWRITER_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_RINGWRITER_TASK, UfoRingwriterTaskClass))
+#define UFO_IS_RINGWRITER_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_RINGWRITER_TASK))
+#define UFO_RINGWRITER_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_RINGWRITER_TASK, UfoRingwriterTaskClass))
+
+typedef struct _UfoRingwriterTask           UfoRingwriterTask;
+typedef struct _UfoRingwriterTaskClass      UfoRingwriterTaskClass;
+typedef struct _UfoRingwriterTaskPrivate    UfoRingwriterTaskPrivate;
+
+/**
+ * UfoRingwriterTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoRingwriterTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoRingwriterTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoRingwriterTaskPrivate *priv;
+};
+
+/**
+ * UfoRingwriterTaskClass:
+ *
+ * #UfoRingwriterTask class
+ */
+struct _UfoRingwriterTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_ringwriter_task_new       (void);
+GType     ufo_ringwriter_task_get_type  (void);
+
+G_END_DECLS
+
+#endif
\ No newline at end of file