Add image-based padding

src/CMakeLists.txt

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

src/kernels/ankapadding.cl

@@ -0,0 +1,31 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+__kernel void pad (image2d_t in_image,
+                   sampler_t sampler,
+                   global float *result,
+                   const uint2 input_shape,
+                   const int2 offset)
+{
+    int2 pixel = (int2) (get_global_id (0), get_global_id (1));
+    float2 norm_pixel = (float2) (((float) pixel.x - offset.x) / input_shape.x,
+                                  ((float) pixel.y - offset.y) / input_shape.y);
+
+    result[pixel.y * get_global_size(0) + pixel.x] = read_imagef(in_image, sampler, norm_pixel).x;
+}

src/ufo-anka-padding-task.c

@@ -0,0 +1,391 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "ufo-anka-padding-task.h"
+
+/**
+ * SECTION:ufo-anka-padding-task
+ * @Short_description: Pad images to some extent
+ * @Title: anka_padding
+ *
+ */
+
+struct _UfoAnkaPaddingTaskPrivate {
+    /* OpenCL */
+    cl_context context;
+    cl_kernel kernel;
+    cl_sampler sampler;
+
+    /* properties */
+    guint width, height, x_0, y_0;
+    cl_addressing_mode addressing_mode;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoAnkaPaddingTask, ufo_anka_padding_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_ANKA_PADDING_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_ANKA_PADDING_TASK, UfoAnkaPaddingTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_WIDTH,
+    PROP_HEIGHT,
+    PROP_X_0,
+    PROP_Y_0,
+    PROP_ADDRESSING_MODE,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+static void
+change_sampler (UfoAnkaPaddingTaskPrivate *priv)
+{
+    cl_int err;
+
+    if (priv->sampler) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseSampler (priv->sampler));
+    }
+
+    priv->sampler = clCreateSampler (priv->context,
+                                     (cl_bool) TRUE,
+                                     priv->addressing_mode,
+                                     CL_FILTER_NEAREST,
+                                     &err);
+
+    UFO_RESOURCES_CHECK_CLERR (err);
+}
+
+UfoNode *
+ufo_anka_padding_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_ANKA_PADDING_TASK, NULL));
+}
+
+static void
+ufo_anka_padding_task_setup (UfoTask *task,
+                       UfoResources *resources,
+                       GError **error)
+{
+    UfoAnkaPaddingTaskPrivate *priv;
+
+    priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (task);
+    priv->context = ufo_resources_get_context (resources);
+    priv->kernel = ufo_resources_get_kernel (resources, "ankapadding.cl", "pad", error);
+    change_sampler (priv);
+
+    UFO_RESOURCES_CHECK_CLERR (clRetainContext (priv->context));
+    if (priv->kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+    }
+}
+
+static void
+ufo_anka_padding_task_get_requisition (UfoTask *task,
+                                 UfoBuffer **inputs,
+                                 UfoRequisition *requisition)
+{
+    UfoAnkaPaddingTaskPrivate *priv;
+    UfoRequisition in_req;
+
+    priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (task);
+    ufo_buffer_get_requisition (inputs[0], &in_req);
+
+    /* if width and height are not set make them as large as input */
+    if (!priv->width) {
+        priv->width = (guint) in_req.dims[0];
+    }
+    if (!priv->height) {
+        priv->height = (guint) in_req.dims[1];
+    }
+
+    requisition->n_dims = 2;
+    requisition->dims[0] = priv->width;
+    requisition->dims[1] = priv->height;
+}
+
+static guint
+ufo_anka_padding_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_anka_padding_task_get_num_dimensions (UfoTask *task,
+                                             guint input)
+{
+    g_return_val_if_fail (input == 0, 0);
+
+    return 2;
+}
+
+static UfoTaskMode
+ufo_anka_padding_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+static gboolean
+ufo_anka_padding_task_process (UfoTask *task,
+                         UfoBuffer **inputs,
+                         UfoBuffer *output,
+                         UfoRequisition *requisition)
+{
+    UfoAnkaPaddingTaskPrivate *priv;
+    UfoGpuNode *node;
+    UfoProfiler *profiler;
+    UfoRequisition in_req;
+    cl_command_queue cmd_queue;
+    cl_mem in_image, out_mem;
+    cl_addressing_mode current_addressing_mode;
+    guint input_shape[2];
+    gint offset[2];
+
+    priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (task);
+    ufo_buffer_get_requisition (inputs[0], &in_req);
+    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+    out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+
+    /* change sampler mode if necessary */
+    UFO_RESOURCES_CHECK_CLERR (clGetSamplerInfo (priv->sampler,
+                                                 CL_SAMPLER_ADDRESSING_MODE,
+                                                 sizeof (cl_addressing_mode),
+                                                 &current_addressing_mode,
+                                                 NULL));
+
+    if (priv->addressing_mode != current_addressing_mode) {
+        change_sampler (priv);
+    }
+
+    input_shape[0] = in_req.dims[0];
+    input_shape[1] = in_req.dims[1];
+
+    offset[0] = priv->x_0;
+    offset[1] = priv->y_0;
+
+    in_image = ufo_buffer_get_device_image (inputs[0], cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_image));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_sampler), &priv->sampler));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_mem), &out_mem));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 3, sizeof (cl_uint2), input_shape));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 4, sizeof (cl_int2), offset));
+
+    profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+    ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL);
+
+    return TRUE;
+}
+
+static void
+ufo_anka_padding_task_set_property (GObject *object,
+                              guint property_id,
+                              const GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoAnkaPaddingTaskPrivate *priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_WIDTH:
+            priv->width = g_value_get_uint (value);
+            break;
+        case PROP_HEIGHT:
+            priv->height = g_value_get_uint (value);
+            break;
+        case PROP_X_0:
+            priv->x_0 = g_value_get_uint (value);
+            break;
+        case PROP_Y_0:
+            priv->y_0 = g_value_get_uint (value);
+            break;
+        case PROP_ADDRESSING_MODE:
+            if (!g_strcmp0 (g_value_get_string (value), "none")) {
+                priv->addressing_mode = CL_ADDRESS_NONE;
+            }
+            else if (!g_strcmp0 (g_value_get_string (value), "clamp")) {
+                priv->addressing_mode = CL_ADDRESS_CLAMP;
+            }
+            else if (!g_strcmp0 (g_value_get_string (value), "clamp_to_edge")) {
+                priv->addressing_mode = CL_ADDRESS_CLAMP_TO_EDGE;
+            }
+            else if (!g_strcmp0 (g_value_get_string (value), "repeat")) {
+                priv->addressing_mode = CL_ADDRESS_REPEAT;
+            } else {
+                g_log ("Ufo", G_LOG_LEVEL_CRITICAL,
+                       "Error <%s:%i>: Invalid addressing mode \"%s\", "\
+                       "it has to be one of [\"none\", \"clamp\", \"clamp_to_edge\", \"repeat\"]",
+                       __FILE__,
+                       __LINE__,
+                       g_value_get_string (value));
+            }
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_anka_padding_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    UfoAnkaPaddingTaskPrivate *priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_WIDTH:
+            g_value_set_uint (value, priv->width);
+            break;
+        case PROP_HEIGHT:
+            g_value_set_uint (value, priv->height);
+            break;
+        case PROP_X_0:
+            g_value_set_uint (value, priv->x_0);
+            break;
+        case PROP_Y_0:
+            g_value_set_uint (value, priv->y_0);
+            break;
+        case PROP_ADDRESSING_MODE:
+            switch (priv->addressing_mode) {
+                case CL_ADDRESS_NONE:
+                    g_value_set_string (value, "none");
+                    break;
+                case CL_ADDRESS_CLAMP:
+                    g_value_set_string (value, "clamp");
+                    break;
+                case CL_ADDRESS_CLAMP_TO_EDGE:
+                    g_value_set_string (value, "clamp_to_edge");
+                    break;
+                case CL_ADDRESS_REPEAT:
+                    g_value_set_string (value, "repeat");
+                    break;
+            }
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_anka_padding_task_finalize (GObject *object)
+{
+    UfoAnkaPaddingTaskPrivate *priv;
+
+    priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE (object);
+
+    if (priv->kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel));
+        priv->kernel = NULL;
+    }
+    if (priv->context) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseContext (priv->context));
+        priv->context = NULL;
+    }
+    if (priv->sampler) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseSampler (priv->sampler));
+        priv->sampler = NULL;
+    }
+
+    G_OBJECT_CLASS (ufo_anka_padding_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_anka_padding_task_setup;
+    iface->get_num_inputs = ufo_anka_padding_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_anka_padding_task_get_num_dimensions;
+    iface->get_mode = ufo_anka_padding_task_get_mode;
+    iface->get_requisition = ufo_anka_padding_task_get_requisition;
+    iface->process = ufo_anka_padding_task_process;
+}
+
+static void
+ufo_anka_padding_task_class_init (UfoAnkaPaddingTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_anka_padding_task_set_property;
+    gobject_class->get_property = ufo_anka_padding_task_get_property;
+    gobject_class->finalize = ufo_anka_padding_task_finalize;
+
+    properties[PROP_WIDTH] =
+        g_param_spec_uint ("width",
+                           "Padded width",
+                           "Padded width",
+                           0, +8192, 0,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_HEIGHT] =
+        g_param_spec_uint ("height",
+                           "Padded height",
+                           "Padded height",
+                           0, +8192, 0,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_X_0] =
+        g_param_spec_uint ("x_0",
+                           "X start index",
+                           "X index for input's 0th column",
+                           0, +8192, 0,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_Y_0] =
+        g_param_spec_uint ("y_0",
+                           "Y start index",
+                           "Y index for input's 0th row",
+                           0, +8192, 0,
+                           G_PARAM_READWRITE);
+
+    properties[PROP_ADDRESSING_MODE] =
+        g_param_spec_string ("addressing_mode",
+                             "Outlier treatment",
+                             "Outlier treatment from : \"none\", \"clamp\", \"clamp_to_edge\", \"repeat\"",
+                             "clamp",
+                             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(UfoAnkaPaddingTaskPrivate));
+}
+
+static void
+ufo_anka_padding_task_init(UfoAnkaPaddingTask *self)
+{
+    self->priv = UFO_ANKA_PADDING_TASK_GET_PRIVATE(self);
+
+    self->priv->width = 0;
+    self->priv->height = 0;
+    self->priv->x_0 = 0;
+    self->priv->y_0 = 0;
+    self->priv->addressing_mode = CL_ADDRESS_CLAMP;
+}

src/ufo-anka-padding-task.h

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