Minor corrections. Convolution filter is added.

src/CMakeLists.txt

@@ -7,6 +7,7 @@ set(ufofilter_SRCS
     ufo-filter-3d-edf-writer.c
     ufo-filter-padding-2d.c
     ufo-filter-lamino-ramp.c
+    ufo-filter-lamino-ft-conv.c
     )
 
 set(ufofilter_KERNELS
@@ -14,6 +15,7 @@ set(ufofilter_KERNELS
     lamino_bp_generic.cl
     padding_2d.cl
     lamino_ramp.cl
+    lamino_ft_conv.cl
     )
 
 set(ufofilter_LIBS

src/lamino_bp_generic.cl

@@ -7,53 +7,6 @@
 
 #include <lamino-filter-def.h>
 
-/*typedef struct
-{
-    // rotation angle
-    float phi;
-        
-    // laminographic angle
-    float theta;
-        
-    // misalignment angle
-    float psi;
-        
-    // modified lamino angle 
-    float alpha;
-
-    // rotation angular step
-    float angle_step;
-
-    // rotation matrix
-    float mat_0;
-    float mat_1;
-    float mat_2;
-    float mat_3;
-    float mat_4;
-    float mat_5;
-
-   // reconstructed volume size
-    int vol_sx;
-    int vol_sy;
-    int vol_sz;
-
-    // volume origin
-    float vol_ox;
-    float vol_oy;
-    float vol_oz;
-
-    // projection origin
-    float proj_ox;
-    float proj_oy;
-
-    // projection sizes
-    int proj_sx;
-    int proj_sy;
-
-
-} CLParameters;
-*/
-
 
 __kernel void lamino_bp_generic ( __global float *proj,
 				  __global float *volume,
@@ -70,6 +23,7 @@ __kernel void lamino_bp_generic ( __global float *proj,
     const long int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
     //const int idx = (vY * vSX) + vX;
 
+
     float newz = (float)vZ - param->vol_oz;
 
     float newz_matr02 = newz * param->mat_2 + param->proj_ox;
@@ -104,6 +58,7 @@ __kernel void lamino_bp_generic ( __global float *proj,
     result += proj[base + param->proj_sx + 1] * xf_1 * yf_1;
     volume[idx] +=  result; 
 
+
 }
 
 __kernel void lamino_clean_vol(__global float *volume)

src/lamino_ft_conv.cl

@@ -0,0 +1,16 @@
+__kernel void lamino_c( __global float * in, __global float * flt,  __global float * out, const unsigned int width)
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int index = idy * width + idx;
+
+
+    if(idx % 2 == 0) // real part
+    {
+         out[index]= in[index] * flt[index] - in[index+1] * flt[index+1];
+    }
+    else // imaginary part
+    {
+        out[index]=  in[index] * flt[index - 1] + in[index-1] * flt[index];
+    }
+}

src/lamino_ramp.cl

@@ -2,6 +2,8 @@
 __kernel void lamino_ramp_create_filter(
                 __global float * flt,
                 const unsigned int width,
+		const unsigned int fwidth,
+		const unsigned int height,
                 const float theta,
 		const float tau)
 {
@@ -9,29 +11,31 @@ __kernel void lamino_ramp_create_filter(
     const int idy = get_global_id(1);
     const int index = idy * width + idx;
 
+    const float scale = (float)width/(float)height;
+
     flt[index] = 0;
     if( idy > 0 ) return;
 
     const float sin2 = -sin(theta) / 2.;
     const float const_pi_tau= M_PI_F * M_PI_F * tau;
-    const uint quatw = width / 4;
+    const uint quatw = fwidth / 2;
 
     if(idx == 0)
     {
-        flt[idx] = -sin2 / ( 4 * tau);
+        flt[idx] =scale* -sin2 / ( 4 * tau);
         return;
     }
 
     if( (idx <= quatw) && ( (idx%2) != 0))
     {
-        flt[idx] = sin2/((float)idx*(float)idx*const_pi_tau);
+        flt[idx] = scale*sin2/((float)idx*(float)idx*const_pi_tau);
         return;       
     }
 
     int x = width - idx; 
     if( (idx >= 3*quatw) && ( (x%2) != 0))
     {
-        flt[idx] = sin2/((float)x*(float)x*const_pi_tau);
+        flt[idx] = scale*sin2/((float)x*(float)x*const_pi_tau);
         return; 
     }
 

src/ufo-filter-lamino-bp-generic.c

@@ -161,12 +161,12 @@ static void ufo_filter_lamino_bp_generic_process(UfoFilter *filter)
     int vSY = priv->params.vol_sy; 
     int vSZ = priv->params.vol_sz; 
 
-g_message("%d %d %d", vSX, vSY, vSZ);    
 
     // shift volume origin to the box center
-    params->vol_ox += (float)vSX/2.0;
-    params->vol_oy += (float)vSY/2.0;
-    params->vol_oz += (float)vSZ/2.0;
+    // moved to user scipt part
+    // params->vol_ox += (float)vSX/2.0;
+    // params->vol_oy += (float)vSY/2.0;
+    // params->vol_oz += (float)vSZ/2.0;
 
 
     // allocate memory for the reconstructed volume
@@ -186,12 +186,12 @@ g_message("%d %d %d", vSX, vSY, vSZ);
     CHECK_OPENCL_ERROR(clSetKernelArg(clean_vol_kernel, 0, sizeof(cl_mem), (void *) &output_mem));
     CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, clean_vol_kernel,
                                               vol_num_dims, NULL, global_work_size, NULL, 0, NULL, &event));
-     // whait until kernel finishes its job
+     // clFinish(command_queue); 
+    // whait until kernel finishes its job
      ufo_buffer_attach_event(output, event);
       
     // setup backprojection kernel    
     cl_kernel kernel = priv->bp_kernel;
-    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &output_mem));
 
     // projection conter
     int proj_idx = 0;
@@ -213,6 +213,9 @@ g_message("%d %d %d", vSX, vSY, vSZ);
         params->mat_4 = -sg * sf + cg * st * cf;
         params->mat_5 =  cg * ct;
 
+
+	//g_message("%d %d %d", vSX, vSY, vSZ);
+        //g_message("%f %f %f", params->vol_ox, params->vol_oy, params->vol_oz);
 	//g_message("%f %f %f %f ", params->theta,  params->phi,  params->psi,  params->alpha);
 	//g_message("%f %f %f", params->proj_ox, params->proj_oy, params->angle_step);
 
@@ -222,11 +225,13 @@ g_message("%d %d %d", vSX, vSY, vSZ);
 	// copy projection to GPU
         cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input,  command_queue);
         CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&input_mem));
+	CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &output_mem));
         CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&param_mem));
         // call backprojection routine
 	g_message("processing of %d-th projection", proj_idx);
         CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
                                               vol_num_dims, NULL, global_work_size, NULL, 0, NULL, &event));
+	 // clFinish(command_queue);
 	// wait untill it finishes its job
 	ufo_buffer_attach_event(output, event);
 
@@ -370,12 +375,6 @@ static void ufo_filter_lamino_bp_generic_initialize(UfoFilter *filter)
 //    ufo_resource_manager_add_program(manager, "lamino_bp_generic.cl",
   //  	 "-I /home/timurttv/_UFO-Project/framework/ufo-filters-am-installed/lib/ufo", &error);
 
-    if (error != NULL) {
-            g_warning("%s", error->message);
-            g_error_free(error);
-            return;
-     }
-
      self->priv->bp_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_bp_generic", &error);
      self->priv->clean_vol_kernel =  ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_clean_vol", &error);
      self->priv->norm_vol_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_norm_vol", &error);

src/ufo-filter-lamino-ft-conv.c

@@ -0,0 +1,192 @@
+#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-lamino-ft-conv.h"
+
+/**
+ * SECTION:ufo-filter-lamino-f-t-conv
+ * @Short_description:
+ * @Title: laminoftconv
+ *
+ * Detailed description.
+ */
+
+struct _UfoFilterLaminoFTConvPrivate {
+ //  float example;	
+    cl_kernel kernel;
+};
+
+G_DEFINE_TYPE(UfoFilterLaminoFTConv, ufo_filter_lamino_ft_conv, UFO_TYPE_FILTER)
+
+#define UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConvPrivate))
+
+enum {
+    PROP_0,
+//    PROP_EXAMPLE, 
+    N_PROPERTIES
+};
+
+// static GParamSpec *lamino_ft_conv_properties[N_PROPERTIES] = { NULL, };
+
+
+static void ufo_filter_lamino_ft_conv_initialize(UfoFilter *filter)
+{
+
+    UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+    self->priv->kernel = ufo_resource_manager_get_kernel(manager, "lamino_ft_conv.cl", "lamino_c", &error);
+ 
+
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+}
+
+
+static void ufo_filter_lamino_ft_conv_process(UfoFilter *flt)
+{
+
+    g_return_if_fail(UFO_IS_FILTER(flt));
+    UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(flt);
+
+
+    UfoChannel *input_channel  = ufo_filter_get_input_channel_by_name(flt, "image");
+    UfoChannel *filter_channel = ufo_filter_get_input_channel_by_name(flt, "filter");
+    UfoChannel *output_channel = ufo_filter_get_output_channel(flt);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(flt);
+
+    UfoBuffer *input = ufo_channel_get_input_buffer(input_channel);
+    UfoBuffer *filter = ufo_channel_get_input_buffer(filter_channel);
+
+    guint num_dims = 0;
+    guint *dim_size = NULL;
+    ufo_buffer_get_dimensions(input, &num_dims, &dim_size);
+    ufo_channel_allocate_output_buffers(output_channel, 2, dim_size);
+
+    guint *fdim_size= NULL;
+    ufo_buffer_get_dimensions(input, &num_dims, &fdim_size);
+
+   if( (dim_size[0] != fdim_size[0]) || (dim_size[1] != fdim_size[1]))
+	   g_error("Filter and image sizes are different");
+
+    size_t global_work_size[2] = { (size_t) dim_size[0], (size_t) dim_size[1] };
+    guint width = dim_size[0];
+    cl_kernel kernel = self->priv->kernel;
+
+    cl_mem filter_mem  = (cl_mem) ufo_buffer_get_device_array(filter,  command_queue);
+
+    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);
+
+        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &input_mem));
+        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &filter_mem));
+        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &output_mem));
+        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(int), &width));
+
+        cl_event event;
+        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
+                                              2, NULL, global_work_size, NULL,
+                                              0, NULL, &event));
+	 clFinish(command_queue);
+        //ufo_buffer_attach_event(output, event);
+
+
+        ufo_channel_finalize_input_buffer(input_channel, input);
+        ufo_channel_finalize_output_buffer(output_channel, output);
+        input = ufo_channel_get_input_buffer(input_channel);
+    }
+
+    ufo_channel_finalize_input_buffer(filter_channel, filter);
+    ufo_channel_finish(output_channel);
+
+    g_free(dim_size);
+    g_free(fdim_size);
+}
+
+static void ufo_filter_lamino_ft_conv_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+
+    // UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(object);
+
+    switch (property_id) {
+        /*case PROP_EXAMPLE:
+            self->priv->example = g_value_get_double(value);
+            break;*/
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_lamino_ft_conv_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    // UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(object);
+
+    switch (property_id) {
+        /*case PROP_EXAMPLE:
+            g_value_set_double(value, self->priv->example);
+            break;*/
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_lamino_ft_conv_class_init(UfoFilterLaminoFTConvClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_lamino_ft_conv_set_property;
+    gobject_class->get_property = ufo_filter_lamino_ft_conv_get_property;
+    filter_class->initialize = ufo_filter_lamino_ft_conv_initialize;
+    filter_class->process = ufo_filter_lamino_ft_conv_process;
+
+      /* lamino_ft_conv_properties[PROP_EXAMPLE] = 
+        g_param_spec_double("example",
+            "This is an example property",
+            "You should definately replace this with some meaningful property",
+            -1.0,   
+             1.0,  
+             1.0,   
+            G_PARAM_READWRITE);
+
+    g_object_class_install_property(gobject_class, PROP_EXAMPLE, lamino_ft_conv_properties[PROP_EXAMPLE]); */
+
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoFTConvPrivate));
+
+}
+
+static void ufo_filter_lamino_ft_conv_init(UfoFilterLaminoFTConv *self)
+{
+    UfoFilterLaminoFTConvPrivate *priv = self->priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(self);
+    // priv->example = 1.0;
+    priv->kernel = NULL;
+
+    ufo_filter_register_input(UFO_FILTER(self),  "filter", 2);
+    ufo_filter_register_input(UFO_FILTER(self),  "image",  2);
+    ufo_filter_register_output(UFO_FILTER(self), "oimage", 2);
+}
+
+G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+{
+    return g_object_new(UFO_TYPE_FILTER_LAMINO_FT_CONV, NULL);
+}

src/ufo-filter-lamino-ft-conv.h

@@ -0,0 +1,40 @@
+#ifndef __UFO_FILTER_LAMINO_FT_CONV_H
+#define __UFO_FILTER_LAMINO_FT_CONV_H
+
+#include <glib.h>
+#include <glib-object.h>
+
+#include <ufo/ufo-filter.h>
+
+#define UFO_TYPE_FILTER_LAMINO_FT_CONV             (ufo_filter_lamino_ft_conv_get_type())
+#define UFO_FILTER_LAMINO_FT_CONV(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConv))
+#define UFO_IS_FILTER_LAMINO_FT_CONV(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV))
+#define UFO_FILTER_LAMINO_FT_CONV_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConvClass))
+#define UFO_IS_FILTER_LAMINO_FT_CONV_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_LAMINO_FT_CONV))
+#define UFO_FILTER_LAMINO_FT_CONV_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConvClass))
+
+typedef struct _UfoFilterLaminoFTConv           UfoFilterLaminoFTConv;
+typedef struct _UfoFilterLaminoFTConvClass      UfoFilterLaminoFTConvClass;
+typedef struct _UfoFilterLaminoFTConvPrivate    UfoFilterLaminoFTConvPrivate;
+
+struct _UfoFilterLaminoFTConv {
+    /*< private >*/
+    UfoFilter parent_instance;
+
+    UfoFilterLaminoFTConvPrivate *priv;
+};
+
+/**
+ * UfoFilterLaminoFTConvClass:
+ *
+ * #UfoFilterLaminoFTConv class
+ */
+struct _UfoFilterLaminoFTConvClass {
+    /*< private >*/
+    UfoFilterClass parent_class;
+};
+
+GType ufo_filter_lamino_ft_conv_get_type(void);
+UfoFilter *ufo_filter_plugin_new(void);
+
+#endif

src/ufo-filter-lamino-ramp.c

@@ -22,6 +22,7 @@ struct _UfoFilterLaminoRampPrivate {
 
     // filter extent
     guint width; // is pow of 2
+    guint fill_width;
     guint height;
     // laminographic angle
     float theta;
@@ -37,6 +38,7 @@ G_DEFINE_TYPE(UfoFilterLaminoRamp, ufo_filter_lamino_ramp, UFO_TYPE_FILTER)
 enum {
     PROP_0,
     PROP_WIDTH,
+    PROP_FILL_WIDTH,
     PROP_HEIGHT,
     PROP_THETA,
     PROP_TAU,
@@ -81,6 +83,8 @@ static void ufo_filter_lamino_ramp_process(UfoFilter *filter)
     if(!is_power_of_two(width))
 	    g_warning("filter width is not the power of two");
 
+    guint fwidth = priv->fill_width;
+
     guint height = priv->height;
     if(!is_power_of_two(height))
 	              g_warning("filter height is not the power of two");
@@ -101,8 +105,10 @@ static void ufo_filter_lamino_ramp_process(UfoFilter *filter)
     
     CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &output_mem));   
     CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(int), &width));
-    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(float), &theta));
-    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(float), &tau)); 
+    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(int), &fwidth));
+    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(int), &height));
+    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 4, sizeof(float), &theta));
+    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 5, sizeof(float), &tau)); 
     
     cl_event event; 
     CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel, 
@@ -126,6 +132,9 @@ static void ufo_filter_lamino_ramp_set_property(GObject *object,
         case PROP_WIDTH:
             self->priv->width = g_value_get_uint(value);
             break;
+        case PROP_FILL_WIDTH:
+            self->priv->fill_width = g_value_get_uint(value);
+            break;
        case PROP_HEIGHT:
             self->priv->height = g_value_get_uint(value);
             break;
@@ -152,6 +161,9 @@ static void ufo_filter_lamino_ramp_get_property(GObject *object,
 	case PROP_WIDTH:
 	     g_value_set_uint(value, self->priv->width);
 	     break;
+        case PROP_FILL_WIDTH:
+	     g_value_set_uint(value, self->priv->fill_width);
+	     break;
         case PROP_HEIGHT:
              g_value_set_uint(value,  self->priv->height);
              break;
@@ -183,6 +195,13 @@ static void ufo_filter_lamino_ramp_class_init(UfoFilterLaminoRampClass *klass)
         "Width of the 2D image filter (power of 2)",
 	1, 32768, 1.0,
 	G_PARAM_READWRITE);
+  
+    lamino_ramp_properties[PROP_FILL_WIDTH] =
+        g_param_spec_uint("fwidth",
+       "Filling width of the 2D image filter",
+       "Filling width of the 2D image filter",
+        1, 32768, 1.0,
+        G_PARAM_READWRITE);
 
      lamino_ramp_properties[PROP_HEIGHT] =
          g_param_spec_uint("height",
@@ -208,6 +227,7 @@ static void ufo_filter_lamino_ramp_class_init(UfoFilterLaminoRampClass *klass)
             G_PARAM_READWRITE);
 
     g_object_class_install_property(gobject_class, PROP_WIDTH, lamino_ramp_properties[PROP_WIDTH]);
+    g_object_class_install_property(gobject_class, PROP_FILL_WIDTH, lamino_ramp_properties[PROP_FILL_WIDTH]);
     g_object_class_install_property(gobject_class, PROP_HEIGHT, lamino_ramp_properties[PROP_HEIGHT]);
     g_object_class_install_property(gobject_class, PROP_THETA, lamino_ramp_properties[PROP_THETA]);
     g_object_class_install_property(gobject_class, PROP_TAU, lamino_ramp_properties[PROP_TAU]);
@@ -219,7 +239,8 @@ static void ufo_filter_lamino_ramp_init(UfoFilterLaminoRamp *self)
 {
     UfoFilterLaminoRampPrivate *priv = self->priv = UFO_FILTER_LAMINO_RAMP_GET_PRIVATE(self);
 
-    priv->width = 1;
+    priv->width = 4;
+    priv->fill_width=2;
     priv->height = 1;
     priv->theta = 0.0;
     priv->tau = 10.0;

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

@@ -130,6 +130,7 @@ static void ufo_filter_padding_2d_process(UfoFilter *filter)
     if(mode== PADDING_ZERO)
 	   pval = 0.0;
 
+    cl_event event;
     while (input != NULL) {
 
   	UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
@@ -154,11 +155,10 @@ static void ufo_filter_padding_2d_process(UfoFilter *filter)
            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);
+           //ufo_buffer_attach_event(output, event);
 
            /// copy old image
 	   CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 0, sizeof(cl_mem), (void *) &input_mem));
@@ -167,11 +167,12 @@ static void ufo_filter_padding_2d_process(UfoFilter *filter)
 	   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, event);
+           //ufo_buffer_attach_event(output, event2);
 	}
 
 	if(mode == PADDING_BREP)
@@ -184,14 +185,14 @@ static void ufo_filter_padding_2d_process(UfoFilter *filter)
             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);
+           // ufo_buffer_attach_event(output, event);
 	}
-	g_message("ufo-filter-padding-2d: done");
-        
+       
+        clFinish(command_queue); // synchro?
+
 	ufo_channel_finalize_input_buffer(input_channel, input);
         ufo_channel_finalize_output_buffer(output_channel, output);