Browse Source

Initial commit

Tomas Farago 9 years ago
commit
63d63fb94d

+ 4 - 0
.gitignore

@@ -0,0 +1,4 @@
+tags
+build/
+tests/venv
+_build/

+ 45 - 0
CMakeLists.txt

@@ -0,0 +1,45 @@
+cmake_minimum_required(VERSION 2.6)
+project(ankafilters C)
+
+set(TARNAME "ufo-ankafilters")
+
+set(UFO_ANKAFILTERS_VERSION_MAJOR "0")
+set(UFO_ANKAFILTERS_VERSION_MINOR "1")
+set(UFO_ANKAFILTERS_VERSION_PATCH "0")
+set(UFO_ANKAFILTERS_VERSION_STRING_LONG "${UFO_ANKAFILTERS_VERSION_MAJOR}.${UFO_ANKAFILTERS_VERSION_MINOR}.${UFO_ANKAFILTERS_VERSION_PATCH}")
+set(UFO_ANKAFILTERS_VERSION_STRING_SHORT "${UFO_ANKAFILTERS_VERSION_MAJOR}.${UFO_ANKAFILTERS_VERSION_MINOR}")
+
+set(UFO_DESCRIPTION "UFO optimized filters for usage at ANKA")
+set(UFO_DESCRIPTION_SUMMARY "UFO optimized filters for usage at ANKA")
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/common/cmake")
+
+#{{{ Variables and Options
+include(ConfigurePaths)
+configure_paths(UFO_ANKAFILTERS)
+
+set(UFO_ANKAFILTERS_PLUGINDIR "${UFO_ANKAFILTERS_LIBDIR}/ufo")
+set(UFO_ANKAFILTERS_KERNELDIR "${UFO_ANKAFILTERS_DATADIR}/ufo")
+set(PKG_UFO_CORE_MIN_REQUIRED "0.6")
+#}}}
+
+#{{{ Common dependencies
+find_package(OpenCL REQUIRED)
+find_package(PkgConfig REQUIRED)
+pkg_check_modules(UFO ufo>=${PKG_UFO_CORE_MIN_REQUIRED} REQUIRED)
+
+link_directories(${UFO_LIBRARY_DIRS})
+#}}}
+
+#{{{ Global definitions
+add_definitions("-std=c99 -pedantic -Wall -Wextra -fPIC")
+add_definitions(-DG_LOG_DOMAIN="Ufo")
+
+if (CMAKE_COMPILER_IS_GNUCC OR ("${CMAKE_C_COMPILER_ID}" STREQUAL "Clang"))
+    add_definitions("-Wno-unused-parameter")
+endif ()
+#}}}
+
+#{{{ Subdirectories
+add_subdirectory(src)
+#}}}

+ 90 - 0
common/cmake/ConfigurePaths.cmake

@@ -0,0 +1,90 @@
+# - pre-configured paths for CMake
+#
+# Usage:
+#   configure_paths(<PREFIX>)
+#
+# Checks if configure-like prefix and installation paths were passed by the user
+# and sets up corresponding variables for use in install() commands and to fill
+# out .pc files:
+#
+#   PREFIX_PREFIX       defaults to ...     CMAKE_INSTALL_PREFIX
+#   PREFIX_EPREFIX                          PREFIX_PREFIX
+#   PREFIX_SBINDIR                          PREFIX_EPREFIX/sbin
+#   PREFIX_SYSCONFDIR                       PREFIX_PREFIX/etc
+#   PREFIX_LOCALSTATEDIR                    PREFIX_PREFIX/var
+#   PREFIX_BINDIR                           PREFIX_EPREFIX/bin
+#   PREFIX_LIBDIR                           PREFIX_EPREFIX/lib
+#   PREFIX_INCLUDEDIR                       PREFIX_PREFIX/include
+#   PREFIX_PKGCONFIGDIR                     PREFIX_LIBDIR/pkgconfig
+#   PREFIX_TYPELIBDIR                       PREFIX_LIBDIR/girepository-1.0
+#   PREFIX_DATAROOTDIR                      PREFIX_PREFIX/share
+#   PREFIX_DATADIR                          PREFIX_DATAROOTDIR
+#   PREFIX_INFODIR                          PREFIX_DATAROOTDIR/info
+#   PREFIX_MANDIR                           PREFIX_DATAROOTDIR/man
+#   PREFIX_LOCALEDIR                        PREFIX_DATAROOTDIR/locale
+#   PREFIX_GIRDIR                           PREFIX_DATAROOTDIR/gir-1.0
+
+# Copyright (C) 2013 Matthias Vogelgesang <matthias.vogelgesang@gmail.com>
+#
+# Redistribution and use, with or without modification, are permitted
+# provided that the following conditions are met:
+# 
+#    1. Redistributions must retain the above copyright notice, this
+#       list of conditions and the following disclaimer.
+#    2. The name of the author may not be used to endorse or promote
+#       products derived from this software without specific prior
+#       written permission.
+# 
+# THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR
+# IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+# WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+# ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY
+# DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE
+# GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER
+# IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
+# OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN
+# IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+if(__configure_paths)
+    return()
+endif()
+
+set(__configure_paths YES)
+
+macro(_set_var _prefix _var _user _override _description)
+    set(_name "${_prefix}_${_var}")
+
+    set("${_name}" "${_user}")
+
+    if("${_name}" STREQUAL "")
+        set("${_name}" "${_override}")
+    endif()
+
+    set(${_name} "${${_name}}" CACHE PATH "${_description}")
+    mark_as_advanced(${_name})
+endmacro()
+
+function(configure_paths _prefix)
+    _set_var("${_prefix}" "PREFIX"          "${PREFIX}"         "${CMAKE_INSTALL_PREFIX}"               "install architecture-independent files in PREFIX")
+    _set_var("${_prefix}" "EPREFIX"         "${EXEC_PREFIX}"    "${${_prefix}_PREFIX}"                  "install architecture-dependent files in EPREFIX")
+
+    _set_var("${_prefix}" "SBINDIR"         "${SBINDIR}"        "${${_prefix}_EPREFIX}/sbin"            "system admin executabls")
+    _set_var("${_prefix}" "SYSCONFDIR"      "${SYSCONFDIR}"     "${${_prefix}_PREFIX}/etc"              "read-only single-machine data")
+    _set_var("${_prefix}" "LOCALSTATEDIR"   "${LOCALSTATEDIR}"  "${${_prefix}_PREFIX}/var"              "modifiable single-machine data")
+    _set_var("${_prefix}" "BINDIR"          "${BINDIR}"         "${${_prefix}_EPREFIX}/bin"             "user executables")
+    _set_var("${_prefix}" "LIBDIR"          "${LIBDIR}"         "${${_prefix}_EPREFIX}/lib"             "object code libraries")
+    _set_var("${_prefix}" "INCLUDEDIR"      "${INCLUDEDIR}"     "${${_prefix}_PREFIX}/include"          "C header files")
+    _set_var("${_prefix}" "PKGCONFIGDIR"    "${PKGCONFIGDIR}"   "${${_prefix}_LIBDIR}/pkgconfig"        "pkg-config files")
+    _set_var("${_prefix}" "TYPELIBDIR"      "${TYPELIBDIR}"     "${${_prefix}_LIBDIR}/girepository-1.0" "GObject run-time introspection data")
+    _set_var("${_prefix}" "DATAROOTDIR"     "${DATAROOTDIR}"    "${${_prefix}_PREFIX}/share"            "read-only arch.-independent data root")
+    _set_var("${_prefix}" "DATADIR"         "${DATADIR}"        "${${_prefix}_DATAROOTDIR}"             "read-only architecture-independent data")
+    _set_var("${_prefix}" "INFODIR"         "${INFODIR}"        "${${_prefix}_DATAROOTDIR}/info"        "info documentation")
+    _set_var("${_prefix}" "MANDIR"          "${MANDIR}"         "${${_prefix}_DATAROOTDIR}/man"         "man documentation")
+    _set_var("${_prefix}" "LOCALEDIR"       "${LOCALEDIR}"      "${${_prefix}_DATAROOTDIR}/locale"      "locale-dependent data")
+    _set_var("${_prefix}" "GIRDIR"          "${GIRDIR}"         "${${_prefix}_DATAROOTDIR}/gir-1.0"     "GObject introspection data")
+endfunction()
+
+# vim: tw=0:

+ 90 - 0
common/cmake/FindOpenCL.cmake

@@ -0,0 +1,90 @@
+# - Try to find OpenCL
+# This module tries to find an OpenCL implementation on your system. It supports
+# AMD / ATI, Apple and NVIDIA implementations, but shoudl work, too.
+#
+# Once done this will define
+#  OPENCL_FOUND        - system has OpenCL
+#  OPENCL_INCLUDE_DIRS  - the OpenCL include directory
+#  OPENCL_LIBRARIES    - link these to use OpenCL
+#
+# WIN32 should work, but is untested
+
+FIND_PACKAGE( PackageHandleStandardArgs )
+
+SET (OPENCL_VERSION_STRING "0.1.0")
+SET (OPENCL_VERSION_MAJOR 0)
+SET (OPENCL_VERSION_MINOR 1)
+SET (OPENCL_VERSION_PATCH 0)
+
+IF (APPLE)
+
+  FIND_LIBRARY(OPENCL_LIBRARIES OpenCL DOC "OpenCL lib for OSX")
+  FIND_PATH(OPENCL_INCLUDE_DIRS OpenCL/cl.h DOC "Include for OpenCL on OSX")
+  FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS OpenCL/cl.hpp DOC "Include for OpenCL CPP bindings on OSX")
+
+ELSE (APPLE)
+
+	IF (WIN32)
+	
+	    FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h)
+	    FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp)
+	
+	    # The AMD SDK currently installs both x86 and x86_64 libraries
+	    # This is only a hack to find out architecture
+	    IF( ${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64" )
+	    	SET(OPENCL_LIB_DIR "$ENV{ATISTREAMSDKROOT}/lib/x86_64")
+	    ELSE (${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64")
+	    	SET(OPENCL_LIB_DIR "$ENV{ATISTREAMSDKROOT}/lib/x86")
+	    ENDIF( ${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64" )
+	    FIND_LIBRARY(OPENCL_LIBRARIES OpenCL.lib ${OPENCL_LIB_DIR})
+	    
+	    GET_FILENAME_COMPONENT(_OPENCL_INC_CAND ${OPENCL_LIB_DIR}/../../include ABSOLUTE)
+	    
+	    # On Win32 search relative to the library
+	    FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS "${_OPENCL_INC_CAND}")
+	    FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS "${_OPENCL_INC_CAND}")
+	
+	ELSE (WIN32)
+
+            # Unix style platforms
+            FIND_LIBRARY(OPENCL_LIBRARIES OpenCL
+              ENV LD_LIBRARY_PATH
+              /usr/lib/nvidia-current
+              /usr/lib64/nvidia
+              /opt/nvidia-current
+              /opt/AMDAPP/lib
+              /opt/AMDAPP/lib/x86_64
+            )
+
+            GET_FILENAME_COMPONENT(OPENCL_LIB_DIR ${OPENCL_LIBRARIES} PATH)
+            GET_FILENAME_COMPONENT(_OPENCL_INC_CAND ${OPENCL_LIB_DIR}/../../include ABSOLUTE)
+
+            # The AMD SDK currently does not place its headers
+            # in /usr/include, therefore also search relative
+            # to the library
+            FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS 
+                ${_OPENCL_INC_CAND} 
+                /usr/local/cuda/include 
+                /opt/cuda/include
+                /opt/AMDAPP/include)
+            FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS ${_OPENCL_INC_CAND})
+
+	ENDIF (WIN32)
+
+ENDIF (APPLE)
+
+FIND_PACKAGE_HANDLE_STANDARD_ARGS( OpenCL DEFAULT_MSG OPENCL_LIBRARIES OPENCL_INCLUDE_DIRS )
+
+IF( _OPENCL_CPP_INCLUDE_DIRS )
+	SET( OPENCL_HAS_CPP_BINDINGS TRUE )
+	LIST( APPEND OPENCL_INCLUDE_DIRS ${_OPENCL_CPP_INCLUDE_DIRS} )
+	# This is often the same, so clean up
+	LIST( REMOVE_DUPLICATES OPENCL_INCLUDE_DIRS )
+ENDIF( _OPENCL_CPP_INCLUDE_DIRS )
+
+MARK_AS_ADVANCED(
+  OPENCL_LIBRARIES
+  OPENCL_INCLUDE_DIRS
+  _OPENCL_CPP_INCLUDE_DIRS
+)
+

+ 63 - 0
src/CMakeLists.txt

@@ -0,0 +1,63 @@
+cmake_minimum_required(VERSION 2.6)
+
+#{{{ Sources
+set(ufofilter_SRCS
+    ufo-anka-backproject-task.c
+    )
+
+file(GLOB ufofilter_KERNELS "kernels/*.cl")
+#}}}
+#{{{ Variables
+if (CMAKE_COMPILER_IS_GNUCC OR ("${CMAKE_C_COMPILER_ID}" STREQUAL "Clang"))
+    add_definitions("-Wcast-align -Wcast-qual -Winline -Wmissing-declarations "
+                    "-Wmissing-prototypes -Wnested-externs -Wno-long-long "
+                    "-Wno-missing-field-initializers -Wpointer-arith "
+                    "-Wredundant-decls -Wshadow -Wstrict-prototypes -Wwrite-strings")
+endif()
+#}}}
+#{{{ Plugin targets
+include_directories(${CMAKE_CURRENT_BINARY_DIR}
+                    ${CMAKE_CURRENT_SOURCE_DIR}
+                    ${OPENCL_INCLUDE_DIRS}
+                    ${UFO_INCLUDE_DIRS})
+
+foreach(_src ${ufofilter_SRCS})
+    # find plugin suffix
+    string(REGEX REPLACE "ufo-([^ \\.]+)-task.*" "\\1" task "${_src}")
+
+    # build string to get miscalleanous sources
+    string(REPLACE "-" "_" _misc ${task})
+    string(TOUPPER ${_misc} _misc_upper)
+
+    # create an option name and add this to disable filters
+    set(target_option "ENABLE_${_misc_upper}")
+    option(${target_option} "Build filter ${task}" ON)
+
+    if (${target_option})
+        set(_misc "${_misc}_misc_SRCS")
+
+        string(REPLACE "-" "" _targetname ${task})
+        set(target "ufofilter${_targetname}")
+
+        # build single shared library per filter
+        if (${CMAKE_SYSTEM_NAME} MATCHES "Darwin")
+            add_library(${target} MODULE ${_src} ${${_misc}})
+        else()
+            add_library(${target} SHARED ${_src} ${${_misc}})
+        endif()
+
+        target_link_libraries(${target} ${ufofilter_LIBS})
+
+        list(APPEND all_targets ${target})
+        
+        install(TARGETS ${target}
+                ARCHIVE DESTINATION ${UFO_ANKAFILTERS_PLUGINDIR}
+                LIBRARY DESTINATION ${UFO_ANKAFILTERS_PLUGINDIR})
+    endif()
+endforeach()
+
+# copy kernels
+foreach(_kernel ${ufofilter_KERNELS})
+    install(FILES ${_kernel} DESTINATION ${UFO_ANKAFILTERS_KERNELDIR})
+endforeach()
+#}}}

+ 49 - 0
src/kernels/ankabackproject.cl

@@ -0,0 +1,49 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+ __kernel void backproject (global float *volume,
+                            __read_only image2d_t projection,
+                            const sampler_t sampler,
+                            const int2 x_region,
+                            const int2 y_region,
+                            const int2 z_region,
+                            const float8 tmatrix,
+                            const int2 offset,
+                            const uint cumulate)
+{
+     int3 id = (int3) (get_global_id (0), get_global_id (1), get_global_id (2));
+     float2 pixel;
+     float3 voxel;
+ 
+     voxel.x = x_region.x + id.x * x_region.y;
+     voxel.y = y_region.x + id.y * y_region.y;
+     voxel.z = z_region.x + id.z * z_region.y;
+ 
+     pixel.x = voxel.x * tmatrix.s0 + voxel.y * tmatrix.s1 + tmatrix.s3 - offset.x;
+     pixel.y = voxel.x * tmatrix.s4 + voxel.y * tmatrix.s5 + voxel.z * tmatrix.s6 + tmatrix.s7 - offset.y;
+ 
+    if (cumulate) {
+         volume[id.z * get_global_size (0) * get_global_size (1) +
+                id.y * get_global_size (0) + id.x] += read_imagef (projection, sampler, pixel).x;
+    }
+    else {
+         volume[id.z * get_global_size (0) * get_global_size (1) +
+                id.y * get_global_size (0) + id.x] = read_imagef (projection, sampler, pixel).x;
+    }
+ }

+ 514 - 0
src/ufo-anka-backproject-task.c

@@ -0,0 +1,514 @@
+/*
+ * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+#include <math.h>
+#include <string.h>
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "ufo-anka-backproject-task.h"
+
+#define EXTRACT_INT(region, index) g_value_get_int (g_value_array_get_nth ((region), (index)))
+#define EXTRACT_FLOAT(region, index) g_value_get_float (g_value_array_get_nth ((region), (index)))
+#define REGION_SIZE(region) ((EXTRACT_INT ((region), 2)) == 0) ? 0 : \
+                            ((EXTRACT_INT ((region), 1) - EXTRACT_INT ((region), 0) - 1) /\
+                            EXTRACT_INT ((region), 2) + 1)
+
+/**
+ * SECTION:ufo-anka-backproject-task
+ * @Short_description: Backproject projection by projection
+ * @Title: anka_backproject
+ *
+ */
+
+struct _UfoAnkaBackprojectTaskPrivate {
+    /* private */
+    gboolean generated;
+    guint count;
+    float tmatrix[8];
+
+    /* OpenCL */
+    cl_context context;
+    cl_kernel bp_kernel;
+    cl_sampler sampler;
+
+    /* properties */
+    GValueArray *x_region;
+    GValueArray *y_region;
+    GValueArray *z_region;
+    GValueArray *center;
+    GValueArray *projection_offset;
+    gboolean tomo_angle_is_absolute;
+    gfloat tomo_angle;
+    gfloat lamino_angle;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoAnkaBackprojectTask, ufo_anka_backproject_task, UFO_TYPE_TASK_NODE,
+                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+                                                ufo_task_interface_init))
+
+#define UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_ANKA_BACKPROJECT_TASK, UfoAnkaBackprojectTaskPrivate))
+
+enum {
+    PROP_0,
+    PROP_X_REGION,
+    PROP_Y_REGION,
+    PROP_Z_REGION,
+    PROP_PROJECTION_OFFSET,
+    PROP_CENTER,
+    PROP_TOMO_ANGLE_IS_ABSOLUTE,
+    PROP_TOMO_ANGLE,
+    PROP_LAMINO_ANGLE,
+    N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+static void
+create_transformation_matrix (UfoAnkaBackprojectTaskPrivate *priv, float tomo_angle)
+{
+    priv->tmatrix[0] = cos (tomo_angle);
+    priv->tmatrix[1] = sin (tomo_angle);
+    priv->tmatrix[2] = 0.0f;
+    priv->tmatrix[3] = EXTRACT_FLOAT (priv->center, 0);
+    priv->tmatrix[4] = cos (priv->lamino_angle) * sin (tomo_angle);
+    priv->tmatrix[5] = -cos (priv->lamino_angle) * cos(tomo_angle);
+    priv->tmatrix[6] = sin(priv->lamino_angle);
+    priv->tmatrix[7] = EXTRACT_FLOAT (priv->center, 1);
+}
+
+UfoNode *
+ufo_anka_backproject_task_new (void)
+{
+    return UFO_NODE (g_object_new (UFO_TYPE_ANKA_BACKPROJECT_TASK, NULL));
+}
+
+static void
+ufo_anka_backproject_task_setup (UfoTask *task,
+                                 UfoResources *resources,
+                                 GError **error)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+    cl_int cl_error;
+
+    priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (task);
+    priv->context = ufo_resources_get_context (resources);
+    priv->bp_kernel = ufo_resources_get_kernel (resources, "ankabackproject.cl", "backproject", error);
+    priv->sampler = clCreateSampler (priv->context,
+                                     (cl_bool) FALSE,
+                                     CL_ADDRESS_CLAMP,
+                                     CL_FILTER_LINEAR,
+                                     &cl_error);
+
+    UFO_RESOURCES_CHECK_CLERR (clRetainContext (priv->context));
+    UFO_RESOURCES_CHECK_CLERR (cl_error);
+    if (priv->bp_kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->bp_kernel));
+    }
+}
+
+static void
+ufo_anka_backproject_task_get_requisition (UfoTask *task,
+                                           UfoBuffer **inputs,
+                                           UfoRequisition *requisition)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+
+    priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (task);
+
+    requisition->n_dims = 3;
+    requisition->dims[0] = REGION_SIZE (priv->x_region);
+    requisition->dims[1] = REGION_SIZE (priv->y_region);
+    requisition->dims[2] = REGION_SIZE (priv->z_region);
+}
+
+static guint
+ufo_anka_backproject_task_get_num_inputs (UfoTask *task)
+{
+    return 1;
+}
+
+static guint
+ufo_anka_backproject_task_get_num_dimensions (UfoTask *task,
+                                              guint input)
+{
+    g_return_val_if_fail (input == 0, 0);
+
+    return 3;
+}
+
+static gboolean
+ufo_anka_backproject_task_equal_real (UfoNode *n1,
+                                      UfoNode *n2)
+{
+    g_return_val_if_fail (UFO_IS_ANKA_BACKPROJECT_TASK (n1) && UFO_IS_ANKA_BACKPROJECT_TASK (n2), FALSE);
+
+    return UFO_ANKA_BACKPROJECT_TASK (n1)->priv->bp_kernel == UFO_ANKA_BACKPROJECT_TASK (n2)->priv->bp_kernel;
+}
+
+static UfoTaskMode
+ufo_anka_backproject_task_get_mode (UfoTask *task)
+{
+    return UFO_TASK_MODE_REDUCTOR | UFO_TASK_MODE_GPU;
+}
+
+static gboolean
+ufo_anka_backproject_task_process (UfoTask *task,
+                                   UfoBuffer **inputs,
+                                   UfoBuffer *output,
+                                   UfoRequisition *requisition)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+    UfoGpuNode *node;
+    UfoProfiler *profiler;
+    gfloat tomo_angle;
+    /* regions stripped off the "to" value */
+    gint x_region[2], y_region[2], z_region[2], proj_offset[2];
+    cl_command_queue cmd_queue;
+    cl_mem image;
+    cl_mem out_mem;
+
+    priv = UFO_ANKA_BACKPROJECT_TASK (task)->priv;
+    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+    out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+    image = ufo_buffer_get_device_image (inputs[0], cmd_queue); 
+
+    x_region[0] = EXTRACT_INT (priv->x_region, 0);
+    x_region[1] = EXTRACT_INT (priv->x_region, 2);
+
+    y_region[0] = EXTRACT_INT (priv->y_region, 0);
+    y_region[1] = EXTRACT_INT (priv->y_region, 2);
+
+    z_region[0] = EXTRACT_INT (priv->z_region, 0);
+    z_region[1] = EXTRACT_INT (priv->z_region, 2);
+
+    proj_offset[0] = EXTRACT_INT (priv->projection_offset, 0);
+    proj_offset[1] = EXTRACT_INT (priv->projection_offset, 1);
+
+
+    tomo_angle = priv->tomo_angle_is_absolute ? priv->tomo_angle : priv->tomo_angle * priv->count;
+    create_transformation_matrix (priv, tomo_angle);
+
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 0, sizeof (cl_mem), &out_mem));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 1, sizeof (cl_mem), &image));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 2, sizeof (cl_sampler), &priv->sampler));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 3, sizeof (cl_int2), x_region));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 4, sizeof (cl_int2), y_region));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 5, sizeof (cl_int2), z_region));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 6, sizeof (cl_float8), priv->tmatrix));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 7, sizeof (cl_int2), proj_offset));
+    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->bp_kernel, 8, sizeof (cl_uint), (cl_uint *) &priv->count));
+
+    profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+    ufo_profiler_call (profiler, cmd_queue, priv->bp_kernel, 3, requisition->dims, NULL);
+
+    priv->count++;
+
+    return TRUE;
+}
+
+static gboolean
+ufo_anka_backproject_task_generate (UfoTask *task,
+                                    UfoBuffer *output,
+                                    UfoRequisition *requisition)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+
+    priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (task);
+
+    if (priv->generated) {
+        return FALSE;
+    }
+
+    priv->generated = TRUE;
+
+    return TRUE;
+}
+
+static void
+ufo_anka_backproject_task_finalize (GObject *object)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+
+    priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (object);
+    g_value_array_free (priv->x_region);
+    g_value_array_free (priv->y_region);
+    g_value_array_free (priv->z_region);
+    g_value_array_free (priv->projection_offset);
+    g_value_array_free (priv->center);
+
+    if (priv->bp_kernel) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->bp_kernel));
+        priv->bp_kernel = NULL;
+    }
+    if (priv->context) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseContext (priv->context));
+        priv->context = NULL;
+    }
+    if (priv->sampler) {
+        UFO_RESOURCES_CHECK_CLERR (clReleaseSampler (priv->sampler));
+        priv->sampler = NULL;
+    }
+
+    G_OBJECT_CLASS (ufo_anka_backproject_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+    iface->setup = ufo_anka_backproject_task_setup;
+    iface->get_requisition = ufo_anka_backproject_task_get_requisition;
+    iface->get_num_inputs = ufo_anka_backproject_task_get_num_inputs;
+    iface->get_num_dimensions = ufo_anka_backproject_task_get_num_dimensions;
+    iface->get_mode = ufo_anka_backproject_task_get_mode;
+    iface->process = ufo_anka_backproject_task_process;
+    iface->generate = ufo_anka_backproject_task_generate;
+}
+
+static void
+ufo_anka_backproject_task_set_property (GObject *object,
+                                        guint property_id,
+                                        const GValue *value,
+                                        GParamSpec *pspec)
+{
+    UfoAnkaBackprojectTaskPrivate *priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (object);
+    GValueArray *array;
+
+
+    switch (property_id) {
+        case PROP_X_REGION:
+            array = (GValueArray *) g_value_get_boxed (value);
+            g_value_array_free (priv->x_region);
+            priv->x_region = g_value_array_copy (array);
+            break;
+        case PROP_Y_REGION:
+            array = (GValueArray *) g_value_get_boxed (value);
+            g_value_array_free (priv->y_region);
+            priv->y_region = g_value_array_copy (array);
+            break;
+        case PROP_Z_REGION:
+            array = (GValueArray *) g_value_get_boxed (value);
+            g_value_array_free (priv->z_region);
+            priv->z_region = g_value_array_copy (array);
+            break;
+        case PROP_PROJECTION_OFFSET:
+            array = (GValueArray *) g_value_get_boxed (value);
+            g_value_array_free (priv->projection_offset);
+            priv->projection_offset = g_value_array_copy (array);
+            break;
+        case PROP_CENTER:
+            array = (GValueArray *) g_value_get_boxed (value);
+            g_value_array_free (priv->center);
+            priv->center = g_value_array_copy (array);
+            break;
+        case PROP_TOMO_ANGLE_IS_ABSOLUTE:
+            priv->tomo_angle_is_absolute = g_value_get_boolean (value);
+            break;
+        case PROP_TOMO_ANGLE:
+            priv->tomo_angle = g_value_get_float (value);
+            break;
+        case PROP_LAMINO_ANGLE:
+            priv->lamino_angle = g_value_get_float (value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_anka_backproject_task_get_property (GObject *object,
+                                        guint property_id,
+                                        GValue *value,
+                                        GParamSpec *pspec)
+{
+    UfoAnkaBackprojectTaskPrivate *priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE (object);
+
+    switch (property_id) {
+        case PROP_X_REGION:
+            g_value_set_boxed (value, priv->x_region);
+            break;
+        case PROP_Y_REGION:
+            g_value_set_boxed (value, priv->y_region);
+            break;
+        case PROP_Z_REGION:
+            g_value_set_boxed (value, priv->z_region);
+            break;
+        case PROP_PROJECTION_OFFSET:
+            g_value_set_boxed (value, priv->projection_offset);
+            break;
+        case PROP_CENTER:
+            g_value_set_boxed (value, priv->center);
+            break;
+        case PROP_TOMO_ANGLE_IS_ABSOLUTE:
+            g_value_set_boolean (value, priv->tomo_angle_is_absolute);
+            break;
+        case PROP_TOMO_ANGLE:
+            g_value_set_float (value, priv->tomo_angle);
+            break;
+        case PROP_LAMINO_ANGLE:
+            g_value_set_float (value, priv->lamino_angle);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+            break;
+    }
+}
+
+static void
+ufo_anka_backproject_task_class_init (UfoAnkaBackprojectTaskClass *klass)
+{
+    GObjectClass *oclass;
+    UfoNodeClass *node_class;
+    
+    oclass = G_OBJECT_CLASS (klass);
+    node_class = UFO_NODE_CLASS (klass);
+
+    oclass->finalize = ufo_anka_backproject_task_finalize;
+    oclass->set_property = ufo_anka_backproject_task_set_property;
+    oclass->get_property = ufo_anka_backproject_task_get_property;
+
+
+    GParamSpec *region_vals = g_param_spec_int ("region_values",
+                                                "Region values",
+                                                "Elements in regions",
+                                                G_MININT,
+                                                G_MAXINT,
+                                                (gint) 0,
+                                                G_PARAM_READWRITE);
+
+    GParamSpec *float_region_vals = g_param_spec_float ("float_region_values",
+                                                        "Float Region values",
+                                                        "Elements in float regions",
+                                                        -G_MAXFLOAT,
+                                                        G_MAXFLOAT,
+                                                        0.0f,
+                                                        G_PARAM_READWRITE);
+
+    properties[PROP_X_REGION] =
+        g_param_spec_value_array ("x_region",
+                                  "X region for reconstruction as (from, to, step)",
+                                  "X region for reconstruction as (from, to, step)",
+                                  region_vals,
+                                  G_PARAM_READWRITE);
+
+    properties[PROP_Y_REGION] =
+        g_param_spec_value_array ("y_region",
+                                  "Y region for reconstruction as (from, to, step)",
+                                  "Y region for reconstruction as (from, to, step)",
+                                  region_vals,
+                                  G_PARAM_READWRITE);
+
+    properties[PROP_Z_REGION] =
+        g_param_spec_value_array ("z_region",
+                                  "Z region for reconstruction as (from, to, step)",
+                                  "Z region for reconstruction as (from, to, step)",
+                                  region_vals,
+                                  G_PARAM_READWRITE);
+
+    properties[PROP_PROJECTION_OFFSET] =
+        g_param_spec_value_array ("projection-offset",
+                                  "Offset to projection data as (x, y)",
+                                  "Offset to projection data as (x, y) for the case input data \
+                                  is cropped to the necessary range of interest",
+                                  region_vals,
+                                  G_PARAM_READWRITE);
+
+    properties[PROP_CENTER] =
+        g_param_spec_value_array ("center",
+                                  "Center of the volume with respect to projections (x, y)",
+                                  "Center of the volume with respect to projections (x, y), (rotation axes)",
+                                  float_region_vals,
+                                  G_PARAM_READWRITE);
+
+    properties[PROP_TOMO_ANGLE_IS_ABSOLUTE] =
+        g_param_spec_boolean ("tomo-angle-is-absolute",
+                              "Tomographic angle is absolute",
+                              "If TRUE, the value stored in tomo-angle property represents \
+                              an absolute angle, relative otherwise",
+                              FALSE,
+                              G_PARAM_READWRITE);
+
+    properties[PROP_TOMO_ANGLE] =
+        g_param_spec_float ("tomo-angle",
+                            "Tomographic rotation angle in radians",
+                            "Tomographic rotation angle in radians (used for acquiring projections)",
+                            -G_MAXFLOAT,
+                            G_MAXFLOAT,
+                            0.0f,
+                            G_PARAM_READWRITE);
+
+    properties[PROP_LAMINO_ANGLE] =
+        g_param_spec_float ("lamino-angle",
+                            "Absolute laminogrpahic angle in radians",
+                            "Absolute laminogrpahic angle in radians determining the sample tilt",
+                            0.0f,
+                            (float) G_PI / 2,
+                            0.0f,
+                            G_PARAM_READWRITE);
+
+    for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+        g_object_class_install_property (oclass, i, properties[i]);
+
+    node_class->equal = ufo_anka_backproject_task_equal_real;
+
+    g_type_class_add_private (klass, sizeof(UfoAnkaBackprojectTaskPrivate));
+}
+
+static void
+ufo_anka_backproject_task_init(UfoAnkaBackprojectTask *self)
+{
+    UfoAnkaBackprojectTaskPrivate *priv;
+    self->priv = priv = UFO_ANKA_BACKPROJECT_TASK_GET_PRIVATE(self);
+    guint i;
+    GValue int_zero = G_VALUE_INIT;
+    GValue float_zero = G_VALUE_INIT;
+
+    g_value_init (&int_zero, G_TYPE_INT);
+    g_value_init (&float_zero, G_TYPE_FLOAT);
+    g_value_set_int (&int_zero, 0);
+    g_value_set_float (&float_zero, 0.0f);
+    self->priv->x_region = g_value_array_new (3);
+    self->priv->y_region = g_value_array_new (3);
+    self->priv->z_region = g_value_array_new (3);
+    self->priv->projection_offset = g_value_array_new (2);
+    self->priv->center = g_value_array_new (2);
+
+    for (i = 0; i < 3; i++) {
+        g_value_array_insert (self->priv->x_region, i, &int_zero);
+        g_value_array_insert (self->priv->y_region, i, &int_zero);
+        g_value_array_insert (self->priv->z_region, i, &int_zero);
+        if (i < 2) {
+            g_value_array_insert (self->priv->projection_offset, i, &int_zero);
+            g_value_array_insert (self->priv->center, i, &float_zero);
+        }
+    }
+
+    self->priv->tomo_angle_is_absolute = FALSE;
+    self->priv->tomo_angle = 0.0f;
+    self->priv->lamino_angle = 0.0f;
+    self->priv->count = 0;
+    self->priv->generated = FALSE;
+}

+ 66 - 0
src/ufo-anka-backproject-task.h

@@ -0,0 +1,66 @@
+/*
+ * Copyright (C) 2011-2013 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifndef __UFO_ANKA_BACKPROJECT_TASK_H
+#define __UFO_ANKA_BACKPROJECT_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_ANKA_BACKPROJECT_TASK             (ufo_anka_backproject_task_get_type())
+#define UFO_ANKA_BACKPROJECT_TASK(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_ANKA_BACKPROJECT_TASK, UfoAnkaBackprojectTask))
+#define UFO_IS_ANKA_BACKPROJECT_TASK(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_ANKA_BACKPROJECT_TASK))
+#define UFO_ANKA_BACKPROJECT_TASK_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_ANKA_BACKPROJECT_TASK, UfoAnkaBackprojectTaskClass))
+#define UFO_IS_ANKA_BACKPROJECT_TASK_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_ANKA_BACKPROJECT_TASK))
+#define UFO_ANKA_BACKPROJECT_TASK_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_ANKA_BACKPROJECT_TASK, UfoAnkaBackprojectTaskClass))
+
+typedef struct _UfoAnkaBackprojectTask           UfoAnkaBackprojectTask;
+typedef struct _UfoAnkaBackprojectTaskClass      UfoAnkaBackprojectTaskClass;
+typedef struct _UfoAnkaBackprojectTaskPrivate    UfoAnkaBackprojectTaskPrivate;
+
+/**
+ * UfoAnkaBackprojectTask:
+ *
+ * [ADD DESCRIPTION HERE]. The contents of the #UfoAnkaBackprojectTask structure
+ * are private and should only be accessed via the provided API.
+ */
+struct _UfoAnkaBackprojectTask {
+    /*< private >*/
+    UfoTaskNode parent_instance;
+
+    UfoAnkaBackprojectTaskPrivate *priv;
+};
+
+/**
+ * UfoAnkaBackprojectTaskClass:
+ *
+ * #UfoAnkaBackprojectTask class
+ */
+struct _UfoAnkaBackprojectTaskClass {
+    /*< private >*/
+    UfoTaskNodeClass parent_class;
+};
+
+UfoNode  *ufo_anka_backproject_task_new       (void);
+GType     ufo_anka_backproject_task_get_type  (void);
+
+G_END_DECLS
+
+#endif