Port padding filter

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

@@ -19,28 +19,35 @@
  */
 
 typedef enum {
-	PADDING_ZERO = 0,
-	PADDING_CONST,
-        PADDING_GAVG,
-        PADDING_BREP	
+    PADDING_ZERO = 0,
+    PADDING_CONST,
+    PADDING_GAVG,
+    PADDING_BREP
 } PaddingMode;
 
 
-struct _UfoFilterPadding2DPrivate 
+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; 
-
+    guint in_width;
+    guint in_height;
+    guint out_width;
+    guint out_height;
+
+    // extent adds
+    guint xl;
+    guint xr;
+    guint yt;
+    guint yb;
+
+    size_t global_work_size_small[2];
+    size_t global_work_size_large[2];
+
+    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)
@@ -48,7 +55,7 @@ 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_0 = 0,
     PROP_XL,
     PROP_XR,
     PROP_YT,
@@ -61,230 +68,198 @@ enum {
 static GParamSpec *padding_2d_properties[N_PROPERTIES] = { NULL, };
 
 
-static void ufo_filter_padding_2d_initialize(UfoFilter *filter)
+static void
+ufo_filter_padding_2d_initialize (UfoFilter *filter, UfoBuffer *params[], guint **dim_sizes, GError **error)
 {
-    
-    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(filter);
+    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(filter);
     UfoResourceManager *manager = ufo_resource_manager();
-    GError *error = NULL;
+    GError *tmp_error = NULL;
 
+    priv->kernel_iconst = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_init_const", &tmp_error);
+    priv->kernel_cpyimg = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_copy_in", &tmp_error);
+    priv->kernel_brep   = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_brep", &tmp_error);
 
-     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);
+    if (tmp_error != NULL) {
+        g_propagate_error (error, tmp_error);
+        return;
     }
 
+    ufo_buffer_get_2d_dimensions (params[0], &priv->in_width, &priv->in_height);
+    dim_sizes[0][0] = priv->out_width  = priv->xl + priv->in_width + priv->xr;
+    dim_sizes[0][1] = priv->out_height = priv->yt + priv->in_height + priv->yb;
+
+    priv->global_work_size_small[0] = (size_t) priv->in_width;
+    priv->global_work_size_small[1] = (size_t) priv->in_height;
+    priv->global_work_size_large[0] = (size_t) dim_sizes[0][0];
+    priv->global_work_size_large[1] = (size_t) dim_sizes[0][1];
 }
 
 /*
  * This is the main method in which the filter processes one buffer after
  * another.
  */
-static void ufo_filter_padding_2d_process(UfoFilter *filter)
+static UfoEventList *
+ufo_filter_padding_2d_process_gpu(UfoFilter *filter, UfoBuffer *input[], UfoBuffer *output[], gpointer cmd_queue, GError **error)
 {
-    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);
+    g_return_val_if_fail (UFO_IS_FILTER (filter), NULL);
+    cl_command_queue command_queue = (cl_command_queue) cmd_queue;
+    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(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;
+    const PaddingMode mode = priv->mode;
+    const guint pxl = priv->xl;
+    const guint pyt = priv->yt;
     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];
+    const guint ixs = priv->in_width;
+    const guint iys = priv->in_height;
+    const guint oxs = priv->out_width;
 
-    ufo_channel_allocate_output_buffers(output_channel, ndims, osize);
+    // init padding value  for different modes
+    if (mode == PADDING_ZERO)
+        pval = 0.0;
 
-    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_event event;
+    cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array (input[0], command_queue);
+    cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array (output[0], command_queue);
 
+    if (mode == PADDING_GAVG) {
+        float *indata = ufo_buffer_get_host_array(input[0], command_queue);
+        float sum = 0;
+        guint psz = ixs * iys;
 
-     cl_kernel k_iconst = self->priv->kernel_iconst;
-     cl_kernel k_cpyimg = self->priv->kernel_cpyimg;
-     cl_kernel k_brep   = self->priv->kernel_brep;
+        for (guint i =0; i < psz; i++) 
+            sum += indata[i];
 
+        pval = sum / (float)psz;
+    }
 
-    // init padding value  for different modes
-    if(mode== PADDING_ZERO)
-	   pval = 0.0;
+    // processing itself
+    if ((mode == PADDING_ZERO) || (mode == PADDING_CONST) || (mode == PADDING_GAVG)) {
+        cl_kernel k_iconst = priv->kernel_iconst;
+        cl_kernel k_cpyimg = priv->kernel_cpyimg;
+
+        /// 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));
+
+        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_iconst,
+                    2, NULL, priv->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));
+
+        //cl_event event2;
+        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_cpyimg,
+                    2, NULL, priv->global_work_size_small, NULL,
+                    0, NULL, &event));
+        //ufo_buffer_attach_event(output, event2);
+    }
 
-    cl_event event;
-    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));
-
-           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));
-       
-	   //cl_event event2; 
-	   CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_cpyimg,
-                                   2, NULL, global_work_size_small, NULL,
-	                           0, NULL, &event));
-           //ufo_buffer_attach_event(output, event2);
-	}
-
-	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));
- 
-            CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_brep,
-                                     2, NULL, global_work_size_large, NULL,
-                                     0, NULL, &event));
-           // ufo_buffer_attach_event(output, event);
-	}
-       
-        clFinish(command_queue); // synchro?
-
-	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);
+    if (mode == PADDING_BREP) {
+        cl_kernel k_brep   = priv->kernel_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));
+
+        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_brep,
+                    2, NULL, priv->global_work_size_large, NULL,
+                    0, NULL, &event));
+        // ufo_buffer_attach_event(output, event);
     }
 
-    /* Tell subsequent filters, that we are finished */
-    ufo_channel_finish(output_channel);
-    g_free(isize);
+    clFinish(command_queue); // synchro?
+    return NULL;
 }
 
-static void ufo_filter_padding_2d_set_property(GObject *object,
-    guint           property_id,
-    const GValue    *value,
-    GParamSpec      *pspec)
+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:
+        case PROP_XL:
             self->priv->xl = g_value_get_uint(value);
             break;
-      case PROP_XR:
-	    self->priv->xr = g_value_get_uint(value);
+        case PROP_XR:
+            self->priv->xr = g_value_get_uint(value);
             break;
-      case PROP_YT:
-	    self->priv->yt = g_value_get_uint(value);
+        case PROP_YT:
+            self->priv->yt = g_value_get_uint(value);
             break;
-      case PROP_YB:
+        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;
+        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;
+                self->priv->mode = PADDING_CONST;
             else if (!g_strcmp0(g_value_get_string(value), "gavg"))
-                    self->priv->mode = PADDING_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->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:
+        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)
+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;
+            g_value_set_uint(value, self->priv->xl);
+            break;
         case PROP_XR:
-             g_value_set_uint(value, self->priv->xr);
-             break;
+            g_value_set_uint(value, self->priv->xr);
+            break;
         case PROP_YT:
-             g_value_set_uint(value, self->priv->yt);
-	     break;
+            g_value_set_uint(value, self->priv->yt);
+            break;
         case PROP_YB:
-             g_value_set_uint(value, self->priv->yb);
-             break;
+            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;
+            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;
@@ -302,52 +277,51 @@ static void ufo_filter_padding_2d_class_init(UfoFilterPadding2DClass *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;
-
+    filter_class->process_gpu = ufo_filter_padding_2d_process_gpu;
 
     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);
+                "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);
+        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);
+        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);
+        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);
+        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] = 
+    padding_2d_properties[PROP_PCONST] =
         g_param_spec_double("pconst",
-            "Padding constant",
-            "Padding constant",
-            -320000.0,   
-             320000.0,  
-             0.0,   
-            G_PARAM_READWRITE);
+                "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]);
@@ -355,11 +329,12 @@ static void ufo_filter_padding_2d_class_init(UfoFilterPadding2DClass *klass)
     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)
+static void
+ufo_filter_padding_2d_init(UfoFilterPadding2D *self)
 {
     UfoFilterPadding2DPrivate *priv = self->priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(self);
 
@@ -370,18 +345,19 @@ static void ufo_filter_padding_2d_init(UfoFilterPadding2D *self)
     priv->yb = 1;
 
     priv->mode   = PADDING_ZERO;
-    priv->pconst = 0.0; 
+    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);
+    // register filter input and output
+    ufo_filter_register_inputs (UFO_FILTER (self), 2, NULL);
+    ufo_filter_register_outputs (UFO_FILTER (self), 2, NULL);
 }
 
-G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+G_MODULE_EXPORT
+UfoFilter *ufo_filter_plugin_new(void)
 {
     return g_object_new(UFO_TYPE_FILTER_PADDING_2D, NULL);
 }