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

#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-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;
	    CLParameters params;
};

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);

#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_process(UfoFilter *filter)
{
    /////////////////// testing OpenCL platform
    // testing_cl_platform();

    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);

    UfoResourceManager *manager = ufo_resource_manager();
    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);
    cl_context context = (cl_context) ufo_resource_manager_get_context(manager);

    UfoBuffer *input  = ufo_channel_get_input_buffer(input_channel);

    ////// init recon-parameters which are common for all projs
    CLParameters * params = &(priv->params);

    // 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];

    int vSX = priv->params.vol_sx; 
    int vSY = priv->params.vol_sy; 
    int vSZ = priv->params.vol_sz; 

    // shift volume origin to the box center
    params->vol_ox += (float)vSX/2.0;
    params->vol_oy += (float)vSY/2.0;
    params->vol_oz += (float)vSZ/2.0;

    // 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};

    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);


    // 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));
     // whait until kernel finishes its job
     ufo_buffer_attach_event(output, event);
      
    // setup backprojection kernel    
    cl_kernel kernel = priv->bp_kernel;
    CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &output_mem));

    // projection conter
    int proj_idx = 0;
    // allocate memory for parameters 
    cl_mem param_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(CLParameters), NULL, NULL);
g_message("loop"); 
    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;

        // 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, 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));
	// 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++;
    }

    // normalize volume after reconstruction
    cl_kernel norm_vol_kernel = priv->norm_vol_kernel;
    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));

    // temporal saving unless 3d writer is not ready
    /*FILE * amfile = fopen("_reco.bin","wb");
    float * recobuf = ufo_buffer_get_host_array(output, command_queue);
    fwrite(recobuf, sizeof(float), vSX*vSY*vSZ, amfile);
    fclose(amfile);*/

    // transfer buffer to 3D write filter 
    ufo_channel_finalize_output_buffer(output_channel, output);
    // realese output buffer
    ufo_channel_finish(output_channel);

    // free all buffers and cl_mems
    CHECK_OPENCL_ERROR(clReleaseMemObject(param_mem));
    g_free(proj_dim_size);
}

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_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?
    ufo_resource_manager_add_program(manager, "lamino_bp_generic.cl",
		    "-I /home/timurttv/_UFO-Project/framework/ufo-filters-am-installed/lib/ufo", &error);

    if (error != NULL) {
            g_warning("%s", error->message);
            g_error_free(error);
            return;
     }

     self->priv->bp_kernel = ufo_resource_manager_get_kernel(manager, "lamino_bp_generic", &error);
     self->priv->clean_vol_kernel =  ufo_resource_manager_get_kernel(manager, "lamino_clean_vol", &error);
     self->priv->norm_vol_kernel = ufo_resource_manager_get_kernel(manager, "lamino_norm_vol", &error);
     if (error != NULL) {
            g_warning("%s", error->message);
            g_error_free(error);
    }
}

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);

    gobject_class->set_property = ufo_filter_lamino_bp_generic_set_property;
    gobject_class->get_property = ufo_filter_lamino_bp_generic_get_property;
    filter_class->initialize    = ufo_filter_lamino_bp_generic_initialize;
    filter_class->process       = ufo_filter_lamino_bp_generic_process;


    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)
{
    // initialize parameters here
    self->priv = UFO_FILTER_LAMINO_BP_GENERIC_GET_PRIVATE(self);

    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_input(UFO_FILTER(self), "projection", 2);
    ufo_filter_register_output(UFO_FILTER(self), "volume", 3); 
}

G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
{
    return g_object_new(UFO_TYPE_FILTER_LAMINO_BP_GENERIC, NULL);
}