Port laminographic backprojection

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

@@ -1,11 +1,11 @@
 #include <gmodule.h>
 #ifdef __APPLE__
-	#include <OpenCL/cl.h>
+#include <OpenCL/cl.h>
 #else
-	#include <CL/cl.h>
+#include <CL/cl.h>
 #endif
 
-#include <ufo/ufo-filter.h>
+#include <ufo/ufo-filter-reduce.h>
 #include <ufo/ufo-buffer.h>
 #include <ufo/ufo-resource-manager.h>
 
@@ -18,17 +18,19 @@
 
 
 struct _UfoFilterLaminoBPGenericPrivate {
-	    // float theta;
-            cl_kernel bp_kernel;
-	    cl_kernel clean_vol_kernel;
-	    cl_kernel norm_vol_kernel;
-	    CLParameters params;
+    // 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;
+    size_t          global_work_size[3];
 };
 
 GType ufo_filter_lamino_bp_generic_get_type(void) G_GNUC_CONST;
 
-/* Inherit from UFO_TYPE_FILTER */
-G_DEFINE_TYPE(UfoFilterLaminoBPGeneric, ufo_filter_lamino_bp_generic, UFO_TYPE_FILTER);
+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))
 
@@ -46,431 +48,417 @@ enum {
     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");
+//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);
 
-static void ufo_filter_lamino_bp_generic_process(UfoFilter *filter)
-{
-    /////////////////// testing OpenCL platform
-    // testing_cl_platform();
+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);
 
-    g_return_if_fail(UFO_IS_FILTER(filter));
-    UfoFilterLaminoBPGeneric 		*self	= UFO_FILTER_LAMINO_BP_GENERIC(filter);
-    UfoFilterLaminoBPGenericPrivate 	*priv 	= UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(self);
+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_resource_manager();
-    UfoChannel *input_channel  = ufo_filter_get_input_channel(filter);
-    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    GError *tmp_error = NULL;
+    cl_context context;
+    guint width, height;
 
-    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
-    cl_context context = (cl_context) ufo_resource_manager_get_context(manager);
+    // TODO: how to solve 'include problem' for cl-files?
+    // resource_manager_add_program
+    //	 (UfoResourceManager *manager,
+    //	          const gchar *filename, const gchar *options, GError **error)
 
-    UfoBuffer *input  = ufo_channel_get_input_buffer(input_channel);
+    //    ufo_resource_manager_add_program(manager, "lamino_bp_generic.cl",
+    //  	 "-I /home/timurttv/_UFO-Project/framework/ufo-filters-am-installed/lib/ufo", &error);
 
-    ////// init recon-parameters which are common for all projs
-    CLParameters * params = &(priv->params);
+    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);
 
-    // get size of a projection
-    guint proj_num_dims = 0;
-    guint * proj_dim_size = NULL;
-    ufo_buffer_get_dimensions(input, &proj_num_dims, &proj_dim_size);
-    params->proj_sx  = proj_dim_size[0];
-    params->proj_sy  = proj_dim_size[1];
+    /* 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); */
 
-    int vSX = priv->params.vol_sx; 
-    int vSY = priv->params.vol_sy; 
-    int vSZ = priv->params.vol_sz; 
+    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[], gpointer cmd_queue, GError **error)
+{
+    /////////////////// testing OpenCL platform
+    // testing_cl_platform();
 
-    // allocate memory for the reconstructed volume
-    guint vol_num_dims = 3; 
-    guint vol_dim_size[3] = {vSX, vSY, vSZ};
-    size_t global_work_size[3] = { (size_t)vSX,  (size_t) vSY,  (size_t)vSZ};
+    g_return_if_fail(UFO_IS_FILTER(filter));
 
-    ufo_channel_allocate_output_buffers(output_channel, vol_num_dims, vol_dim_size);
-    UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
-    cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output, command_queue);
+    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(filter);
+    CLParameters *params = &(priv->params);
 
+    cl_command_queue command_queue = (cl_command_queue) cmd_queue;
+    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 
+    // 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); 
+    /* 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    
+    /* ufo_buffer_attach_event(output, event); */
+
+    // setup backprojection kernel
     cl_kernel kernel = priv->bp_kernel;
 
     // projection conter
-    int proj_idx = 0;
-    // allocate memory for parameters 
-    cl_mem param_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(CLParameters), NULL, NULL);
-    while(input != NULL)
-    {
-        ////// init recon-parameters which are individual  for each projs
-        params->alpha = - 3 * G_PI/2 + params->theta;
-        params->phi   = params->angle_step*(float)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, param_mem, CL_TRUE, 0,  sizeof(CLParameters), params, 0, NULL, &event);
-
-	// copy projection to GPU
-        cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input,  command_queue);
-        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&input_mem));
-	CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &output_mem));
-        CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&param_mem));
-        // call backprojection routine
-	g_message("processing of %d-th projection", proj_idx);
-        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
-                                              vol_num_dims, NULL, global_work_size, NULL, 0, NULL, &event));
-	 // clFinish(command_queue);
-	// wait untill it finishes its job
-	ufo_buffer_attach_event(output, event);
-
-	// release input buffer
-        ufo_channel_finalize_input_buffer(input_channel, input);
-	// get next projection
-        input = ufo_channel_get_input_buffer(input_channel);
-	
-	proj_idx++;
-    }
+    // 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 void
+ufo_filter_lamino_bp_generic_reduce (UfoFilterReduce *filter, UfoBuffer *output[], gpointer cmd_queue, GError **error)
+{
+    UfoFilterLaminoBPGenericPrivate *priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(filter);
 
     // normalize volume after reconstruction
     cl_kernel norm_vol_kernel = priv->norm_vol_kernel;
+    cl_mem output_mem = ufo_buffer_get_device_array (output[0], (cl_command_queue) 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));
-    // setup the normalization factor
-    float stepPhi = params->angle_step;
     CHECK_OPENCL_ERROR(clSetKernelArg(norm_vol_kernel, 1, sizeof(float), &stepPhi));
+
     // call normalization kernel
     g_message("volume post-processing");
-    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, norm_vol_kernel,
-	                                            vol_num_dims, NULL, global_work_size, NULL, 0, NULL, &event));
-
-
-    // transfer buffer to 3D write filter 
-    ufo_channel_finalize_output_buffer(output_channel, output);
-    // realese output buffer
-    ufo_channel_finish(output_channel);
+    CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel((cl_command_queue) cmd_queue, norm_vol_kernel,
+                vol_num_dims, NULL, priv->global_work_size, NULL, 0, NULL, NULL));
+}
 
-    // free all buffers and cl_mems
-    CHECK_OPENCL_ERROR(clReleaseMemObject(param_mem));
-    g_free(proj_dim_size);
+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)
+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:
+    switch (property_id) {
+        case PROP_THETA:
             self->priv->params.theta = (float) g_value_get_double(value);
             break;
-      case PROP_PSI:
+        case PROP_PSI:
             self->priv->params.psi = (float) g_value_get_double(value);
-	    break;
-      case PROP_ANGLE_STEP:
+            break;
+        case PROP_ANGLE_STEP:
             self->priv->params.angle_step = (float) g_value_get_double(value);
             break;
-      case PROP_VOL_SX:
+        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:
+        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)
+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_initialize(UfoFilter *filter)
-{
-    UfoFilterLaminoBPGeneric *self = UFO_FILTER_LAMINO_BP_GENERIC(filter);
-    UfoResourceManager *manager = ufo_resource_manager();
-    GError *error = NULL;
-    self->priv->bp_kernel = NULL;
-
-    // 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);
-
-     self->priv->bp_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_bp_generic", &error);
-     self->priv->clean_vol_kernel =  ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_clean_vol", &error);
-     self->priv->norm_vol_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic.cl", "lamino_norm_vol", &error);
-     if (error != NULL) {
-            g_warning("%s", error->message);
-            g_error_free(error);
+    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)
+static void
+ufo_filter_lamino_bp_generic_class_init(UfoFilterLaminoBPGenericClass *klass)
 {
     GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
-    UfoFilterClass *filter_class = UFO_FILTER_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->process       = ufo_filter_lamino_bp_generic_process;
-
+    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);
+                "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_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);
+                "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] =
+    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] =
+                "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);
+                "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]);
@@ -489,7 +477,8 @@ static void ufo_filter_lamino_bp_generic_class_init(UfoFilterLaminoBPGenericClas
     g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoBPGenericPrivate));
 }
 
-static void ufo_filter_lamino_bp_generic_init(UfoFilterLaminoBPGeneric *self)
+static void
+ufo_filter_lamino_bp_generic_init(UfoFilterLaminoBPGeneric *self)
 {
     // initialize parameters here
     self->priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(self);
@@ -510,11 +499,12 @@ static void ufo_filter_lamino_bp_generic_init(UfoFilterLaminoBPGeneric *self)
     prms->proj_ox = 0.0;
     prms->proj_oy = 0.0;
 
-    ufo_filter_register_input(UFO_FILTER(self), "projection", 2);
-    ufo_filter_register_output(UFO_FILTER(self), "volume", 3); 
+    ufo_filter_register_inputs (UFO_FILTER (self), 2, NULL);
+    ufo_filter_register_outputs (UFO_FILTER (self), 3, NULL);
 }
 
-G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+G_MODULE_EXPORT UfoFilter *
+ufo_filter_plugin_new(void)
 {
     return g_object_new(UFO_TYPE_FILTER_LAMINO_BP_GENERIC, NULL);
 }

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

@@ -2,7 +2,7 @@
 #define __UFO_FILTER_LAMINO_BP_GENERIC_H__
 
 #include <glib.h>
-#include <ufo/ufo-filter.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))	
@@ -17,7 +17,7 @@ typedef struct _UfoFilterLaminoBPGenericPrivate    UfoFilterLaminoBPGenericPriva
 
 struct _UfoFilterLaminoBPGeneric {
 	   /*< private >*/
-	   UfoFilter parent_instance;
+	   UfoFilterReduce parent_instance;
 	   UfoFilterLaminoBPGenericPrivate *priv;
 };
 
@@ -28,7 +28,7 @@ struct _UfoFilterLaminoBPGeneric {
  */
 struct _UfoFilterLaminoBPGenericClass {
 	   /*< private >*/
-	   UfoFilterClass parent_class;
+	   UfoFilterReduceClass parent_class;
 };
 
 GType ufo_filter_lamino_bp_generic_get_type(void);