Update remaining nodes conv and ramp

src/CMakeLists.txt

@@ -2,14 +2,11 @@ cmake_minimum_required(VERSION 2.6)
 
 # --- Set sources -------------------------------------------------------------
 set(ufofilter_SRCS 
-    ufo-scale-task.c
-    ufo-padding-2d-task.c
     ufo-lamino-bp-task.c
-    #    ufo-filter-lamino-bp-generic.c
-    #ufo-filter-3d-edf-writer.c
-    #ufo-filter-padding-2d.c
-    #ufo-filter-lamino-ramp.c
-    #ufo-filter-lamino-ft-conv.c
+    ufo-lamino-conv-task.c
+    ufo-lamino-ramp-task.c
+    ufo-padding-2d-task.c
+    ufo-scale-task.c
     )
 
 set(ufofilter_KERNELS

src/lamino_ft_conv.cl

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

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

@@ -1,524 +0,0 @@
-#include <gmodule.h>
-#ifdef __APPLE__
-#include <OpenCL/cl.h>
-#else
-#include <CL/cl.h>
-#endif
-
-#include <ufo/ufo-filter-reduce.h>
-#include <ufo/ufo-buffer.h>
-#include <ufo/ufo-resource-manager.h>
-
-#include "ufo-filter-lamino-bp-generic.h"
-
-//#include <stdio.h> // TODO remove later
-
-#include "lamino-filter-def.h"
-#include <math.h>
-
-
-struct _UfoFilterLaminoBPGenericPrivate {
-    // float theta;
-    cl_kernel       bp_kernel;
-    cl_kernel       clean_vol_kernel;
-    cl_kernel       norm_vol_kernel;
-    cl_mem          param_mem;
-    gint            proj_idx;
-    CLParameters    params;
-    gboolean        finished;
-    size_t          global_work_size[3];
-};
-
-GType ufo_filter_lamino_bp_generic_get_type(void) G_GNUC_CONST;
-
-G_DEFINE_TYPE(UfoFilterLaminoBPGeneric, ufo_filter_lamino_bp_generic, UFO_TYPE_FILTER_REDUCE);
-
-#define UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_LAMINO_BP_GENERIC, UfoFilterLaminoBPGenericPrivate))
-
-enum {
-    PROP_0 = 0,
-    PROP_THETA,
-    PROP_PSI,
-    PROP_ANGLE_STEP,
-    PROP_VOL_SX,
-    PROP_VOL_SY,
-    PROP_VOL_SZ,
-    PROP_VOL_OX,
-    PROP_VOL_OY,
-    PROP_VOL_OZ,
-    PROP_PROJ_OX,
-    PROP_PROJ_OY,
-    N_PROPERTIES
-};
-
-static GParamSpec *lamino_bp_generic_properties[N_PROPERTIES] = { NULL, };
-
-
-/* static void testing_cl_platform()
-   {
-   cl_int error = 0;
-
-//Platform Information
-cl_uint numPlatforms;
-cl_platform_id* clSelectedPlatformID = NULL;
-//get the number of available platforms
-
-clGetPlatformIDs(0, NULL, &numPlatforms);
-//alloc memory so we can get the whole list
-clSelectedPlatformID = (cl_platform_id*)malloc(sizeof(cl_platform_id)*numPlatforms);
-//get the list of available platforms
-error = clGetPlatformIDs(numPlatforms, clSelectedPlatformID, NULL);
-
-g_message("Available platforms number: %d", numPlatforms);
-
-char platform_info[128];
-clGetPlatformInfo(clSelectedPlatformID[0], CL_PLATFORM_NAME, sizeof(char)*128, platform_info, NULL);
-g_message("CL_PLATFORM_NAME: %s", platform_info);
-
-cl_uint ciDeviceCount;
-cl_device_id* clDevices =  NULL;
-error = clGetDeviceIDs(clSelectedPlatformID[0], CL_DEVICE_TYPE_GPU, 0, NULL, &ciDeviceCount);
-clDevices = (cl_device_id*) malloc(sizeof(cl_device_id) * ciDeviceCount);
-error = clGetDeviceIDs(clSelectedPlatformID[0], CL_DEVICE_TYPE_GPU, ciDeviceCount, clDevices, &ciDeviceCount);
-
-g_message("Available Devices: %d.",ciDeviceCount);
-char device_info[128];
-cl_uint device_value = 0;
-cl_ulong device_value_ulong = 0;
-size_t device_sizet = 0;
-size_t dimsz[3];
-
-for(unsigned int i=0; i <  ciDeviceCount ; i++)
-{
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_NAME, sizeof(char)*128, device_info, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_NAME: %s", i+1, device_info);
-
-if(clGetDeviceInfo(clDevices[i], CL_DRIVER_VERSION, sizeof(char)*128, device_info, NULL) == CL_SUCCESS)
-g_message("#%d CL_DRIVER_VERSION: %s", i+1, device_info);
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(cl_uint), &device_value, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_CLOCK_FREQUENCY: %dMHz", i+1, device_value);
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(cl_ulong), &device_value_ulong, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: %d kB", i+1, device_value_ulong/(1024));
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &device_value_ulong, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_GLOBAL_MEM_SIZE: %d MB", i+1, device_value_ulong/(1024*1024));
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &device_value_ulong, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_LOCAL_MEM_SIZE: %d kB", i+1, device_value_ulong/(1024));
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &device_value_ulong, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_MEM_ALLOC_SIZE: %d MB", i+1, device_value_ulong/(1024*1024));
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &device_sizet, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_WORK_GROUP_SIZE: %d", i+1, device_sizet);
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uint), &device_value, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: %d", i+1, device_value);
-
-if(clGetDeviceInfo(clDevices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, dimsz, NULL) == CL_SUCCESS)
-g_message("#%d CL_DEVICE_MAX_WORK_ITEM_SIZES: %d x %d x %d", i+1, dimsz[0], dimsz[1], dimsz[2]);
-
-g_message(" ");
-
-
-}
-
-if(error != CL_SUCCESS)
-    g_message("OpenCL testing failed");
-    else
-    g_message("OpenCL testing succeded");
-
-    }*/
-
-static void
-ufo_filter_lamino_bp_generic_initialize (UfoFilterReduce *filter, UfoBuffer *input[], guint **output_dims, gfloat *default_value, GError **error)
-{
-    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(filter);
-    UfoResourceManager *manager = ufo_filter_get_resource_manager(UFO_FILTER(filter));
-    GError *tmp_error = NULL;
-    cl_context context;
-    guint width, height;
-
-    // TODO: how to solve 'include problem' for cl-files?
-    // resource_manager_add_program
-    //	 (UfoResourceManager *manager,
-    //	          const gchar *filename, const gchar *options, GError **error)
-
-    //    ufo_resource_manager_add_program(manager, "lamino_bp_generic.cl",
-    //  	 "-I /home/timurttv/_UFO-Project/framework/ufo-filters-am-installed/lib/ufo", &error);
-
-    priv->bp_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_bp_generic", &tmp_error);
-    priv->norm_vol_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_norm_vol", &tmp_error);
-
-    /* We don't need to clean the volume manually. All output buffers are
-     * initialized with `default_value`. */
-    /* priv->clean_vol_kernel =  ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_clean_vol", &tmp_error); */
-
-    if (tmp_error != NULL) {
-        g_propagate_error (error, tmp_error);
-        return;
-    }
-
-    ufo_buffer_get_2d_dimensions (input[0], &width, &height);
-    priv->params.proj_sx = (int) width;
-    priv->params.proj_sy = (int) height;
-
-    ////// init recon-parameters which are common for all projs
-    // shift volume origin to the box center
-    // 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;
-    context = (cl_context) ufo_resource_manager_get_context(manager);
-    priv->param_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(CLParameters), NULL, NULL);
-    priv->proj_idx = 0;
-
-    output_dims[0][0] = priv->params.vol_sx;
-    output_dims[0][1] = priv->params.vol_sy;
-    output_dims[0][2] = priv->params.vol_sz;
-
-    priv->global_work_size[0] = (size_t) priv->params.vol_sx;
-    priv->global_work_size[1] = (size_t) priv->params.vol_sy;
-    priv->global_work_size[2] = (size_t) priv->params.vol_sz;
-
-    *default_value = 0.0f;
-}
-
-static void
-ufo_filter_lamino_bp_generic_collect(UfoFilterReduce *filter, UfoBuffer *input[], UfoBuffer *output[], GError **error)
-{
-    /////////////////// testing OpenCL platform
-    // testing_cl_platform();
-
-    g_return_if_fail(UFO_IS_FILTER(filter));
-
-    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(filter);
-    CLParameters *params = &(priv->params);
-
-    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue (UFO_FILTER (filter));
-    cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output[0], command_queue);
-    const guint vol_num_dims = 3;
-
-    // clean volume before reconstruction on the GPU side
-    g_message("prepare the volume");
-    /* cl_event event; */
-    /* cl_kernel clean_vol_kernel = priv->clean_vol_kernel; */
-    /* 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)); */
-    // clFinish(command_queue);
-    // whait until kernel finishes its job
-    /* ufo_buffer_attach_event(output, event); */
-
-    // setup backprojection kernel
-    cl_kernel kernel = priv->bp_kernel;
-
-    // projection conter
-    // allocate memory for parameters
-    //
-    ////// init recon-parameters which are individual  for each projs
-    params->alpha = - 3 * G_PI/2 + params->theta;
-    params->phi   = params->angle_step* ((float) priv->proj_idx);
-
-    float sf = sin(params->phi),   cf = cos(params->phi), ct=cos(params->alpha),
-          st = sin(params->alpha), cg = cos(params->psi), sg=sin(params->psi);
-
-    params->mat_0 =  cg * cf - sg * st * sf;
-    params->mat_1 = -cg * sf - sg * st * cf;
-    params->mat_2 = -sg * ct;
-    params->mat_3 =  sg * cf + cg * st * sf;
-    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);
-
-    // send parameters to GPU
-    clEnqueueWriteBuffer(command_queue, priv->param_mem, CL_TRUE, 0, sizeof(CLParameters), params, 0, NULL, NULL);
-
-    // copy projection to GPU
-    cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input[0], 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 *) &priv->param_mem));
-    // call backprojection routine
-    g_message("processing of %d-th projection", priv->proj_idx);
-    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
-                vol_num_dims, NULL, priv->global_work_size, NULL, 0, NULL, NULL));
-    // clFinish(command_queue);
-
-    priv->proj_idx++;
-}
-
-static gboolean
-ufo_filter_lamino_bp_generic_reduce (UfoFilterReduce *filter, UfoBuffer *output[], GError **error)
-{
-    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(filter);
-
-    if (priv->finished)
-        return FALSE;
-
-    // normalize volume after reconstruction
-    cl_command_queue cmd_queue = ufo_filter_get_command_queue (UFO_FILTER (filter));
-    cl_kernel norm_vol_kernel = priv->norm_vol_kernel;
-    cl_mem output_mem = ufo_buffer_get_device_array (output[0], cmd_queue);
-    float stepPhi = priv->params.angle_step;
-    const guint vol_num_dims = 3;
-
-    CHECK_OPENCL_ERROR(clSetKernelArg(norm_vol_kernel, 0, sizeof(cl_mem), (void *) &output_mem));
-    CHECK_OPENCL_ERROR(clSetKernelArg(norm_vol_kernel, 1, sizeof(float), &stepPhi));
-
-    // call normalization kernel
-    g_message("volume post-processing");
-    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel((cl_command_queue) cmd_queue, norm_vol_kernel,
-                vol_num_dims, NULL, priv->global_work_size, NULL, 0, NULL, NULL));
-
-    priv->finished = TRUE;
-
-    return TRUE;
-}
-
-static void
-ufo_filter_lamino_bp_generic_finalize (GObject *object)
-{
-    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE (object);
-    CHECK_OPENCL_ERROR (clReleaseMemObject (priv->param_mem));
-    G_OBJECT_CLASS (ufo_filter_lamino_bp_generic_parent_class)->finalize (object);
-}
-
-static void
-ufo_filter_lamino_bp_generic_set_property(GObject *object, guint property_id, const GValue *value, GParamSpec *pspec)
-{
-    UfoFilterLaminoBPGeneric *self = UFO_FILTER_LAMINO_BP_GENERIC(object);
-    switch (property_id) {
-        case PROP_THETA:
-            self->priv->params.theta = (float) g_value_get_double(value);
-            break;
-        case PROP_PSI:
-            self->priv->params.psi = (float) g_value_get_double(value);
-            break;
-        case PROP_ANGLE_STEP:
-            self->priv->params.angle_step = (float) g_value_get_double(value);
-            break;
-        case PROP_VOL_SX:
-            self->priv->params.vol_sx = g_value_get_uint(value);
-            break;
-        case PROP_VOL_SY:
-            self->priv->params.vol_sy = g_value_get_uint(value);
-            break;
-        case PROP_VOL_SZ:
-            self->priv->params.vol_sz = g_value_get_uint(value);
-            break;
-        case PROP_VOL_OX:
-            self->priv->params.vol_ox = (float)g_value_get_double(value);
-            break;
-        case PROP_VOL_OY:
-            self->priv->params.vol_oy = (float)g_value_get_double(value);
-            break;
-        case PROP_VOL_OZ:
-            self->priv->params.vol_oz = (float)g_value_get_double(value);
-            break;
-        case PROP_PROJ_OX:
-            self->priv->params.proj_ox = (float)g_value_get_double(value);
-            break;
-        case PROP_PROJ_OY:
-            self->priv->params.proj_oy = (float)g_value_get_double(value);
-            break;
-        default:
-            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
-            break;
-    }
-}
-
-static void
-ufo_filter_lamino_bp_generic_get_property (GObject *object, guint property_id, GValue *value, GParamSpec*pspec)
-{
-    UfoFilterLaminoBPGeneric *self = UFO_FILTER_LAMINO_BP_GENERIC(object);
-    switch (property_id) {
-        case PROP_THETA:
-            g_value_set_double(value, (double) self->priv->params.theta);
-            break;
-        case PROP_PSI:
-            g_value_set_double(value, (double) self->priv->params.psi);
-            break;
-        case PROP_ANGLE_STEP:
-            g_value_set_double(value, (double) self->priv->params.angle_step);
-            break;
-        case PROP_VOL_SX:
-            g_value_set_uint(value, self->priv->params.vol_sx);
-            break;
-        case PROP_VOL_SY:
-            g_value_set_uint(value, self->priv->params.vol_sy);
-            break;
-        case PROP_VOL_SZ:
-            g_value_set_uint(value, self->priv->params.vol_sz);
-            break;
-        case PROP_VOL_OX:
-            g_value_set_double(value, (double)self->priv->params.vol_ox);
-            break;
-        case PROP_VOL_OY:
-            g_value_set_double(value, (double)self->priv->params.vol_oy);
-            break;
-        case PROP_VOL_OZ:
-            g_value_set_double(value, (double)self->priv->params.vol_oz);
-            break;
-        case PROP_PROJ_OX:
-            g_value_set_double(value, (double)self->priv->params.proj_ox);
-            break;
-        case PROP_PROJ_OY:
-            g_value_set_double(value, (double)self->priv->params.proj_oy);
-            break;
-        default:
-            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
-            break;
-    }
-}
-
-static void
-ufo_filter_lamino_bp_generic_class_init(UfoFilterLaminoBPGenericClass *klass)
-{
-    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
-    UfoFilterReduceClass *filter_class = UFO_FILTER_REDUCE_CLASS(klass);
-
-    gobject_class->set_property = ufo_filter_lamino_bp_generic_set_property;
-    gobject_class->get_property = ufo_filter_lamino_bp_generic_get_property;
-    gobject_class->finalize     = ufo_filter_lamino_bp_generic_finalize;
-    filter_class->initialize    = ufo_filter_lamino_bp_generic_initialize;
-    filter_class->collect       = ufo_filter_lamino_bp_generic_collect;
-    filter_class->reduce        = ufo_filter_lamino_bp_generic_reduce;
-
-    lamino_bp_generic_properties[PROP_THETA] =
-        g_param_spec_double("theta",
-                "Laminographic angle in radians",
-                "Laminographic angle in radians",
-                -4.0 * G_PI, +4.0 * G_PI, 0.0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_PSI] =
-        g_param_spec_double("psi",
-                "Axis misalignment angle in radians",
-                "Axis misalignment angle in radians",
-                -4.0 * G_PI, +4.0 * G_PI, 0.0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_ANGLE_STEP] =
-        g_param_spec_double("angle-step",
-                "Increment of rotation angle phi in radians",
-                "Increment of rotation angle phi in radians",
-                -4.0 * G_PI, +4.0 * G_PI, 0.0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_VOL_SX] =
-        g_param_spec_uint("vol-sx",
-                "Size of reconstructed volume along the 0X-axis in voxels",
-                "Size of reconstructed volume along the 0X-axis in voxels",
-                0, 1024*8, 512,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_VOL_SY] =
-        g_param_spec_uint("vol-sy",
-                "Size of reconstructed volume along the 0Y-axis in voxels",
-                "Size of reconstructed volume along the 0Y-axis in voxels",
-                0, 1024*8, 512,
-                G_PARAM_READWRITE);
-
-
-    lamino_bp_generic_properties[PROP_VOL_SZ] =
-        g_param_spec_uint("vol-sz",
-                "Size of reconstructed volume along the 0Z-axis in voxels",
-                "Size of reconstructed volume along the 0Z-axis in voxels",
-                0, 1024*8, 512,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_VOL_OX] =
-        g_param_spec_double("vol-ox",
-                "Volume origin offset from the center of a reco-box along the OX-axis in voxels",
-                "Volume origin offset from the center of a reco-box along the OX-axis in voxels",
-                -1024*8, 1024*8, 0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_VOL_OY] =
-        g_param_spec_double("vol-oy",
-                "Volume origin offset from the center of a reco-box along the OY-axis in voxels",
-                "Volume origin offset from the center of a reco-box along the OY-axis in voxels",
-                -1024*8, 1024*8, 0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_VOL_OZ] =
-        g_param_spec_double("vol-oz",
-                "Volume origin offset from the center of a reco-box along the OZ-axis in voxels",
-                "Volume origin offset from the center of a reco-box along the OZ-axis in voxels",
-                -1024*8, 1024*8, 0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_PROJ_OX] =
-        g_param_spec_double("proj-ox",
-                "Projection of the rotation center on the radiograph origin on the OX-axis",
-                "Projection of the rotation center on the radiograph origin on the OX-axis",
-                -1024*8, 1024*8, 0,
-                G_PARAM_READWRITE);
-
-    lamino_bp_generic_properties[PROP_PROJ_OY] =
-        g_param_spec_double("proj-oy",
-                "Projection of the rotation center on the radiograph origin on the OY-axis",
-                "Projection of the rotation center on the radiograph origin on the OY-axis",
-                -1024*8, 1024*8, 0,
-                G_PARAM_READWRITE);
-
-    g_object_class_install_property(gobject_class, PROP_THETA,    lamino_bp_generic_properties[PROP_THETA]);
-    g_object_class_install_property(gobject_class, PROP_PSI,      lamino_bp_generic_properties[PROP_PSI]);
-    g_object_class_install_property(gobject_class, PROP_ANGLE_STEP,   lamino_bp_generic_properties[PROP_ANGLE_STEP]);
-    g_object_class_install_property(gobject_class, PROP_VOL_SX,   lamino_bp_generic_properties[PROP_VOL_SX]);
-    g_object_class_install_property(gobject_class, PROP_VOL_SY,   lamino_bp_generic_properties[PROP_VOL_SY]);
-    g_object_class_install_property(gobject_class, PROP_VOL_SZ,   lamino_bp_generic_properties[PROP_VOL_SZ]);
-    g_object_class_install_property(gobject_class, PROP_VOL_OX,   lamino_bp_generic_properties[PROP_VOL_OX]);
-    g_object_class_install_property(gobject_class, PROP_VOL_OY,   lamino_bp_generic_properties[PROP_VOL_OY]);
-    g_object_class_install_property(gobject_class, PROP_VOL_OZ,   lamino_bp_generic_properties[PROP_VOL_OZ]);
-    g_object_class_install_property(gobject_class, PROP_PROJ_OX,  lamino_bp_generic_properties[PROP_PROJ_OX]);
-    g_object_class_install_property(gobject_class, PROP_PROJ_OY,  lamino_bp_generic_properties[PROP_PROJ_OY]);
-
-
-    /* install private data */
-    g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoBPGenericPrivate));
-}
-
-static void
-ufo_filter_lamino_bp_generic_init(UfoFilterLaminoBPGeneric *self)
-{
-    UfoInputParameter input_params[] = {{2, UFO_FILTER_INFINITE_INPUT}};
-    UfoOutputParameter output_params[] = {{3}};
-
-    // initialize parameters here
-    self->priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(self);
-    self->priv->finished = FALSE;
-
-    CLParameters * prms = &(self->priv->params);
-    prms->theta = 0.0;
-    prms->psi   = 0.0;
-    prms->angle_step = 0.0;
-
-    prms->vol_sx = 512;
-    prms->vol_sy = 512;
-    prms->vol_sz = 512;
-
-    prms->vol_ox = 0.0;
-    prms->vol_oy = 0.0;
-    prms->vol_oz = 0.0;
-
-    prms->proj_ox = 0.0;
-    prms->proj_oy = 0.0;
-
-    ufo_filter_register_inputs (UFO_FILTER (self), 1, input_params);
-    ufo_filter_register_outputs (UFO_FILTER (self), 1, output_params);
-}
-
-G_MODULE_EXPORT UfoFilter *
-ufo_filter_plugin_new(void)
-{
-    return g_object_new(UFO_TYPE_FILTER_LAMINO_BP_GENERIC, NULL);
-}
-

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

@@ -1,37 +0,0 @@
-#ifndef __UFO_FILTER_LAMINO_BP_GENERIC_H__
-#define __UFO_FILTER_LAMINO_BP_GENERIC_H__
-
-#include <glib.h>
-#include <ufo/ufo-filter-reduce.h>
-
-#define UFO_TYPE_FILTER_LAMINO_BP_GENERIC 	(ufo_filter_lamino_bp_generic_get_type())
-#define UFO_FILTER_LAMINO_BP_GENERIC(obj)	(G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_LAMINO_BP_GENERIC, UfoFilterLaminoBPGeneric))	
-#define UFO_IS_FILTER_LAMINO_BP_GENERIC(obj)	(G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_LAMINO_BP_GENERIC))
-#define UFO_FILTER_LAMINO_BP_GENERIC_CLASS(klass) 	(G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_LAMINO_BP_GENERIC, UfoFilterLaminoBPGenericClass))
-#define UFO_IS_FILTER_LAMINO_BP_GENERIC_CLASS(klass)	(G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_LAMINO_BP_GENERIC))
-#define UFO_FILTER_LAMINO_BP_GENERI_GET_CLASS(obj)   	(G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_LAMINO_BP_GENERIC, UfoFilterLaminoBPGenericClass))
-
-typedef struct _UfoFilterLaminoBPGeneric           UfoFilterLaminoBPGeneric;
-typedef struct _UfoFilterLaminoBPGenericClass      UfoFilterLaminoBPGenericClass;
-typedef struct _UfoFilterLaminoBPGenericPrivate    UfoFilterLaminoBPGenericPrivate;
-
-struct _UfoFilterLaminoBPGeneric {
-	   /*< private >*/
-	   UfoFilterReduce parent_instance;
-	   UfoFilterLaminoBPGenericPrivate *priv;
-};
-
-/*
- * UfoFilterLaminoBPGenericClass:
- *
- * #UfoFilterLaminoBPGenericClass class
- */
-struct _UfoFilterLaminoBPGenericClass {
-	   /*< private >*/
-	   UfoFilterReduceClass parent_class;
-};
-
-GType ufo_filter_lamino_bp_generic_get_type(void);
-
-
-#endif //__UFO_FILTER_LAMINO_BP_GENERIC_H__

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

@@ -1,122 +0,0 @@
-#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-ft-conv
- * @Short_description:
- * @Title: laminoftconv
- *
- * Detailed description.
- */
-
-struct _UfoFilterLaminoFTConvPrivate {
-    cl_kernel   kernel;
-    size_t      global_work_size[2];
-    guint       img_width;
-    guint       img_height;
-};
-
-enum {
-    INPUT_FILTER,
-    INPUT_IMAGE
-};
-
-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))
-
-static void
-ufo_filter_lamino_ft_conv_initialize (UfoFilter *filter, UfoBuffer *params[], guint **dim_sizes, GError **error)
-{
-    UfoFilterLaminoFTConvPrivate *priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE (filter);
-    UfoResourceManager *manager = ufo_filter_get_resource_manager (filter);
-    guint f_width, f_height;
-    GError *tmp_error = NULL;
-
-    priv->kernel = ufo_resource_manager_get_kernel (manager, "lamino_ft_conv.cl", "lamino_c", &tmp_error);
-
-    if (tmp_error != NULL) {
-        g_propagate_error (error, tmp_error);
-        return;
-    }
-
-    ufo_buffer_get_2d_dimensions (params[0], &priv->img_width, &priv->img_height);
-    ufo_buffer_get_2d_dimensions (params[1], &f_width, &f_height);
-
-    if ((priv->img_width != f_width ) || (priv->img_height != f_height)) {
-        g_set_error (error, UFO_FILTER_ERROR, UFO_FILTER_ERROR_INITIALIZATION,
-                "Filter and image sizes are different");
-        return;
-    }
-
-    dim_sizes[0][0] = priv->img_width;
-    dim_sizes[0][1] = priv->img_height;
-    priv->global_work_size[0] = (size_t) priv->img_width;
-    priv->global_work_size[1] = (size_t) priv->img_height;
-}
-
-static void
-ufo_filter_lamino_ft_conv_process_gpu(UfoFilter *flt, UfoBuffer *input[], UfoBuffer *output[], GError **error)
-{
-    g_return_if_fail(UFO_IS_FILTER(flt));
-
-    UfoFilterLaminoFTConvPrivate *priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE (flt);
-    UfoProfiler *profiler = ufo_filter_get_profiler(flt);
-
-    cl_command_queue command_queue = ufo_filter_get_command_queue (flt);
-    cl_kernel kernel = priv->kernel;
-    cl_mem filter_mem = (cl_mem) ufo_buffer_get_device_array(input[INPUT_FILTER], command_queue);
-    cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input[INPUT_IMAGE],  command_queue);
-    cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output[0], 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), &priv->img_width));
-
-    ufo_profiler_call (profiler, command_queue, kernel, 2, priv->global_work_size, NULL);
-
-    clFinish(command_queue);
-}
-
-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);
-
-    filter_class->initialize = ufo_filter_lamino_ft_conv_initialize;
-    filter_class->process_gpu = ufo_filter_lamino_ft_conv_process_gpu;
-
-    g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoFTConvPrivate));
-}
-
-static void
-ufo_filter_lamino_ft_conv_init(UfoFilterLaminoFTConv *self)
-{
-    UfoOutputParameter output_params[] = {{2}};
-    UfoInputParameter input_params[] = {
-        {2, 1},
-        {2, UFO_FILTER_INFINITE_INPUT}};
-
-    UfoFilterLaminoFTConvPrivate *priv = self->priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(self);
-    priv->kernel = NULL;
-
-    ufo_filter_register_inputs (UFO_FILTER (self), 2, input_params);
-    ufo_filter_register_outputs (UFO_FILTER (self), 1, output_params);
-}
-
-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

@@ -1,40 +0,0 @@
-#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

@@ -1,252 +0,0 @@
-#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-source.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 fill_width;
-    guint height;
-    // laminographic angle
-    float theta;
-    // pixel resolution (in um)
-    float tau;
-    cl_kernel kernel_cf;
-    gboolean done;
-};
-
-G_DEFINE_TYPE(UfoFilterLaminoRamp, ufo_filter_lamino_ramp, UFO_TYPE_FILTER_SOURCE)
-
-#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_FILL_WIDTH,
-    PROP_HEIGHT,
-    PROP_THETA,
-    PROP_TAU,
-    N_PROPERTIES
-};
-
-static GParamSpec *lamino_ramp_properties[N_PROPERTIES] = { NULL, };
-
-static int
-is_power_of_two (guint x)
-{
-    return ((x != 0) && !(x & (x - 1)));
-}
-
-static void
-ufo_filter_lamino_ramp_initialize (UfoFilterSource *filter, guint **dim_sizes, GError **error)
-{
-    UfoFilterLaminoRampPrivate *priv = UFO_FILTER_LAMINO_RAMP_GET_PRIVATE(filter);
-    UfoResourceManager *manager = ufo_filter_get_resource_manager(UFO_FILTER(filter));
-    GError *tmp_error = NULL;
-
-    priv->kernel_cf = ufo_resource_manager_get_kernel(manager, "lamino_ramp.cl", "lamino_ramp_create_filter", &tmp_error);
-
-    if (tmp_error != NULL) {
-        g_propagate_error (error, tmp_error);
-        return;
-    }
-
-    if (!is_power_of_two (priv->width)) {
-        g_set_error (error, UFO_FILTER_ERROR, UFO_FILTER_ERROR_INITIALIZATION,
-                "Filter width `%i` is not a power of two", priv->width);
-        return;
-    }
-
-    if (!is_power_of_two (priv->height)) {
-        g_set_error (error, UFO_FILTER_ERROR, UFO_FILTER_ERROR_INITIALIZATION,
-                "Filter height `%i` is not a power of two", priv->height);
-        return;
-    }
-
-    dim_sizes[0][0] = priv->width;
-    dim_sizes[0][1] = priv->height;
-}
-
-static gboolean
-ufo_filter_lamino_ramp_generate (UfoFilterSource *filter, UfoBuffer *output[], GError **error)
-{
-    g_return_val_if_fail (UFO_IS_FILTER (filter), FALSE);
-    UfoFilterLaminoRampPrivate *priv = UFO_FILTER_LAMINO_RAMP_GET_PRIVATE (filter);
-    cl_command_queue cmd_queue = ufo_filter_get_command_queue (UFO_FILTER (filter));
-
-    if (priv->done)
-        return FALSE;
-
-    guint width  = priv->width;
-    guint fwidth = priv->fill_width;
-    guint height = priv->height;
-    float theta = priv->theta;
-    float tau   = priv->tau;
-    size_t global_work_size[2] = { (size_t) priv->width, (size_t) priv->height };
-
-    cl_kernel kernel = priv->kernel_cf;
-    cl_mem output_mem  = (cl_mem) ufo_buffer_get_device_array(output[0], cmd_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(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));
-
-    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(cmd_queue, kernel,
-                2, NULL, global_work_size, NULL,
-                0, NULL, NULL));
-
-    priv->done = TRUE;
-
-    return TRUE;
-}
-
-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);
-
-    switch (property_id) {
-        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;
-        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_FILL_WIDTH:
-            g_value_set_uint(value, self->priv->fill_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);
-    UfoFilterSourceClass *filter_class = UFO_FILTER_SOURCE_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->generate = ufo_filter_lamino_ramp_generate;
-
-    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_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",
-                "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_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]);
-
-    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);
-    UfoOutputParameter output_params[] = {{2}};
-
-    priv->width = 4;
-    priv->fill_width=2;
-    priv->height = 1;
-    priv->theta = 0.0;
-    priv->tau = 10.0;
-    priv->done = FALSE;
-
-    ufo_filter_register_outputs (UFO_FILTER (self), 1, output_params);
-}
-
-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

@@ -1,39 +0,0 @@
-#ifndef __UFO_FILTER_LAMINO_RAMP_H
-#define __UFO_FILTER_LAMINO_RAMP_H
-
-#include <glib.h>
-#include <glib-object.h>
-#include <ufo/ufo-filter-source.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 >*/
-    UfoFilterSource parent_instance;
-
-    UfoFilterLaminoRampPrivate *priv;
-};
-
-/**
- * UfoFilterLaminoRampClass:
- *
- * #UfoFilterLaminoRamp class
- */
-struct _UfoFilterLaminoRampClass {
-    /*< private >*/
-    UfoFilterSourceClass parent_class;
-};
-
-GType ufo_filter_lamino_ramp_get_type(void);
-UfoFilter *ufo_filter_plugin_new(void);
-
-#endif

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

@@ -1,364 +0,0 @@
-#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
-{
-    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)
-
-#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, UfoBuffer *params[], guint **dim_sizes, GError **error)
-{
-    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(filter);
-    UfoResourceManager *manager = ufo_filter_get_resource_manager(filter);
-    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);
-
-    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_gpu(UfoFilter *filter, UfoBuffer *input[], UfoBuffer *output[], GError **error)
-{
-    g_return_if_fail (UFO_IS_FILTER (filter));
-    cl_command_queue command_queue = ufo_filter_get_command_queue (filter);
-    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(filter);
-
-    const PaddingMode mode = priv->mode;
-    const guint pxl = priv->xl;
-    const guint pyt = priv->yt;
-    float pval = priv->pconst;
-
-    const guint ixs = priv->in_width;
-    const guint iys = priv->in_height;
-    const guint oxs = priv->out_width;
-
-    // init padding value  for different modes
-    if (mode == PADDING_ZERO)
-        pval = 0.0;
-
-    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;
-
-        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)) {
-        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);
-    }
-
-    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);
-    }
-
-    clFinish(command_queue); // synchro?
-}
-
-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_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);
-
-    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);
-    UfoInputParameter input_params[] = {{2, UFO_FILTER_INFINITE_INPUT}};
-    UfoOutputParameter output_params[] = {{2}};
-
-    // 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_inputs (UFO_FILTER (self), 1, input_params);
-    ufo_filter_register_outputs (UFO_FILTER (self), 1, output_params);
-}
-
-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

@@ -1,40 +0,0 @@
-#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.c

@@ -1,172 +0,0 @@
-#include <gmodule.h>
-#ifdef __APPLE__
-#include <OpenCL/cl.h>
-#else
-#include <CL/cl.h>
-#endif
-
-#include <ufo/ufo-filter.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;
-
-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, UfoBuffer *params[], guint **dim_sizes, GError **error)
-{
-    UfoFilterScalePrivate *priv = UFO_FILTER_SCALE_GET_PRIVATE(filter);
-    UfoResourceManager *manager = ufo_filter_get_resource_manager(filter);
-    GError *tmp_error = NULL;
-
-    priv->kernel = ufo_resource_manager_get_kernel(manager, "scale.cl", "scale", &tmp_error);
-    if (tmp_error != NULL) {
-	g_propagate_error (error, tmp_error);
-        return;
-    }
-   
-    guint xs, ys; 
-    ufo_buffer_get_2d_dimensions (params[0], &xs, &ys);
-    dim_sizes[0][0] = xs;
-    dim_sizes[0][1] = ys;
-}
-
-/*
- * This is the main method in which the filter processes one buffer after
- * another.
- */
-static void ufo_filter_scale_process_gpu(UfoFilter *filter, UfoBuffer *input[], UfoBuffer *output[], GError **error)
-{
-    g_return_if_fail(UFO_IS_FILTER(filter));
-    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
-    UfoFilterScalePrivate *priv = UFO_FILTER_SCALE_GET_PRIVATE(filter);
-
-    guint xs, ys;
-    ufo_buffer_get_2d_dimensions (input[0], &xs, &ys);
-    size_t global_work_size[2] = {(size_t) xs, (size_t) ys};
-
-
-     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);
-
-     float scale = (float) priv->scale;
-     cl_kernel kernel = priv->kernel;
-
-     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(int),   &xs));
-     CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(float), &scale));
-
-     cl_event event;	 
-     CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
-					      2, NULL, global_work_size, NULL,
-					      0, NULL, &event));
-     clFinish(command_queue);	 
-}
-
-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_gpu = ufo_filter_scale_process_gpu;
-
-    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);
-    UfoInputParameter input_params[] = {{2, UFO_FILTER_INFINITE_INPUT}};
-    UfoOutputParameter output_params[] = {{2}};
-
-    priv->scale = 1.0;
-    priv->kernel = NULL;
-    
-    ufo_filter_register_inputs (UFO_FILTER(self), 1, input_params);
-    ufo_filter_register_outputs(UFO_FILTER(self), 1, output_params);
-}
-
-G_MODULE_EXPORT UfoFilter  *ufo_filter_plugin_new(void) 
-{
-    return g_object_new(UFO_TYPE_FILTER_SCALE, NULL);
-}

src/ufo-filter-scale.h

@@ -1,38 +0,0 @@
-#ifndef __UFO_FILTER_SCALE_H
-#define __UFO_FILTER_SCALE_H
-
-#include <glib.h>
-#include <glib-object.h>
-
-#include <ufo/ufo-filter.h>
-
-#define UFO_TYPE_FILTER_SCALE             (ufo_filter_scale_get_type())
-#define UFO_FILTER_SCALE(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_SCALE, UfoFilterScale))
-#define UFO_IS_FILTER_SCALE(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_SCALE))
-#define UFO_FILTER_SCALE_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_SCALE, UfoFilterScaleClass))
-#define UFO_IS_FILTER_SCALE_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_SCALE))
-#define UFO_FILTER_SCALE_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_SCALE, UfoFilterScaleClass))
-
-typedef struct _UfoFilterScale           UfoFilterScale;
-typedef struct _UfoFilterScaleClass      UfoFilterScaleClass;
-typedef struct _UfoFilterScalePrivate    UfoFilterScalePrivate;
-
-struct _UfoFilterScale {
-   /*< private >*/
-   UfoFilter parent_instance;
-
-   UfoFilterScalePrivate *priv;
-};
-
-/*
- * UfoFilterScaleClass:
- *
- * #UfoFilterScaleClass class
- */
-struct _UfoFilterScaleClass {
-   /*< private >*/
-   UfoFilterClass parent_class;
-};
-
-GType ufo_filter_scale_get_type(void);
-#endif

src/ufo-filter-scale.host

@@ -1,212 +0,0 @@
-#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);
-}

src/ufo-lamino-bp-task.c

@@ -109,13 +109,14 @@ ufo_lamino_bp_task_get_requisition (UfoTask *task,
 static void
 ufo_lamino_bp_task_get_structure (UfoTask *task,
                                   guint *n_inputs,
-                                  guint **n_dims,
+                                  UfoInputParam **input_params,
                                   UfoTaskMode *mode)
 {
     *mode = UFO_TASK_MODE_REDUCE;
     *n_inputs = 1;
-    *n_dims = g_new0 (guint, 1);
-    (*n_dims)[0] = 2;
+    *input_params = g_new0 (UfoInputParam, 1);
+    (*input_params)[0].n_dims = 2;
+    (*input_params)[0].n_expected = -1;
 }
 
 static gboolean

src/ufo-lamino-conv-task.c

@@ -0,0 +1,207 @@
+/**
+ * SECTION:ufo-filter-task
+ * @Short_description: Process arbitrary Filter kernels
+ * @Title: filter
+ *
+ * This module is used to load an arbitrary #UfoLaminoConvTask:kernel from
+ * #UfoLaminoConvTask:filename and execute it on each input. The kernel must have
+ * only two global float array parameters, the first represents the input, the
+ * second one the output. #UfoLaminoConvTask:num-dims must be changed, if the kernel
+ * accesses either one or three dimensional index spaces.
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+#include <ufo-gpu-task-iface.h>
+#include "ufo-lamino-conv-task.h"
+
+struct _UfoLaminoConvTaskPrivate {
+    cl_kernel kernel;
+    guint img_width;
+    guint img_height;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+static void ufo_gpu_task_interface_init (UfoGpuTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoLaminoConvTask, ufo_lamino_conv_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init)
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_GPU_TASK,
+                                                ufo_gpu_task_interface_init))
+
+#define UFO_LAMINO_CONV_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_LAMINO_CONV_TASK, UfoLaminoConvTaskPrivate))
+
+UfoNode *
+ufo_lamino_conv_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_LAMINO_CONV_TASK, NULL));
+}
+
+static gboolean
+ufo_lamino_conv_task_process (UfoGpuTask *task,
+                              UfoBuffer **inputs,
+                              UfoBuffer *output,
+                              UfoRequisition *requisition,
+                              UfoGpuNode *node)
+{
+    UfoLaminoConvTaskPrivate *priv;
+    cl_command_queue cmd_queue;
+    cl_mem in1_mem;
+    cl_mem in2_mem;
+    cl_mem out_mem;
+
+    priv = UFO_LAMINO_CONV_TASK (task)->priv;
+
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+    in1_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
+    in2_mem = ufo_buffer_get_device_array (inputs[1], cmd_queue);
+    out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof(cl_mem), (void *) &in1_mem));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof(cl_mem), (void *) &in2_mem));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof(cl_mem), (void *) &out_mem));
+
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, priv->kernel,
+                                                       2, NULL, requisition->dims, NULL,
+                                                       0, NULL, NULL));
+
+    return TRUE;
+}
+
+static void
+ufo_lamino_conv_task_setup (UfoTask *task,
+                           UfoResources *resources,
+                           GError **error)
+{
+    UfoLaminoConvTaskPrivate *priv;
+
+    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (task);
+
+    priv->kernel = ufo_resources_get_kernel (resources, "lamino_ft_conv.cl", "lamino_c", error);
+
+    if (priv->kernel)
+        clRetainKernel (priv->kernel);
+}
+
+static void
+ufo_lamino_conv_task_get_requisition (UfoTask *task,
+                                      UfoBuffer **inputs,
+                                      UfoRequisition *requisition)
+{
+    ufo_buffer_get_requisition (inputs[0], requisition);
+}
+
+static void
+ufo_lamino_conv_task_get_structure (UfoTask *task,
+                                    guint *n_inputs,
+                                    UfoInputParam **in_params,
+                                    UfoTaskMode *mode)
+{
+    UfoLaminoConvTaskPrivate *priv;
+
+    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (task);
+    *mode = UFO_TASK_MODE_SINGLE;
+    *n_inputs = 2;
+    *in_params = g_new0 (UfoInputParam, 2);
+    (*in_params)[0].n_dims = 2;
+    (*in_params)[0].n_expected = -1;
+    (*in_params)[1].n_dims = 2;
+    (*in_params)[1].n_expected = 1;
+}
+
+static UfoNode *
+ufo_lamino_conv_task_copy_real (UfoNode *node,
+                                GError **error)
+{
+    UfoLaminoConvTask *copy;
+    copy = UFO_LAMINO_CONV_TASK (ufo_lamino_conv_task_new ());
+    return UFO_NODE (copy);
+}
+
+static gboolean
+ufo_lamino_conv_task_equal_real (UfoNode *n1,
+                            UfoNode *n2)
+{
+    g_return_val_if_fail (UFO_IS_LAMINO_CONV_TASK (n1) && UFO_IS_LAMINO_CONV_TASK (n2), FALSE);
+    return TRUE;
+}
+
+static void
+ufo_lamino_conv_task_finalize (GObject *object)
+{
+    UfoLaminoConvTaskPrivate *priv;
+
+    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (object);
+
+    G_OBJECT_CLASS (ufo_lamino_conv_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_lamino_conv_task_setup;
+    iface->get_requisition = ufo_lamino_conv_task_get_requisition;
+    iface->get_structure = ufo_lamino_conv_task_get_structure;
+}
+
+static void
+ufo_gpu_task_interface_init (UfoGpuTaskIface *iface)
+{
+    iface->process = ufo_lamino_conv_task_process;
+}
+
+static void
+ufo_lamino_conv_task_set_property (GObject *object,
+                                  guint property_id,
+                                  const GValue *value,
+                                  GParamSpec *pspec)
+{
+    switch (property_id) {
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_lamino_conv_task_get_property (GObject *object,
+                              guint property_id,
+                              GValue *value,
+                              GParamSpec *pspec)
+{
+    switch (property_id) {
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_lamino_conv_task_class_init (UfoLaminoConvTaskClass *klass)
+{
+    GObjectClass *oclass;
+    UfoNodeClass *node_class;
+    
+    oclass = G_OBJECT_CLASS (klass);
+    node_class = UFO_NODE_CLASS (klass);
+
+    oclass->finalize = ufo_lamino_conv_task_finalize;
+    oclass->set_property = ufo_lamino_conv_task_set_property;
+    oclass->get_property = ufo_lamino_conv_task_get_property;
+
+    node_class->copy = ufo_lamino_conv_task_copy_real;
+    node_class->equal = ufo_lamino_conv_task_equal_real;
+
+    g_type_class_add_private(klass, sizeof(UfoLaminoConvTaskPrivate));
+}
+
+static void
+ufo_lamino_conv_task_init (UfoLaminoConvTask *self)
+{
+    UfoLaminoConvTaskPrivate *priv;
+    self->priv = priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (self);
+}

src/ufo-lamino-conv-task.h

@@ -0,0 +1,47 @@
+#ifndef __UFO_LAMINO_CONV_TASK_H
+#define __UFO_LAMINO_CONV_TASK_H
+
+#include <ufo-task-node.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_LAMINO_CONV_TASK             (ufo_lamino_conv_task_get_type())
+#define UFO_LAMINO_CONV_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_LAMINO_CONV_TASK, UfoLaminoConvTask))
+#define UFO_IS_LAMINO_CONV_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_LAMINO_CONV_TASK))
+#define UFO_LAMINO_CONV_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_LAMINO_CONV_TASK, UfoLaminoConvTaskClass))
+#define UFO_IS_LAMINO_CONV_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_LAMINO_CONV_TASK))
+#define UFO_LAMINO_CONV_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_LAMINO_CONV_TASK, UfoLaminoConvTaskClass))
+
+typedef struct _UfoLaminoConvTask           UfoLaminoConvTask;
+typedef struct _UfoLaminoConvTaskClass      UfoLaminoConvTaskClass;
+typedef struct _UfoLaminoConvTaskPrivate    UfoLaminoConvTaskPrivate;
+
+/**
+ * UfoLaminoConvTask:
+ *
+ * Main object for organizing filters. The contents of the #UfoLaminoConvTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoLaminoConvTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoLaminoConvTaskPrivate *priv;
+};
+
+/**
+ * UfoLaminoConvTaskClass:
+ *
+ * #UfoLaminoConvTask class
+ */
+struct _UfoLaminoConvTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_lamino_conv_task_new       (void);
+GType     ufo_lamino_conv_task_get_type  (void);
+
+G_END_DECLS
+
+#endif

src/ufo-lamino-ramp-task.c

@@ -0,0 +1,313 @@
+/**
+ * SECTION:ufo-reader-task
+ * @Short_description: Read TIFF and EDF files
+ * @Title: reader
+ *
+ * The reader node loads single files from disk and provides them as a stream
+ * The nominal resolution can be decreased by specifying the #UfoLaminoRampTask:x
+ * and #UfoLaminoRampTask:y coordinates, and the #UfoLaminoRampTask:width and
+ * #UfoLaminoRampTask:height of a region of interest.
+ */
+
+#include <gmodule.h>
+#include <stdlib.h>
+#include <string.h>
+#include <tiffio.h>
+#include <glob.h>
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <ufo-gpu-task-iface.h>
+#include "ufo-lamino-ramp-task.h"
+
+struct _UfoLaminoRampTaskPrivate {
+    guint width;
+    guint height;
+    guint fill_width;
+    gfloat theta;
+    gfloat tau;
+    cl_kernel kernel;
+    gboolean done;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+static void ufo_gpu_task_interface_init (UfoGpuTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoLaminoRampTask, ufo_lamino_ramp_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init)
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_GPU_TASK,
+                                                ufo_gpu_task_interface_init))
+
+#define UFO_LAMINO_RAMP_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_LAMINO_RAMP_TASK, UfoLaminoRampTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_WIDTH,
+    PROP_FILL_WIDTH,
+    PROP_HEIGHT,
+    PROP_THETA,
+    PROP_TAU,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_lamino_ramp_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_LAMINO_RAMP_TASK, NULL));
+}
+
+static int
+is_power_of_two (guint x)
+{
+    return ((x != 0) && !(x & (x - 1)));
+}
+
+static void
+ufo_lamino_ramp_task_setup (UfoTask *task,
+                            UfoResources *resources,
+                            GError **error)
+{
+    UfoLaminoRampTask *node;
+    UfoLaminoRampTaskPrivate *priv;
+
+    node = UFO_LAMINO_RAMP_TASK (task);
+    priv = node->priv;
+
+    if (!is_power_of_two (priv->width)) {
+        g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
+                     "Filter width `%i' is not a power of two", priv->width);
+        return;
+    }
+
+    if (!is_power_of_two (priv->height)) {
+        g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
+                     "Filter height `%i' is not a power of two", priv->height);
+        return;
+    }
+
+    priv->kernel = ufo_resources_get_kernel (resources,
+                                             "lamino_ramp.cl",
+                                             "lamino_ramp_create_filter",
+                                             error);
+
+    if (priv->kernel != NULL)
+        clRetainKernel (priv->kernel);
+}
+
+static void
+ufo_lamino_ramp_task_get_requisition (UfoTask *task,
+                                      UfoBuffer **inputs,
+                                      UfoRequisition *requisition)
+{
+    UfoLaminoRampTaskPrivate *priv;
+
+    priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (UFO_LAMINO_RAMP_TASK (task));
+
+    requisition->n_dims = 2;
+    requisition->dims[0] = priv->width;
+    requisition->dims[1] = priv->height;
+}
+
+static void
+ufo_lamino_ramp_task_get_structure (UfoTask *task,
+                                    guint *n_inputs,
+                                    UfoInputParam **in_params,
+                                    UfoTaskMode *mode)
+{
+    *n_inputs = 0;
+    *mode = UFO_TASK_MODE_SINGLE;
+}
+
+static gboolean
+ufo_lamino_ramp_task_process (UfoGpuTask *task,
+                              UfoBuffer **inputs,
+                              UfoBuffer *output,
+                              UfoRequisition *requisition,
+                              UfoGpuNode *node)
+{
+    UfoLaminoRampTaskPrivate *priv;
+    cl_command_queue cmd_queue;
+    cl_mem out_mem;
+
+    priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (UFO_LAMINO_RAMP_TASK (task));
+
+    if (priv->done)
+        return FALSE;
+
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+    out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof(cl_mem), (void *) &out_mem));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof(int), &priv->width));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof(int), &priv->fill_width));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 3, sizeof(int), &priv->height));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 4, sizeof(float), &priv->theta));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 5, sizeof(float), &priv->tau));
+
+    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, priv->kernel,
+                                                       2, NULL, requisition->dims, NULL,
+                                                       0, NULL, NULL));
+
+    priv->done = TRUE;
+    return TRUE;
+}
+
+static void
+ufo_lamino_ramp_task_set_property (GObject *object,
+                                   guint property_id,
+                                   const GValue *value,
+                                   GParamSpec *pspec)
+{
+    UfoLaminoRampTaskPrivate *priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_WIDTH:
+            priv->width = g_value_get_uint (value);
+            break;
+        case PROP_FILL_WIDTH:
+            priv->fill_width = g_value_get_uint (value);
+            break;
+        case PROP_HEIGHT:
+            priv->height = g_value_get_uint (value);
+            break;
+        case PROP_THETA:
+            priv->theta = (gfloat) g_value_get_double (value);
+            break;
+        case PROP_TAU:
+            priv->tau = (gfloat) g_value_get_double (value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_lamino_ramp_task_get_property (GObject *object,
+                                   guint property_id,
+                                   GValue *value,
+                                   GParamSpec *pspec)
+{
+    UfoLaminoRampTaskPrivate *priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_WIDTH:
+            g_value_set_uint (value, priv->width);
+            break;
+        case PROP_FILL_WIDTH:
+            g_value_set_uint (value, priv->fill_width);
+            break;
+        case PROP_HEIGHT:
+            g_value_set_uint (value, priv->height);
+            break;
+        case PROP_THETA:
+            g_value_set_double (value, priv->theta);
+            break;
+        case PROP_TAU:
+            g_value_set_double (value, priv->tau);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_lamino_ramp_task_finalize (GObject *object)
+{
+    UfoLaminoRampTaskPrivate *priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (object);
+
+    if (priv->kernel != NULL) {
+        clReleaseKernel (priv->kernel);
+        priv->kernel = NULL;
+    }
+
+    G_OBJECT_CLASS (ufo_lamino_ramp_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_lamino_ramp_task_setup;
+    iface->get_structure = ufo_lamino_ramp_task_get_structure;
+    iface->get_requisition = ufo_lamino_ramp_task_get_requisition;
+}
+
+static void
+ufo_gpu_task_interface_init (UfoGpuTaskIface *iface)
+{
+    iface->process = ufo_lamino_ramp_task_process;
+}
+
+static void
+ufo_lamino_ramp_task_class_init (UfoLaminoRampTaskClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
+
+    gobject_class->set_property = ufo_lamino_ramp_task_set_property;
+    gobject_class->get_property = ufo_lamino_ramp_task_get_property;
+    gobject_class->finalize = ufo_lamino_ramp_task_finalize;
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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(UfoLaminoRampTaskPrivate));
+}
+
+static void
+ufo_lamino_ramp_task_init(UfoLaminoRampTask *self)
+{
+    UfoLaminoRampTaskPrivate *priv = NULL;
+
+    self->priv = priv = UFO_LAMINO_RAMP_TASK_GET_PRIVATE (self);
+    priv->width = 4;
+    priv->fill_width=2;
+    priv->height = 1;
+    priv->theta = 0.0;
+    priv->tau = 10.0;
+    priv->kernel = NULL;
+    priv->done = FALSE;
+}

src/ufo-lamino-ramp-task.h

@@ -0,0 +1,47 @@
+#ifndef __UFO_LAMINO_RAMP_TASK_H
+#define __UFO_LAMINO_RAMP_TASK_H
+
+#include <ufo-task-node.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_LAMINO_RAMP_TASK             (ufo_lamino_ramp_task_get_type())
+#define UFO_LAMINO_RAMP_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_LAMINO_RAMP_TASK, UfoLaminoRampTask))
+#define UFO_IS_LAMINO_RAMP_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_LAMINO_RAMP_TASK))
+#define UFO_LAMINO_RAMP_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_LAMINO_RAMP_TASK, UfoLaminoRampTaskClass))
+#define UFO_IS_LAMINO_RAMP_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_LAMINO_RAMP_TASK))
+#define UFO_LAMINO_RAMP_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_LAMINO_RAMP_TASK, UfoLaminoRampTaskClass))
+
+typedef struct _UfoLaminoRampTask           UfoLaminoRampTask;
+typedef struct _UfoLaminoRampTaskClass      UfoLaminoRampTaskClass;
+typedef struct _UfoLaminoRampTaskPrivate    UfoLaminoRampTaskPrivate;
+
+/**
+ * UfoLaminoRampTask:
+ *
+ * Main object for organizing filters. The contents of the #UfoLaminoRampTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoLaminoRampTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoLaminoRampTaskPrivate *priv;
+};
+
+/**
+ * UfoLaminoRampTaskClass:
+ *
+ * #UfoLaminoRampTask class
+ */
+struct _UfoLaminoRampTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_lamino_ramp_task_new       (void);
+GType     ufo_lamino_ramp_task_get_type  (void);
+
+G_END_DECLS
+
+#endif

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

@@ -202,7 +202,7 @@ ufo_padding_2d_task_get_requisition (UfoTask *task,
 static void
 ufo_padding_2d_task_get_structure (UfoTask *task,
                                    guint *n_inputs,
-                                   guint **n_dims,
+                                   UfoInputParam **in_params,
                                    UfoTaskMode *mode)
 {
     UfoPadding2DTaskPrivate *priv;
@@ -210,8 +210,9 @@ ufo_padding_2d_task_get_structure (UfoTask *task,
     priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task);
     *mode = UFO_TASK_MODE_SINGLE;
     *n_inputs = 1;
-    *n_dims = g_new0 (guint, 1);
-    (*n_dims)[0] = 2;
+    *in_params = g_new0 (UfoInputParam, 1);
+    (*in_params)[0].n_dims = 2;
+    (*in_params)[0].n_expected = -1;
 }
 
 static UfoNode *

src/ufo-scale-task.c

@@ -79,17 +79,18 @@ ufo_scale_task_process (UfoGpuTask *task,
 
 static void
 ufo_scale_task_get_structure (UfoTask *task,
-                               guint *n_inputs,
-                               guint **n_dims,
-                               UfoTaskMode *mode)
+                              guint *n_inputs,
+                              UfoInputParam **in_params,
+                              UfoTaskMode *mode)
 {
     UfoScaleTaskPrivate *priv;
 
     priv = UFO_SCALE_TASK_GET_PRIVATE (task);
     *mode = UFO_TASK_MODE_SINGLE;
     *n_inputs = 1;
-    *n_dims = g_new0 (guint, 1);
-    (*n_dims)[0] = 2;
+    *in_params = g_new0 (UfoInputParam, 1);
+    (*in_params)[0].n_dims = 2;
+    (*in_params)[0].n_expected = -1;
 }
 
 static void
@@ -113,8 +114,8 @@ ufo_scale_task_setup (UfoTask *task,
 
 static void
 ufo_scale_task_get_requisition (UfoTask *task,
-                                 UfoBuffer **inputs,
-                                 UfoRequisition *requisition)
+                                UfoBuffer **inputs,
+                                UfoRequisition *requisition)
 {
     UfoScaleTaskPrivate *priv;
 
@@ -124,7 +125,7 @@ ufo_scale_task_get_requisition (UfoTask *task,
 
 static UfoNode *
 ufo_scale_task_copy_real (UfoNode *node,
-                           GError **error)
+                          GError **error)
 {
     UfoScaleTask *orig;
     UfoScaleTask *copy;