code cleen up

src/padding_2d.cl

@@ -0,0 +1,81 @@
+__kernel void padding_2d_init_const(
+                __global float * out_img,
+                   const unsigned int nxs,
+		   const float pval)
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int lidx = idy * nxs + idx;
+    out_img[lidx] = pval;
+}
+
+
+__kernel void padding_2d_copy_in(
+		__global float * inp_img,
+		__global float * out_img,
+                const unsigned int sxs,
+                const unsigned int lxs,
+	       const unsigned int xl,
+               const unsigned int yt)
+{
+
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int lidx = (idy + yt) * lxs + (idx + xl);
+    const int sidx = idy * sxs + idx; 
+
+    out_img[lidx] = inp_img[sidx]; 
+}
+
+__kernel void padding_2d_brep(
+               __global float * inp_img,
+               __global float * out_img,
+               const unsigned int sxs,
+	       const unsigned int syz,
+               const unsigned int lxs,
+               const unsigned int xl,
+               const unsigned int yt)
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int lidx = idy * lxs + idx;
+
+    const int sidx = (idy - yt) * sxs + (idx - xl); 
+
+    if( idy < yt)
+    {
+        if(idx < xl) 
+            out_img[lidx] = inp_img[0]; // top left
+        else if(idx >= (xl+sxs))
+            out_img[lidx] = inp_img[sxs-1]; // top right
+        else 
+            out_img[lidx] = inp_img[idx-xl]; // top center
+        return;
+    }
+
+    if( idy >= (yt + syz))
+    {
+        if(idx < xl) 
+            out_img[lidx] =  inp_img[(syz-1)*sxs]; // bottom left
+        else if(idx >= (xl+sxs)) 
+            out_img[lidx] = inp_img[syz*sxs-1]; // bottom right
+        else
+            out_img[lidx] = inp_img[(syz-1)*sxs +(idx-xl)]; // bottom center
+        return;
+    }
+
+    if(idx < xl)
+    {
+       out_img[lidx] = inp_img[(idy-yt)*sxs]; // center left
+       return;
+    }
+   
+    if(idx >= (xl+sxs))
+    {
+       out_img[lidx] = inp_img[(idy-yt)*sxs + (sxs -1)]; // center right
+       return;
+    }
+
+    out_img[lidx] = inp_img[sidx]; // center center          
+
+}

src/ufo-filter-padding-2d.c

@@ -0,0 +1,387 @@
+#include <gmodule.h>
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <ufo/ufo-resource-manager.h>
+#include <ufo/ufo-filter.h>
+#include <ufo/ufo-buffer.h>
+#include "ufo-filter-padding-2d.h"
+
+/**
+ * SECTION:ufo-filter-padding-2d
+ * @Short_description:
+ * @Title: padding2d
+ *
+ * Detailed description.
+ */
+
+typedef enum {
+	PADDING_ZERO = 0,
+	PADDING_CONST,
+        PADDING_GAVG,
+        PADDING_BREP	
+} PaddingMode;
+
+
+struct _UfoFilterPadding2DPrivate 
+{
+   // extent adds	
+   guint xl;
+   guint xr;
+   guint yt;
+   guint yb;
+   //
+   PaddingMode mode; 
+   // padding constant
+   float pconst;
+   cl_kernel kernel_iconst;
+   cl_kernel kernel_cpyimg;
+   cl_kernel kernel_brep; 
+
+};
+
+G_DEFINE_TYPE(UfoFilterPadding2D, ufo_filter_padding_2d, UFO_TYPE_FILTER)
+
+#define UFO_FILTER_PADDING_2D_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_PADDING_2D, UfoFilterPadding2DPrivate))
+
+enum {
+    PROP_0=0,
+    PROP_XL,
+    PROP_XR,
+    PROP_YT,
+    PROP_YB,
+    PROP_MODE,
+    PROP_PCONST,
+    N_PROPERTIES
+};
+
+static GParamSpec *padding_2d_properties[N_PROPERTIES] = { NULL, };
+
+
+static void ufo_filter_padding_2d_initialize(UfoFilter *filter)
+{
+    /* Here you can code, that is called for each newly instantiated filter */
+    
+    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+
+
+     self->priv->kernel_iconst = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_init_const", &error);
+     self->priv->kernel_cpyimg = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_copy_in", &error);
+     self->priv->kernel_brep   = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_brep", &error);
+  
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+
+}
+
+/*
+ * This is the main method in which the filter processes one buffer after
+ * another.
+ */
+static void ufo_filter_padding_2d_process(UfoFilter *filter)
+{
+    g_return_if_fail(UFO_IS_FILTER(filter));
+    UfoChannel *input_channel = ufo_filter_get_input_channel(filter);
+    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
+
+    UfoBuffer *input = ufo_channel_get_input_buffer(input_channel);
+
+    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(filter);
+    UfoFilterPadding2DPrivate *priv = self->priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(self);
+    //g_message("%d %d %d %d %d %f", priv->xl, priv->xr, priv->yt, priv->yb, priv->mode,priv->pconst);
+
+    guint pxl = priv->xl;
+    guint pxr = priv->xr;
+    guint pyt = priv->yt;
+    guint pyb = priv->yb; 
+    PaddingMode mode = priv->mode;
+    float pval = priv->pconst;
+
+    guint   ndims = 0;
+    guint * isize = NULL;
+    ufo_buffer_get_dimensions(input, &ndims, &isize);
+
+    guint ixs = isize[0];
+    guint iys = isize[1];
+
+    guint osize[2] = {ixs+pxl+pxr, iys+pyt+pyb};
+    guint oxs = osize[0];
+    guint oys = osize[1];
+
+    ufo_channel_allocate_output_buffers(output_channel, ndims, osize);
+
+    size_t global_work_size_small[2] = { (size_t) ixs, (size_t) iys };
+    size_t global_work_size_large[2] = { (size_t) oxs, (size_t) oys };
+
+
+     cl_kernel k_iconst = self->priv->kernel_iconst;
+     cl_kernel k_cpyimg = self->priv->kernel_cpyimg;
+     cl_kernel k_brep   = self->priv->kernel_brep;
+
+
+    // init padding value  for different modes
+    if(mode== PADDING_ZERO)
+	   pval = 0.0;
+
+    while (input != NULL) {
+
+  	UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
+	cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input, command_queue);
+	cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output, command_queue);
+
+        if(mode==PADDING_GAVG)
+	{
+            float *indata = ufo_buffer_get_host_array(input, command_queue);
+
+	    guint psz = ixs * iys;
+            float sum = 0;
+            for(guint i =0; i < psz; i++) sum += indata[i]; 
+            pval = sum/(float)psz;
+        }
+
+         // processing itself
+	if((mode == PADDING_ZERO) || (mode == PADDING_CONST) || (mode == PADDING_GAVG))
+	{	
+	    /// fill with constant
+           CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 0, sizeof(cl_mem), (void *) &output_mem));
+           CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 1, sizeof(int),   &oxs));
+           CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 2, sizeof(float), &pval));
+
+           cl_event event;
+           CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_iconst,
+                                          2, NULL, global_work_size_large, NULL,
+                                          0, NULL, &event));
+           ufo_buffer_attach_event(output, event);
+
+           /// copy old image
+	   CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 0, sizeof(cl_mem), (void *) &input_mem));
+           CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 1, sizeof(cl_mem), (void *) &output_mem));
+           CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 2, sizeof(int),   &ixs));
+	   CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 3, sizeof(int),   &oxs));
+	   CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 4, sizeof(int),   &pxl));
+	   CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 5, sizeof(int),   &pyt));
+        
+	   CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_cpyimg,
+                                   2, NULL, global_work_size_small, NULL,
+	                           0, NULL, &event));
+           ufo_buffer_attach_event(output, event);
+	}
+
+	if(mode == PADDING_BREP)
+	{
+	    CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 0, sizeof(cl_mem), (void *) &input_mem));
+            CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 1, sizeof(cl_mem), (void *) &output_mem));
+            CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 2, sizeof(int),   &ixs));
+	    CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 3, sizeof(int),   &iys));
+            CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 4, sizeof(int),   &oxs));
+            CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 5, sizeof(int),   &pxl));
+            CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 6, sizeof(int),   &pyt));
+ 
+	    cl_event event;
+            CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_brep,
+                                     2, NULL, global_work_size_large, NULL,
+                                     0, NULL, &event));
+            ufo_buffer_attach_event(output, event);
+	}
+	g_message("ufo-filter-padding-2d: done");
+        
+	ufo_channel_finalize_input_buffer(input_channel, input);
+        ufo_channel_finalize_output_buffer(output_channel, output);
+
+        /* Get new input */
+        input = ufo_channel_get_input_buffer(input_channel);
+    }
+
+    /* Tell subsequent filters, that we are finished */
+    ufo_channel_finish(output_channel);
+    g_free(isize);
+}
+
+static void ufo_filter_padding_2d_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(object);
+
+    switch (property_id) {
+      case PROP_XL:
+            self->priv->xl = g_value_get_uint(value);
+            break;
+      case PROP_XR:
+	    self->priv->xr = g_value_get_uint(value);
+            break;
+      case PROP_YT:
+	    self->priv->yt = g_value_get_uint(value);
+            break;
+      case PROP_YB:
+            self->priv->yb = g_value_get_uint(value);
+            break;
+      case PROP_MODE:
+                 if (!g_strcmp0(g_value_get_string(value), "zero"))
+                   self->priv->mode = PADDING_ZERO;
+            else if (!g_strcmp0(g_value_get_string(value), "const"))
+                   self->priv->mode = PADDING_CONST;
+            else if (!g_strcmp0(g_value_get_string(value), "gavg"))
+                    self->priv->mode = PADDING_GAVG;
+            else if (!g_strcmp0(g_value_get_string(value), "brep"))
+                    self->priv->mode = PADDING_BREP;
+ 	    else 
+        	    G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+	    break;
+      case PROP_PCONST:
+            self->priv->pconst = (float) g_value_get_double(value);
+            break;
+      default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_padding_2d_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(object);
+
+    switch (property_id) {
+        case PROP_XL:
+             g_value_set_uint(value, self->priv->xl);
+	     break;
+        case PROP_XR:
+             g_value_set_uint(value, self->priv->xr);
+             break;
+        case PROP_YT:
+             g_value_set_uint(value, self->priv->yt);
+	     break;
+        case PROP_YB:
+             g_value_set_uint(value, self->priv->yb);
+             break;
+        case PROP_MODE:
+             switch (self->priv->mode) {
+             case PADDING_ZERO:
+                 g_value_set_string(value, "zero");
+                 break;
+	     case PADDING_CONST:
+                 g_value_set_string(value, "const");
+                 break;
+	     case PADDING_GAVG:
+                 g_value_set_string(value, "gavg");
+                 break;
+             case PADDING_BREP:
+                 g_value_set_string(value, "brep");
+                 break;
+	     default:
+                 G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+                 break;
+	     }
+	     break;
+        case PROP_PCONST:
+            g_value_set_double(value, self->priv->pconst);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_padding_2d_class_init(UfoFilterPadding2DClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_padding_2d_set_property;
+    gobject_class->get_property = ufo_filter_padding_2d_get_property;
+    filter_class->initialize = ufo_filter_padding_2d_initialize;
+    filter_class->process = ufo_filter_padding_2d_process;
+
+
+    padding_2d_properties[PROP_XL] =
+        g_param_spec_uint("xl",
+		         "Number of additional pixel on the left hand image side",
+                         "Number of additional pixel on the left hand image side",
+	                  0, 16384, 1,
+			  G_PARAM_READWRITE);
+
+    padding_2d_properties[PROP_XR] =
+         g_param_spec_uint("xr",
+			 "Number of additional pixel on the right hand image side",
+                         "Number of additional pixel on the right hand image side",
+			 0, 16384, 1,                                            
+			 G_PARAM_READWRITE);
+
+    padding_2d_properties[PROP_YT] =
+         g_param_spec_uint("yt",
+			 "Number of additional pixel on the top image side",
+                         "Number of additional pixel on the top image side",
+                          0, 16384, 1,
+			  G_PARAM_READWRITE);
+
+    padding_2d_properties[PROP_YB] =
+         g_param_spec_uint("yb",                            
+	                   "Number of additional pixel on the bottom image side",
+                           "Number of additional pixel on the bottom image side",
+			   0, 16384, 1,                                                                              
+                           G_PARAM_READWRITE);
+
+    padding_2d_properties[PROP_MODE] =
+         g_param_spec_string("mode",
+ 		            "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ",
+                            "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ",
+			    "zero",
+    		           G_PARAM_READWRITE);
+
+    padding_2d_properties[PROP_PCONST] = 
+        g_param_spec_double("pconst",
+            "Padding constant",
+            "Padding constant",
+            -320000.0,   
+             320000.0,  
+             0.0,   
+            G_PARAM_READWRITE);
+
+    g_object_class_install_property(gobject_class, PROP_XL, padding_2d_properties[PROP_XL]);
+    g_object_class_install_property(gobject_class, PROP_XR, padding_2d_properties[PROP_XR]);
+    g_object_class_install_property(gobject_class, PROP_YT, padding_2d_properties[PROP_YT]);
+    g_object_class_install_property(gobject_class, PROP_YB, padding_2d_properties[PROP_YB]);
+    g_object_class_install_property(gobject_class, PROP_MODE, padding_2d_properties[PROP_MODE]);
+    g_object_class_install_property(gobject_class, PROP_PCONST, padding_2d_properties[PROP_PCONST]);
+ 
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterPadding2DPrivate));
+}
+
+static void ufo_filter_padding_2d_init(UfoFilterPadding2D *self)
+{
+    UfoFilterPadding2DPrivate *priv = self->priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(self);
+
+    // set default values
+    priv->xl = 1;
+    priv->xr = 1;
+    priv->yt = 1;
+    priv->yb = 1;
+
+    priv->mode   = PADDING_ZERO;
+    priv->pconst = 0.0; 
+
+    priv->kernel_iconst = NULL;
+    priv->kernel_cpyimg = NULL;
+    priv->kernel_brep   = NULL;
+
+    // register filter input and output 
+    ufo_filter_register_input(UFO_FILTER(self), "image", 2);
+    ufo_filter_register_output(UFO_FILTER(self), "pimage", 2);
+}
+
+G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+{
+    return g_object_new(UFO_TYPE_FILTER_PADDING_2D, NULL);
+}

src/ufo-filter-padding-2d.h

@@ -0,0 +1,40 @@
+#ifndef __UFO_FILTER_PADDING_2D_H
+#define __UFO_FILTER_PADDING_2D_H
+
+#include <glib.h>
+#include <glib-object.h>
+
+#include <ufo/ufo-filter.h>
+
+#define UFO_TYPE_FILTER_PADDING_2D             (ufo_filter_padding_2d_get_type())
+#define UFO_FILTER_PADDING_2D(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_PADDING_2D, UfoFilterPadding2D))
+#define UFO_IS_FILTER_PADDING_2D(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_PADDING_2D))
+#define UFO_FILTER_PADDING_2D_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_PADDING_2D, UfoFilterPadding2DClass))
+#define UFO_IS_FILTER_PADDING_2D_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_PADDING_2D))
+#define UFO_FILTER_PADDING_2D_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_PADDING_2D, UfoFilterPadding2DClass))
+
+typedef struct _UfoFilterPadding2D           UfoFilterPadding2D;
+typedef struct _UfoFilterPadding2DClass      UfoFilterPadding2DClass;
+typedef struct _UfoFilterPadding2DPrivate    UfoFilterPadding2DPrivate;
+
+struct _UfoFilterPadding2D {
+    /*< private >*/
+    UfoFilter parent_instance;
+
+    UfoFilterPadding2DPrivate *priv;
+};
+
+/**
+ * UfoFilterPadding2DClass:
+ *
+ * #UfoFilterPadding2D class
+ */
+struct _UfoFilterPadding2DClass {
+    /*< private >*/
+    UfoFilterClass parent_class;
+};
+
+GType ufo_filter_padding_2d_get_type(void);
+UfoFilter *ufo_filter_plugin_new(void);
+
+#endif

src/ufo-filter-scale.host

@@ -0,0 +1,212 @@
+#include <gmodule.h>
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <ufo/ufo-filter.h>
+// #include <ufo/ufo-element.h>
+#include <ufo/ufo-buffer.h>
+#include <ufo/ufo-resource-manager.h>
+
+#include "ufo-filter-scale.h"
+
+/**
+ * SECTION:ufo-filter-scale
+ * @Short_description: Scale image values 
+ * @Title: scale
+ *
+ * Scale input image values. The output
+ * is a new image.
+ * #UfoFilterScale: params.
+*/
+
+struct _UfoFilterScalePrivate {
+    float scale;
+    cl_kernel kernel;
+};
+
+GType ufo_filter_scale_get_type(void) G_GNUC_CONST;
+
+/* Inherit from UFO_TYPE_FILTER */
+G_DEFINE_TYPE(UfoFilterScale, ufo_filter_scale, UFO_TYPE_FILTER);
+
+#define UFO_FILTER_SCALE_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_SCALE, UfoFilterScalePrivate))
+
+enum {
+    PROP_0 =0,
+    PROP_SCALE,
+    N_PROPERTIES
+};
+
+static GParamSpec *scale_properties[N_PROPERTIES] = { NULL, };
+
+
+/* 
+ * virtual methods 
+ */
+static void ufo_filter_scale_initialize(UfoFilter *filter)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+    self->priv->kernel = NULL;
+
+    ufo_resource_manager_add_program(manager, "scale.cl", NULL,  &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+        return;
+    }
+
+    self->priv->kernel = ufo_resource_manager_get_kernel(manager, "scale", &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+}
+
+static void ufo_filter_scale_process(UfoFilter *filter)
+{
+    g_return_if_fail(UFO_IS_FILTER(filter));
+   UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+    UfoChannel *input_channel = ufo_filter_get_input_channel(filter);
+    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
+
+     UfoBuffer *input  = ufo_channel_get_input_buffer(input_channel);
+     ufo_channel_allocate_output_buffers_like(output_channel, input);
+  
+     const gint32 num_elements = ufo_buffer_get_size(input) / sizeof(float);
+     float scale = (float) self->priv->scale;
+     while(input !=NULL)
+     {
+         float *in_data = ufo_buffer_get_host_array(input, command_queue);
+
+	 UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
+	  
+	 /* This avoids an unneccessary GPU-to-host transfer */
+	  ufo_buffer_invalidate_gpu_data(output);
+
+         float *out_data = ufo_buffer_get_host_array(output, command_queue);
+
+	 for (int i = 0; i < num_elements; i++) 
+		out_data[i] = scale * in_data[i];				         
+
+         ufo_channel_finalize_input_buffer(input_channel, input);
+         ufo_channel_finalize_output_buffer(output_channel, output);
+	 input = ufo_channel_get_input_buffer(input_channel); // read next
+	 g_message("ufo-filter-scale: processing is completed");     
+     }	     
+
+/*     UfoBuffer *oimage = NULL;
+
+
+    gint32 width, height;
+    UfoBuffer *buffer = (UfoBuffer *) g_async_queue_pop(input_queue);
+    while (!ufo_buffer_is_finished(buffer)) {
+        if (self->priv->kernel != NULL) {
+            float scale
+		    = (float) self->priv->scale;
+            size_t global_work_size[2];
+
+            ufo_buffer_get_dimensions(buffer, &width, &height);
+            global_work_size[0] = width;
+            global_work_size[1] = height;
+
+            cl_mem buffer_mem = (cl_mem) ufo_buffer_get_gpu_data(buffer, command_queue);
+            cl_int err = CL_SUCCESS;
+
+            err = clSetKernelArg(self->priv->kernel, 0, sizeof(float), &scale);
+            err = clSetKernelArg(self->priv->kernel, 1, sizeof(cl_mem), (void *) &buffer_mem);
+            err = clEnqueueNDRangeKernel(command_queue,
+                self->priv->kernel,
+                2, NULL, global_work_size, NULL,
+                0, NULL, &event);
+
+            ufo_filter_account_gpu_time(filter, (void **) &event);
+        }
+        g_async_queue_push(output_queue, buffer);
+        buffer = (UfoBuffer *) g_async_queue_pop(input_queue);
+    }
+   // g_message("ufo-filter-scale: 0s/%fs", ufo_filter_get_gpu_time(filter));
+   // g_async_queue_push(output_queue, buffer);
+*/
+    ufo_channel_finish(output_channel);
+
+}
+
+static void ufo_filter_scale_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            self->priv->scale = (float) g_value_get_double(value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            g_value_set_double(value, (double) self->priv->scale);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_class_init(UfoFilterScaleClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_scale_set_property;
+    gobject_class->get_property = ufo_filter_scale_get_property;
+    filter_class->initialize = ufo_filter_scale_initialize;
+    filter_class->process = ufo_filter_scale_process;
+
+    scale_properties[PROP_SCALE] = 
+      g_param_spec_double("scale",
+        "Scale",
+        "Scale for each pixel",
+        -5.0,   /* minimum */
+         10.0,   /* maximum */
+         1.0,   /* default */
+        G_PARAM_READWRITE);
+
+    g_object_class_install_property(gobject_class, PROP_SCALE, scale_properties[PROP_SCALE]);
+
+    /* install private data */
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterScalePrivate));
+}
+
+static void ufo_filter_scale_init(UfoFilterScale *self)
+{
+    UfoFilterScalePrivate *priv = self->priv = UFO_FILTER_SCALE_GET_PRIVATE(self);
+    priv->scale = 1.0;
+    priv->kernel = NULL;
+    
+    ufo_filter_register_input (UFO_FILTER(self), "image", 2);
+    ufo_filter_register_output(UFO_FILTER(self), "image", 2);
+}
+
+G_MODULE_EXPORT  *ufo_filter_plugin_new(void) 
+{
+    return g_object_new(UFO_TYPE_FILTER_SCALE, NULL);
+}