lamino_ramp filter was developed. Used in the projections filtering stage of the

laminographic reconstruction. Parameters are width x height, detector pixel size,
laminographic angle. The values are given in 2D image (real space).

src/CMakeLists.txt

@@ -6,12 +6,14 @@ set(ufofilter_SRCS
     ufo-filter-lamino-bp-generic.c
     ufo-filter-3d-edf-writer.c
     ufo-filter-padding-2d.c
+    ufo-filter-lamino-ramp.c
     )
 
 set(ufofilter_KERNELS
     scale.cl
     lamino_bp_generic.cl
     padding_2d.cl
+    lamino_ramp.cl
     )
 
 set(ufofilter_LIBS

src/lamino_ramp.cl

@@ -0,0 +1,39 @@
+
+__kernel void lamino_ramp_create_filter(
+                __global float * flt,
+                const unsigned int width,
+                const float theta,
+		const float tau)
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int index = idy * width + idx;
+
+    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;
+
+    if(idx == 0)
+    {
+        flt[idx] = -sin2 / ( 4 * tau);
+        return;
+    }
+
+    if( (idx <= quatw) && ( (idx%2) != 0))
+    {
+        flt[idx] = 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);
+        return; 
+    }
+
+}
+

src/ufo-filter-lamino-ramp.c

@@ -0,0 +1,234 @@
+#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-ramp.h"
+
+/**
+ * SECTION:ufo-filter-lamino-ramp
+ * @Short_description:
+ * @Title: laminoramp
+ *
+ * Detailed description.
+ */
+
+struct _UfoFilterLaminoRampPrivate {
+
+    // filter extent
+    guint width; // is pow of 2
+    guint height;
+    // laminographic angle
+    float theta;
+    // pixel resolution (in um)
+    float tau;	
+    cl_kernel kernel_cf;
+};
+
+G_DEFINE_TYPE(UfoFilterLaminoRamp, ufo_filter_lamino_ramp, UFO_TYPE_FILTER)
+
+#define UFO_FILTER_LAMINO_RAMP_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_LAMINO_RAMP, UfoFilterLaminoRampPrivate))
+
+enum {
+    PROP_0,
+    PROP_WIDTH,
+    PROP_HEIGHT,
+    PROP_THETA,
+    PROP_TAU,
+    N_PROPERTIES
+};
+
+static GParamSpec *lamino_ramp_properties[N_PROPERTIES] = { NULL, };
+
+
+static void ufo_filter_lamino_ramp_initialize(UfoFilter *filter)
+{
+    
+    UfoFilterLaminoRamp *self = UFO_FILTER_LAMINO_RAMP(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+
+    self->priv->kernel_cf = ufo_resource_manager_get_kernel(manager, "lamino_ramp.cl", "lamino_ramp_create_filter", &error);
+
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+}
+
+static int is_power_of_two (guint x)
+{
+    return ((x != 0) && !(x & (x - 1)));
+}
+
+static void ufo_filter_lamino_ramp_process(UfoFilter *filter)
+{
+    g_return_if_fail(UFO_IS_FILTER(filter));
+    UfoFilterLaminoRamp *self = UFO_FILTER_LAMINO_RAMP(filter);
+    UfoFilterLaminoRampPrivate *priv = self->priv = UFO_FILTER_LAMINO_RAMP_GET_PRIVATE(self);  
+    
+    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
+    
+    UfoBuffer *output = NULL;
+
+    guint width  = priv->width; // check width for power of 2
+    if(!is_power_of_two(width))
+	    g_warning("filter width is not the power of two");
+
+    guint height = priv->height;
+    if(!is_power_of_two(height))
+	              g_warning("filter height is not the power of two");
+
+    float theta = priv-> theta;
+    float tau   = priv->tau;
+
+    guint dim_size[2] = {width, height};
+    size_t global_work_size[2] = { (size_t) width, (size_t) height };
+
+    ufo_channel_allocate_output_buffers(output_channel, 2, dim_size);
+
+    cl_kernel kernel = priv->kernel_cf;
+
+    output = ufo_channel_get_output_buffer(output_channel);
+    cl_mem output_mem  = (cl_mem) ufo_buffer_get_device_array(output,  command_queue);     
+        
+    
+    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)); 
+    
+    cl_event event; 
+    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel, 
+			     2, NULL, global_work_size, NULL,
+        		     0, NULL, &event)); 
+    ufo_buffer_attach_event(output, event);
+    
+    ufo_channel_finalize_output_buffer(output_channel, output);
+    ufo_channel_finish(output_channel);
+}
+
+static void ufo_filter_lamino_ramp_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilterLaminoRamp *self = UFO_FILTER_LAMINO_RAMP(object);
+
+    /* Handle all properties accordingly */
+    switch (property_id) {
+        case PROP_WIDTH:
+            self->priv->width = g_value_get_uint(value);
+            break;
+       case PROP_HEIGHT:
+            self->priv->height = g_value_get_uint(value);
+            break;
+       case PROP_THETA:
+            self->priv->theta = (float) g_value_get_double(value);
+            break;
+       case PROP_TAU:
+	    self->priv->tau = (float) g_value_get_double(value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_lamino_ramp_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilterLaminoRamp *self = UFO_FILTER_LAMINO_RAMP(object);
+
+    switch (property_id) {
+	case PROP_WIDTH:
+	     g_value_set_uint(value, self->priv->width);
+	     break;
+        case PROP_HEIGHT:
+             g_value_set_uint(value,  self->priv->height);
+             break;
+	case PROP_THETA:
+	    g_value_set_double(value, self->priv->theta);
+            break;
+        case PROP_TAU:
+            g_value_set_double(value, self->priv->tau);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_lamino_ramp_class_init(UfoFilterLaminoRampClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_lamino_ramp_set_property;
+    gobject_class->get_property = ufo_filter_lamino_ramp_get_property;
+    filter_class->initialize = ufo_filter_lamino_ramp_initialize;
+    filter_class->process = ufo_filter_lamino_ramp_process;
+
+    lamino_ramp_properties[PROP_WIDTH] =
+	g_param_spec_uint("width",
+        "Width of the 2D image filter (power of 2)",
+        "Width of the 2D image filter (power of 2)",
+	1, 32768, 1.0,
+	G_PARAM_READWRITE);
+
+     lamino_ramp_properties[PROP_HEIGHT] =
+         g_param_spec_uint("height",
+        "Height of the 2D image filter",
+	"Height of the 2D image filter",
+	1, 16384, 1.0,
+	G_PARAM_READWRITE);
+
+    lamino_ramp_properties[PROP_THETA] =    
+         g_param_spec_double("theta",
+           "Laminographic angle in radians",
+	   "Resolution (pixel size) in microns",
+	        -4.0 * G_PI, +4.0 * G_PI, 0.0,
+             G_PARAM_READWRITE);
+
+    lamino_ramp_properties[PROP_TAU] = 
+        g_param_spec_double("tau",
+            "Resolution (pixel size) in microns",
+            "Resolution (pixel size) in microns",
+             0.0,        /* minimum */
+             100000.0,   /* maximum */
+             10.0,       /* default */
+            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_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]);
+
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoRampPrivate));
+}
+
+static void ufo_filter_lamino_ramp_init(UfoFilterLaminoRamp *self)
+{
+    UfoFilterLaminoRampPrivate *priv = self->priv = UFO_FILTER_LAMINO_RAMP_GET_PRIVATE(self);
+
+    priv->width = 1;
+    priv->height = 1;
+    priv->theta = 0.0;
+    priv->tau = 10.0;
+
+    // ufo_filter_register_input(UFO_FILTER(self), "input0", 2);
+    ufo_filter_register_output(UFO_FILTER(self), "filter", 2);
+}
+
+G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+{
+    return g_object_new(UFO_TYPE_FILTER_LAMINO_RAMP, NULL);
+}

src/ufo-filter-lamino-ramp.h

@@ -0,0 +1,40 @@
+#ifndef __UFO_FILTER_LAMINO_RAMP_H
+#define __UFO_FILTER_LAMINO_RAMP_H
+
+#include <glib.h>
+#include <glib-object.h>
+
+#include <ufo/ufo-filter.h>
+
+#define UFO_TYPE_FILTER_LAMINO_RAMP             (ufo_filter_lamino_ramp_get_type())
+#define UFO_FILTER_LAMINO_RAMP(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_LAMINO_RAMP, UfoFilterLaminoRamp))
+#define UFO_IS_FILTER_LAMINO_RAMP(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_LAMINO_RAMP))
+#define UFO_FILTER_LAMINO_RAMP_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_LAMINO_RAMP, UfoFilterLaminoRampClass))
+#define UFO_IS_FILTER_LAMINO_RAMP_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_LAMINO_RAMP))
+#define UFO_FILTER_LAMINO_RAMP_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_LAMINO_RAMP, UfoFilterLaminoRampClass))
+
+typedef struct _UfoFilterLaminoRamp           UfoFilterLaminoRamp;
+typedef struct _UfoFilterLaminoRampClass      UfoFilterLaminoRampClass;
+typedef struct _UfoFilterLaminoRampPrivate    UfoFilterLaminoRampPrivate;
+
+struct _UfoFilterLaminoRamp {
+    /*< private >*/
+    UfoFilter parent_instance;
+
+    UfoFilterLaminoRampPrivate *priv;
+};
+
+/**
+ * UfoFilterLaminoRampClass:
+ *
+ * #UfoFilterLaminoRamp class
+ */
+struct _UfoFilterLaminoRampClass {
+    /*< private >*/
+    UfoFilterClass parent_class;
+};
+
+GType ufo_filter_lamino_ramp_get_type(void);
+UfoFilter *ufo_filter_plugin_new(void);
+
+#endif

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

@@ -63,7 +63,6 @@ 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();