Browse Source

initial commit for filters developed by Anton Myagotin

Anton Myagotin 12 years ago
commit
d951420c5f

+ 63 - 0
CMakeLists.txt

@@ -0,0 +1,63 @@
+cmake_minimum_required(VERSION 2.6)
+project(ufo)
+
+set(TARNAME "ufo-filters")
+set(UFO_FILTERS_VERSION_MAJOR "0")
+set(UFO_FILTERS_VERSION_MINOR "2")
+set(UFO_FILTERS_VERSION_PATCH "0")
+set(UFO_FILTERS_VERSION_STRING_LONG "${UFO_FILTERS_VERSION_MAJOR}.${UFO_FILTERS_VERSION_MINOR}.${UFO_FILTERS_VERSION_PATCH}")
+set(UFO_FILTERS_VERSION_STRING_SHORT "${UFO_FILTERS_VERSION_MAJOR}.${UFO_FILTERS_VERSION_MINOR}")
+
+set(UFO_DESCRIPTION "UFO good filters")
+set(UFO_DESCRIPTION_SUMMARY "UFO good filters")
+
+set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/common/cmake")
+
+# --- Options -----------------------------------------------------------------
+option(WITH_PROFILING "Enable profiling" OFF)
+if (WITH_PROFILING)
+    add_definitions("-pg")
+    set(CMAKE_C_FLAGS "-pg")
+endif ()
+
+# --- Find packages and libraries ---------------------------------------------
+find_package(OpenCL REQUIRED)
+find_package(PkgConfig REQUIRED)
+pkg_check_modules(GLIB2 glib-2.0>=2.24 REQUIRED)
+pkg_check_modules(GOBJECT2 gobject-2.0>=2.24 REQUIRED)
+pkg_check_modules(UFO ufo>=0.2 REQUIRED)
+
+include_directories(
+    ${GLIB2_INCLUDE_DIRS}
+    ${OPENCL_INCLUDE_DIRS}
+    ${UFO_INCLUDE_DIRS}
+    )
+
+add_definitions("-std=c99 -Wall -fPIC")
+add_definitions(-DG_LOG_DOMAIN="Ufo")
+
+add_subdirectory(src)
+
+# --- Package generation ------------------------------------------------------
+set(CPACK_PACKAGE_DESCRIPTION ${UFO_DESCRIPTION})
+set(CPACK_PACKAGE_DESCRIPTION_SUMMARY ${UFO_DESCRIPTION_SUMMARY})
+set(CPACK_PACKAGE_NAME ${TARNAME})
+
+set(CPACK_PACKAGE_CONTACT "matthias.vogelgesang@kit.edu")
+set(CPACK_PACKAGE_VENDOR "Karlsruhe Institute of Technology/IPE")
+set(CPACK_PACKAGE_VERSION ${UCA_FILTERS_VERSION_STRING_LONG})
+set(CPACK_PACKAGE_VERSION_MAJOR ${UFO_FILTERS_VERSION_MAJOR})
+set(CPACK_PACKAGE_VERSION_MINOR ${UFO_FILTERS_VERSION_MINOR})
+set(CPACK_PACKAGE_VERSION_PATCH ${UFO_FILTERS_VERSION_PATCH})
+set(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_NAME}-${UFO_FILTERS_VERSION_STRING_LONG}-${CMAKE_SYSTEM_PROCESSOR}")
+set(VERSION ${UFO_FILTERS_VERSION})
+
+set(CPACK_GENERATOR "DEB;RPM;")
+set(CPACK_SOURCE_GENERATOR "TGZ")
+set(CPACK_SOURCE_IGNORE_FILES "tags" ".bzr" ".swp" "~1~")
+set(CPACK_SOURCE_PACKAGE_FILE_NAME "${TARNAME}-${UFO_FILTERS_VERSION_STRING_LONG}" CACHE INTERNAL "tarball basename")
+
+# --- Distro specific
+set(CPACK_DEBIAN_PACKAGE_DEPENDS "libc6 (>= 2.3.6), libgcc1 (>= 1:4.1)")
+
+include(CPack)

+ 133 - 0
common/cmake/FindFFTW3.cmake

@@ -0,0 +1,133 @@
+# 
+# Try to find FFTW3  library  
+# (see www.fftw.org)
+# Once run this will define: 
+# 
+# FFTW3_FOUND
+# FFTW3_INCLUDE_DIR 
+# FFTW3_LIBRARIES
+# FFTW3_LINK_DIRECTORIES
+#
+# You may set one of these options before including this file:
+#  FFTW3_USE_SSE2
+#
+#  TODO: _F_ versions.
+#
+# Jan Woetzel 05/2004
+# www.mip.informatik.uni-kiel.de
+# --------------------------------
+
+ FIND_PATH(FFTW3_INCLUDE_DIR fftw3.h
+   ${FFTW3_DIR}/include
+   ${FFTW3_HOME}/include
+   ${FFTW3_DIR}
+   ${FFTW3_HOME}
+   $ENV{FFTW3_DIR}/include
+   $ENV{FFTW3_HOME}/include
+   $ENV{FFTW3_DIR}
+   $ENV{FFTW3_HOME}
+   /usr/include
+   /usr/local/include
+   $ENV{SOURCE_DIR}/fftw3
+   $ENV{SOURCE_DIR}/fftw3/include
+   $ENV{SOURCE_DIR}/fftw
+   $ENV{SOURCE_DIR}/fftw/include
+ )
+#MESSAGE("DBG FFTW3_INCLUDE_DIR=${FFTW3_INCLUDE_DIR}")  
+
+
+SET(FFTW3_POSSIBLE_LIBRARY_PATH
+  ${FFTW3_DIR}/lib
+  ${FFTW3_HOME}/lib
+  ${FFTW3_DIR}
+  ${FFTW3_HOME}  
+  $ENV{FFTW3_DIR}/lib
+  $ENV{FFTW3_HOME}/lib
+  $ENV{FFTW3_DIR}
+  $ENV{FFTW3_HOME}  
+  /usr/lib
+  /usr/local/lib
+  $ENV{SOURCE_DIR}/fftw3
+  $ENV{SOURCE_DIR}/fftw3/lib
+  $ENV{SOURCE_DIR}/fftw
+  $ENV{SOURCE_DIR}/fftw/lib
+)
+
+  
+# the lib prefix is containe din filename onf W32, unfortuantely. JW
+# teh "general" lib: 
+FIND_LIBRARY(FFTW3_FFTW_LIBRARY
+  NAMES fftw3 libfftw libfftw3 libfftw3-3
+  PATHS 
+  ${FFTW3_POSSIBLE_LIBRARY_PATH}
+  )
+#MESSAGE("DBG FFTW3_FFTW_LIBRARY=${FFTW3_FFTW_LIBRARY}")
+
+FIND_LIBRARY(FFTW3_FFTWF_LIBRARY
+  NAMES fftwf3 fftw3f fftwf libfftwf libfftwf3 libfftw3f libfftw3f-3
+  PATHS 
+  ${FFTW3_POSSIBLE_LIBRARY_PATH}
+  )
+#MESSAGE("DBG FFTW3_FFTWF_LIBRARY=${FFTW3_FFTWF_LIBRARY}")
+
+FIND_LIBRARY(FFTW3_FFTWL_LIBRARY
+  NAMES fftwl3 fftw3l fftwl libfftwl libfftwl3 libfftw3l libfftw3l-3
+  PATHS 
+  ${FFTW3_POSSIBLE_LIBRARY_PATH}
+  )
+#MESSAGE("DBG FFTW3_FFTWF_LIBRARY=${FFTW3_FFTWL_LIBRARY}")
+
+
+FIND_LIBRARY(FFTW3_FFTW_SSE2_LIBRARY
+  NAMES fftw_sse2 fftw3_sse2 libfftw_sse2 libfftw3_sse2
+  PATHS 
+  ${FFTW3_POSSIBLE_LIBRARY_PATH}
+  )
+#MESSAGE("DBG FFTW3_FFTW_SSE2_LIBRARY=${FFTW3_FFTW_SSE2_LIBRARY}")
+
+FIND_LIBRARY(FFTW3_FFTWF_SSE_LIBRARY
+  NAMES fftwf_sse fftwf3_sse fftw3f_sse libfftwf_sse libfftwf3_sse libfftw3f_sse
+  PATHS 
+  ${FFTW3_POSSIBLE_LIBRARY_PATH}
+  )
+#MESSAGE("DBG FFTW3_FFTWF_SSE_LIBRARY=${FFTW3_FFTWF_SSE_LIBRARY}")
+
+
+# --------------------------------
+# select one of the above
+# default: 
+IF (FFTW3_FFTW_LIBRARY)
+  SET(FFTW3_LIBRARIES ${FFTW3_FFTW_LIBRARY})
+ENDIF (FFTW3_FFTW_LIBRARY)
+# specialized: 
+IF (FFTW3_USE_SSE2 AND FFTW3_FFTW_SSE2_LIBRARY)
+  SET(FFTW3_LIBRARIES ${FFTW3_FFTW_SSE2_LIBRARY})
+ENDIF (FFTW3_USE_SSE2 AND FFTW3_FFTW_SSE2_LIBRARY)
+
+# --------------------------------
+
+IF(FFTW3_LIBRARIES)
+  IF (FFTW3_INCLUDE_DIR)
+
+    # OK, found all we need
+    SET(FFTW3_FOUND TRUE)
+    GET_FILENAME_COMPONENT(FFTW3_LINK_DIRECTORIES ${FFTW3_LIBRARIES} PATH)
+    
+  ELSE (FFTW3_INCLUDE_DIR)
+    MESSAGE("FFTW3 include dir not found. Set FFTW3_DIR to find it.")
+  ENDIF(FFTW3_INCLUDE_DIR)
+ELSE(FFTW3_LIBRARIES)
+  MESSAGE("FFTW3 lib not found. Set FFTW3_DIR to find it.")
+ENDIF(FFTW3_LIBRARIES)
+
+
+MARK_AS_ADVANCED(
+  FFTW3_INCLUDE_DIR
+  FFTW3_LIBRARIES
+  FFTW3_FFTW_LIBRARY
+  FFTW3_FFTW_SSE2_LIBRARY
+  FFTW3_FFTWF_LIBRARY
+  FFTW3_FFTWF_SSE_LIBRARY
+  FFTW3_FFTWL_LIBRARY
+  FFTW3_LINK_DIRECTORIES
+)

+ 61 - 0
common/cmake/FindGObjectIntrospection.cmake

@@ -0,0 +1,61 @@
+# - try to find gobject-introspection
+#
+# Once done this will define
+#
+#  INTROSPECTION_FOUND - system has gobject-introspection
+#  INTROSPECTION_SCANNER - the gobject-introspection scanner, g-ir-scanner
+#  INTROSPECTION_COMPILER - the gobject-introspection compiler, g-ir-compiler
+#  INTROSPECTION_GENERATE - the gobject-introspection generate, g-ir-generate
+#  INTROSPECTION_GIRDIR
+#  INTROSPECTION_TYPELIBDIR
+#  INTROSPECTION_CFLAGS
+#  INTROSPECTION_LIBS
+#
+# Copyright (C) 2010, Pino Toscano, <pino@kde.org>
+#
+# Redistribution and use is allowed according to the terms of the BSD license.
+# For details see the accompanying COPYING-CMAKE-SCRIPTS file.
+
+macro(_GIR_GET_PKGCONFIG_VAR _outvar _varname)
+  execute_process(
+    COMMAND ${PKG_CONFIG_EXECUTABLE} --variable=${_varname} gobject-introspection-1.0
+    OUTPUT_VARIABLE _result
+    RESULT_VARIABLE _null
+  )
+
+  if (_null)
+  else()
+    string(REGEX REPLACE "[\r\n]" " " _result "${_result}")
+    string(REGEX REPLACE " +$" ""  _result "${_result}")
+    separate_arguments(_result)
+    set(${_outvar} ${_result} CACHE INTERNAL "")
+  endif()
+endmacro(_GIR_GET_PKGCONFIG_VAR)
+
+find_package(PkgConfig)
+if(PKG_CONFIG_FOUND)
+  if(PACKAGE_FIND_VERSION_COUNT GREATER 0)
+    set(_gir_version_cmp ">=${PACKAGE_FIND_VERSION}")
+  endif()
+  pkg_check_modules(_pc_gir gobject-introspection-1.0${_gir_version_cmp})
+  if(_pc_gir_FOUND)
+    set(INTROSPECTION_FOUND TRUE)
+    _gir_get_pkgconfig_var(INTROSPECTION_SCANNER "g_ir_scanner")
+    _gir_get_pkgconfig_var(INTROSPECTION_COMPILER "g_ir_compiler")
+    _gir_get_pkgconfig_var(INTROSPECTION_GENERATE "g_ir_generate")
+    _gir_get_pkgconfig_var(INTROSPECTION_GIRDIR "girdir")
+    _gir_get_pkgconfig_var(INTROSPECTION_TYPELIBDIR "typelibdir")
+    set(INTROSPECTION_CFLAGS "${_pc_gir_CFLAGS}")
+    set(INTROSPECTION_LIBS "${_pc_gir_LIBS}")
+  endif()
+endif()
+
+mark_as_advanced(
+  INTROSPECTION_SCANNER
+  INTROSPECTION_COMPILER
+  INTROSPECTION_GENERATE
+  INTROSPECTION_GIRDIR
+  INTROSPECTION_TYPELIBDIR
+  INTROSPECTION_CFLAGS
+  INTROSPECTION_LIBS
+)

+ 14 - 0
common/cmake/FindOCLFFT.cmake

@@ -0,0 +1,14 @@
+# Try to find liboclfft and clFFT.h. Once found the following variables will be
+# defined:
+#
+# OCLFFT_FOUND
+# OCLFFT_INCLUDE_DIRS
+# OCLFFT_LIBRARIES
+
+find_path(OCLFFT_INCLUDE_DIRS clFFT.h)
+find_library(OCLFFT_LIBRARIES oclfft)
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(OCLFFT DEFAULT_MSG OCLFFT_INCLUDE_DIRS OCLFFT_LIBRARIES)
+
+mark_as_advanced(OCLFFT_INCLUDE_DIRS OCLFFT_LIBRARIES)

+ 88 - 0
common/cmake/FindOpenCL.cmake

@@ -0,0 +1,88 @@
+# - 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
+              /opt/nvidia-current
+              /opt/AMDAPP/lib
+            )
+
+            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
+)
+

+ 65 - 0
common/cmake/FindVala.cmake

@@ -0,0 +1,65 @@
+##
+# Copyright 2009 Jakob Westhoff. All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+#    1. Redistributions of source code must retain the above copyright notice,
+#       this list of conditions and the following disclaimer.
+#
+#    2. Redistributions in binary form must reproduce the above copyright notice,
+#       this list of conditions and the following disclaimer in the documentation
+#       and/or other materials provided with the distribution.
+#
+# THIS SOFTWARE IS PROVIDED BY JAKOB WESTHOFF ``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 JAKOB WESTHOFF OR CONTRIBUTORS 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.
+#
+# The views and conclusions contained in the software and documentation are those
+# of the authors and should not be interpreted as representing official policies,
+# either expressed or implied, of Jakob Westhoff
+##
+
+##
+# Find module for the Vala compiler (valac)
+#
+# This module determines wheter a Vala compiler is installed on the current
+# system and where its executable is.
+#
+# Call the module using "find_package(Vala) from within your CMakeLists.txt.
+#
+# The following variables will be set after an invocation:
+#
+#  VALA_FOUND       Whether the vala compiler has been found or not
+#  VALA_EXECUTABLE  Full path to the valac executable if it has been found
+#  VALA_VERSION     Version number of the available valac
+##
+
+
+# Search for the valac executable in the usual system paths.
+find_program(VALA_EXECUTABLE
+  NAMES valac)
+
+# Handle the QUIETLY and REQUIRED arguments, which may be given to the find call.
+# Furthermore set VALA_FOUND to TRUE if Vala has been found (aka.
+# VALA_EXECUTABLE is set)
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(Vala DEFAULT_MSG VALA_EXECUTABLE)
+
+mark_as_advanced(VALA_EXECUTABLE)
+
+# Determine the valac version
+if(VALA_FOUND)
+    execute_process(COMMAND ${VALA_EXECUTABLE} "--version"
+                    OUTPUT_VARIABLE "VALA_VERSION")
+    string(REPLACE "Vala" "" "VALA_VERSION" ${VALA_VERSION})
+    string(STRIP ${VALA_VERSION} "VALA_VERSION")
+endif(VALA_FOUND)

+ 202 - 0
src/CMakeLists.txt

@@ -0,0 +1,202 @@
+cmake_minimum_required(VERSION 2.6)
+
+# --- Set sources -------------------------------------------------------------
+set(ufofilter_SRCS 
+    ufo-filter-scale.c
+    ufo-filter-lamino-bp-generic.c
+    ufo-filter-3d-edf-writer.c
+    )
+
+set(ufofilter_KERNELS
+    scale.cl
+    lamino_bp_generic.cl
+    )
+
+set(ufofilter_LIBS
+    ${UFO_LIBRARIES}
+    ${GLIB2_LIBRARIES}
+    ${GOBJECT2_LIBRARIES}
+    )
+
+set(ufofilter_HEADERS
+   lamino-filter-def.h
+)
+
+link_directories(${UFO_LIBRARY_DIRS})
+
+
+# --- Add filters that depend on other libraries ------------------------------
+#find_package(TIFF)
+#find_package(OCLFFT)
+pkg_check_modules(UCA uca)
+#pkg_check_modules(OPENCV opencv)
+
+if (UCA_INCLUDE_DIRS AND UCA_LIBRARIES)
+#   set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-cam-access.c)
+    set(ufofilter_LIBS ${ufofilter_LIBS} ${UCA_LIBRARIES})
+    include_directories(${UCA_INCLUDE_DIRS})
+endif ()
+
+# if (TIFF_FOUND)
+#     set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-reader.c)
+#    set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-writer.c)
+#    set(ufofilter_LIBS ${ufofilter_LIBS} ${TIFF_LIBRARIES})
+#    include_directories(${TIFF_INCLUDE_DIRS})
+# endif ()
+
+# if (OCLFFT_FOUND)
+#    set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-fft.c)
+#    set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-ifft.c)
+#    set(ufofilter_LIBS ${ufofilter_LIBS} ${OCLFFT_LIBRARIES})
+#    include_directories(${OCLFFT_INCLUDE_DIRS})
+# endif ()
+
+# if (OPENCV_FOUND)
+#    set(ufofilter_SRCS ${ufofilter_SRCS} ufo-filter-cv-show.c)
+#    set(ufofilter_LIBS ${ufofilter_LIBS} ${OPENCV_LIBRARIES})
+#    include_directories(${OPENCV_INCLUDE_DIRS})
+#endif ()
+
+
+# --- Add sources that belong to other libraries ------------------------------
+# set(optical_flow_lucas_kanade_misc_SRCS
+#    oflk_cl_buffer.c
+#    oflk_cl_image.c
+#    oflk_pyramid.c)
+
+
+# --- Target ------------------------------------------------------------------
+include_directories(${CMAKE_CURRENT_BINARY_DIR})
+
+get_property(LIB64 GLOBAL PROPERTY FIND_LIBRARY_USE_LIB64_PATHS)
+
+set(LIB_INSTALL_DIR "lib${LIB_SUFFIX}/ufo")
+set(shared_objects "")
+set(documented_types "")
+set(all_targets)
+
+foreach(_src ${ufofilter_SRCS})
+    # find plugin suffix
+    string(REGEX REPLACE "ufo-filter-([^ \\.]+).*" "\\1" filter "${_src}")
+
+    # build string to get miscalleanous sources
+    string(REPLACE "-" "_" _misc ${filter})
+    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 ${filter}" ON)
+
+    if(${target_option})
+        set(documented_types "${documented_types}\nufo_filter_${_misc}_get_type")
+        set(_misc "${_misc}_misc_SRCS")
+
+        string(REPLACE "-" "" _targetname ${filter})
+        set(target "ufofilter${_targetname}")
+        set(shared_objects "${shared_objects} -l${target}")
+
+        # 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 ${LIB_INSTALL_DIR}
+            LIBRARY DESTINATION ${LIB_INSTALL_DIR})
+    endif()
+endforeach()
+
+# copy kernels into $LIB_INSTALL_DIR
+foreach(_kernel ${ufofilter_KERNELS})
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${_kernel}
+        ${CMAKE_CURRENT_BINARY_DIR}/${_kernel})
+
+    install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/${_kernel}
+        DESTINATION ${LIB_INSTALL_DIR})
+endforeach()
+
+#copy aux headers into $LIB_INSTALL_DIR
+foreach(_header ${ufofilter_HEADERS})
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${_header}
+            ${CMAKE_CURRENT_BINARY_DIR}/${_header})
+		        
+    install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/${_header}
+        DESTINATION ${LIB_INSTALL_DIR})
+endforeach()
+
+
+# --- Build filter reference ---------------------------------------------------
+pkg_check_modules(GTK_DOC gtk-doc)
+if(GTK_DOC_FOUND)
+    option(WITH_GTK_DOC "Build filter documentation" ON)
+    if (WITH_GTK_DOC)
+        get_directory_property(_current_include_dirs INCLUDE_DIRECTORIES)
+
+        set(GTK_DOC_CFLAGS)
+        foreach(_incl ${_current_include_dirs})
+            set(GTK_DOC_CFLAGS "-I${_incl} ${GTK_DOC_CFLAGS}")
+        endforeach()
+
+        set(GTK_DOC_LDFLAGS)
+        foreach(_lib ${ufofilter_LIBS})
+            # check if we have an absolute library path
+            if (NOT ${_lib} MATCHES "^[/]")
+                set(GTK_DOC_LDFLAGS "-l${_lib} ${GTK_DOC_LDFLAGS}")
+            endif()
+        endforeach()
+
+        find_program(GTK_DOC_SCAN gtkdoc-scan REQUIRED)
+        find_program(GTK_DOC_SCANGOBJ gtkdoc-scangobj REQUIRED)
+        find_program(GTK_DOC_MKDB gtkdoc-mkdb REQUIRED)
+        find_program(GTK_DOC_MKHTML gtkdoc-mkhtml REQUIRED)
+
+        set(doc_base "${CMAKE_CURRENT_BINARY_DIR}/../docs")
+        set(doc_out "${doc_base}/reference")
+
+        set(_xml_doc_input)
+
+        foreach (_src ${ufofilter_SRCS})
+            string(REPLACE ".c" ".xml" _xml_doc ${_src})
+            list(APPEND _xml_doc_input "<xi:include href=\"xml/${_xml_doc}\"/>")
+        endforeach()
+
+        string(REPLACE ";" "\n" _xml_doc_input ${_xml_doc_input})
+        configure_file("${CMAKE_CURRENT_SOURCE_DIR}/UfoFilters-docs.xml.in" "${doc_out}/UfoFilters-docs.xml")
+        configure_file("${CMAKE_CURRENT_SOURCE_DIR}/scangobj.sh.in" "${doc_out}/scangobj.sh")
+        configure_file("${CMAKE_CURRENT_SOURCE_DIR}/UfoFilters.types.in" "${doc_out}/UfoFilters.types")
+
+        add_custom_command(OUTPUT ${doc_out}/UfoFilters-decl.txt
+            COMMAND ${GTK_DOC_SCAN}
+                    --module=UfoFilters
+                    --source-dir=${CMAKE_CURRENT_SOURCE_DIR}/../src
+                    DEPENDS ${ufofilter_SRCS}
+            WORKING_DIRECTORY ${doc_out})
+
+        add_custom_command(OUTPUT ${doc_out}/UfoFilters.args
+            COMMAND sh scangobj.sh
+            DEPENDS ${ufofilter_SRCS} ${doc_out}/UfoFilters-decl.txt ${all_targets}
+            WORKING_DIRECTORY ${doc_out})
+
+        add_custom_command(OUTPUT ${doc_out}/sgml.stamp
+            COMMAND ${GTK_DOC_MKDB}
+                    --module=UfoFilters
+                    --source-dir=${CMAKE_CURRENT_SOURCE_DIR}
+                    --output-format=xml
+            DEPENDS ${ufofilter_SRCS} ${doc_out}/UfoFilters.args
+            WORKING_DIRECTORY ${doc_out})
+
+        add_custom_command(OUTPUT ${doc_base}/html.stamp
+            COMMAND ${GTK_DOC_MKHTML}
+                    UfoFilters
+                    ${doc_out}/UfoFilters-docs.xml
+            DEPENDS ${doc_out}/sgml.stamp
+            WORKING_DIRECTORY ${doc_out})
+
+        add_custom_target(reference ALL DEPENDS ${doc_base}/html.stamp)
+    endif()
+endif(GTK_DOC_FOUND)

+ 38 - 0
src/UfoFilters-docs.xml.in

@@ -0,0 +1,38 @@
+<?xml version="1.0"?>
+<!DOCTYPE book PUBLIC "-//OASIS//DTD DocBook XML V4.3//EN"
+               "http://www.oasis-open.org/docbook/xml/4.3/docbookx.dtd"
+[
+  <!ENTITY % local.common.attrib "xmlns:xi  CDATA  #FIXED 'http://www.w3.org/2003/XInclude'">
+]>
+<book id="index">
+  <bookinfo>
+    <title>UFO Reference Manual</title>
+    <releaseinfo>for the UFO Filters package ${UFO_FILTERS_VERSION_STRING_LONG}.
+      The latest version of this documentation can be found on-line at
+      <ulink role="online-location"
+          url="http://ufo.kit.edu/ufo">http://ufo.kit.edu/ufo/</ulink>.
+    </releaseinfo>
+  </bookinfo>
+
+  <section>
+      <para>This documentation was generated from the source of the UFO
+          filters, thus it looks rather technical. What is most important is,
+          that the filters are instantiated with the shortened form of the name
+          as given on this page rather than the object name UfoFilterSomething.
+          A filter should not be instantiated on its own but by calling an
+          <function>get_filter</function> of an <link linkend="Ufo-Graph">UFO Graph</link>.
+      </para>
+  </section>
+
+  <chapter>
+    <title>UFO API Reference</title>
+        ${_xml_doc_input}
+
+  </chapter>
+  <index id="api-index-full">
+    <title>API Index</title>
+    <xi:include href="xml/api-index-full.xml"><xi:fallback /></xi:include>
+  </index>
+
+  <xi:include href="xml/annotation-glossary.xml"><xi:fallback /></xi:include>
+</book>

+ 1 - 0
src/UfoFilters.types.in

@@ -0,0 +1 @@
+${documented_types}

+ 50 - 0
src/lamino-filter-def.h

@@ -0,0 +1,50 @@
+#ifndef __LAMINO_FILTER_DEF_H__
+#define __LAMINO_FILTER_DEF_H__
+
+typedef struct
+{
+    // rotation angle
+    float phi;
+
+    // laminographic angle
+    float theta;
+
+    // misalignment angle
+    float psi;
+	
+    // modified lamino angle 
+    float alpha;
+
+    // rotation angular step
+    float angle_step;
+
+    // rotation matrix
+    float mat_0;
+    float mat_1;
+    float mat_2;
+    float mat_3;
+    float mat_4;
+    float mat_5;
+
+    // reconstructed volume size
+    int vol_sx;
+    int vol_sy;
+    int vol_sz;
+
+    // volume origin
+    float vol_ox;
+    float vol_oy;
+    float vol_oz;
+     
+    // projection origin
+    float proj_ox;
+    float proj_oy;
+
+    // projection sizes
+    int proj_sx;
+    int proj_sy;
+
+} CLParameters;
+
+#endif
+

+ 138 - 0
src/lamino_bp_generic.cl

@@ -0,0 +1,138 @@
+// all kernels must process volme voxelwise
+// please be careful with local and global workers
+// we need a 3D processing (not 2D)
+
+// using tests we show max volume grosse: 1024 x 1024 x 256 or 2048 x 2048 x 64
+// theory: 223 MB (CL_DEVICE_MAX_MEM_ALLOC_SIZE) x 5 cards / 4 Bytes (Float) = 1024 x 1024 x 278 voxels
+
+#include <lamino-filter-def.h>
+
+/*
+typedef struct
+{
+    // rotation angle
+    float phi;
+        
+    // laminographic angle
+    float theta;
+        
+    // misalignment angle
+    float psi;
+        
+    // modified lamino angle 
+    float alpha;
+
+    // angular step
+    float stepPhi;
+
+    // rotation matrix
+    float mat_0;
+    float mat_1;
+    float mat_2;
+    float mat_3;
+    float mat_4;
+    float mat_5;
+
+    // volume origin
+    float vol_ox;
+    float vol_oy;
+    float vol_oz;
+
+    // projection origin
+    float proj_ox;
+    float proj_oy;
+
+    // projection sizes
+    int proj_sx;
+    int proj_sy;
+
+
+} CLParameters;
+*/
+
+__kernel void lamino_bp_generic ( __global float *proj,
+				  __global float *volume,
+				  __global CLParameters *param)
+{
+    const int vX = get_global_id(0);
+    const int vY = get_global_id(1);
+    const int vZ = get_global_id(2);
+
+    const int vSX = get_global_size(0);
+    const int vSY = get_global_size(1);
+    const int vSZ = get_global_size(2);
+
+    const long int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
+    //const int idx = (vY * vSX) + vX;
+
+    float newz = (float)vZ - param->vol_oz;
+
+    float newz_matr02 = newz * param->mat_2 + param->proj_ox;
+    float newz_matr12 = newz * param->mat_5 + param->proj_oy;
+
+    /// prepare y info
+    float newy = (float)vY - param->vol_oy;
+
+    float newy_matr01 = newy * param->mat_1 + newz_matr02;
+    float newy_matr11 = newy * param->mat_4 + newz_matr12;
+
+    /// prepare x info
+    float newx = (float)vX - param->vol_ox;
+    float oldy = newx * param->mat_3 + newy_matr11;
+    float yo = floor(oldy);
+
+    float oldx = newx*param->mat_0 + newy_matr01;
+    float xo = floor(oldx);
+
+    // bilinear interpolation
+    float yf_1 = oldy - yo;
+    float yf_0 = 1.0f - yf_1;
+    float xf_1 = oldx - xo;
+    float xf_0 = 1.0f  - xf_1;
+
+    int base = (int)xo + (int)yo * param->proj_sx;
+    float result;
+// TODO: check that out of proj plain
+    result  = proj[base    ] * xf_0 * yf_0;
+    result += proj[base + 1] * xf_1 * yf_0;
+    result += proj[base + param->proj_sx    ] * xf_0 * yf_1;
+    result += proj[base + param->proj_sx + 1] * xf_1 * yf_1;
+    volume[idx] +=  result; 
+
+}
+
+__kernel void lamino_clean_vol(__global float *volume)
+{
+    const int vX = get_global_id(0);
+    const int vY = get_global_id(1);
+    const int vZ = get_global_id(2);
+
+    const int vSX = get_global_size(0);
+    const int vSY = get_global_size(1);
+    const int vSZ = get_global_size(2);
+
+    const int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
+    //const int idx = (vY * vSX) + vX;
+
+    volume[idx] = 0; 
+}
+
+__kernel void lamino_norm_vol(__global float *volume, 
+				const float factor)
+{
+    const int vX = get_global_id(0);
+    const int vY = get_global_id(1);
+    const int vZ = get_global_id(2);
+
+    const int vSX = get_global_size(0);
+    const int vSY = get_global_size(1);
+    const int vSZ = get_global_size(2);
+
+    const int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
+    //const int idx = (vY * vSX) + vX;
+
+    float val = volume[idx] * factor;
+    volume[idx] = val; 
+}
+
+

+ 35 - 0
src/reco.help

@@ -0,0 +1,35 @@
+  { "type" : "filter", "plugin" : "backproject",
+                    "properties" : {
+                        "axis-pos"  : 413.0,
+                        "angle-step" : 0.01256637,
+                        "use-texture" : true
+                    }
+                }
+
+
+
+Public Attributes
+bool 	execute
+ 	If false then the algorithm will not be executed.
+Rotation angles.
+TypeIData 	theta
+ 	Laminographic angle (between X-ray direction and rotation axis).
+TypeIData 	phi
+ 	Rotation angle.
+TypeIData 	axisMisaligment
+ 	Angle between projection of rotation axis and column of detectors.
+TypeIData 	volMisalignX
+ 	Volume rotation correction angle (about X-axis).
+TypeIData 	volMisalignY
+ 	Volume rotation correction angle (about Y-axis).
+Volume origin position.
+TypeIData 	vol_ox
+ 	X.
+TypeIData 	vol_oy
+ 	Y.
+TypeIData 	vol_oz
+ 	Z.
+\brief Coordinates of the volume origion on the projection plane
+TypeIData 	proj_ox
+ 	X.
+TypeIData 	proj_oy

+ 8 - 0
src/scale.cl

@@ -0,0 +1,8 @@
+__kernel void scale ( __global float *in, __global float *out, const int width, const float factor )
+{
+    const int idx = get_global_id(0);
+    const int idy = get_global_id(1);
+    const int index = idy * width + idx;
+    out[index] = in[index] * factor; 
+}
+

+ 1 - 0
src/scangobj.sh.in

@@ -0,0 +1 @@
+LD_LIBRARY_PATH=${CMAKE_CURRENT_BINARY_DIR} CC=gcc CFLAGS="${GTK_DOC_CFLAGS}" LDFLAGS="${GTK_DOC_LDFLAGS} -L${CMAKE_CURRENT_BINARY_DIR} -L${UFO_LIBRARY_DIRS} ${shared_objects}" gtkdoc-scangobj --module=UfoFilters

+ 223 - 0
src/ufo-filter-3d-edf-writer.c

@@ -0,0 +1,223 @@
+#include <gmodule.h>
+
+#include <ufo/ufo-filter.h>
+#include <ufo/ufo-buffer.h>
+#include <ufo/ufo-resource-manager.h>
+
+#include "ufo-filter-3d-edf-writer.h"
+
+#include <stdio.h>
+
+/**
+ * SECTION:ufo-filter-3d-edf-writer
+ * @Short_description: Stores 3d buffer as an EDF file
+ * @Title: 3d-edf-writer
+ *
+ * The writer node writes each incoming buffer as an EDF file to disk.
+ * Each file is prefixed with #UfoFilter3DEdfWriter:prefix and written into
+ * #UfoFilter3DEdfWriter:path.
+*/
+
+struct _UfoFilter3DEdfWriterPrivate {
+    gchar *path;
+    gchar *prefix;
+};
+
+GType ufo_filter_3d_edf_writer_get_type(void) G_GNUC_CONST;
+
+/* Inherit from UFO_TYPE_FILTER */
+G_DEFINE_TYPE(UfoFilter3DEdfWriter, ufo_filter_3d_edf_writer, UFO_TYPE_FILTER);
+
+#define UFO_FILTER_3D_EDF_WRITER_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_3D_EDF_WRITER, UfoFilter3DEdfWriterPrivate))
+
+enum {
+    PROP_0,
+    PROP_PATH,
+    PROP_PREFIX,
+    N_PROPERTIES
+};
+
+static GParamSpec * edfwriter3d_properties[N_PROPERTIES] = { NULL, };
+
+static gboolean filter_3d_edf_write_header(FILE * file, 
+		                           const guint32 xs,  
+					   const guint32 ys,  
+					   const guint32 zs )
+{
+/*
+   {
+   ByteOrder = LowByteFirst;
+   DataType = FloatValue;
+   Dim_1 = 1024;
+   Dim_2 = 1024;
+   Dim_3 = 256;
+   Size = 268435456;
+
+*/
+
+    const guint32 headsz = 512;	
+    gchar  * header  =  g_strnfill (headsz, ' ');
+
+    sprintf(header, "{\nByteOrder = LowByteFirst;\nDataType = FloatValue;\nDim_1 = %i;\nDim_2 = %i;\nDim_3 = %i;\nSize = %li;\n",
+		       xs, ys, zs, xs*ys*zs*sizeof(float));
+
+    header[510] = '}';
+    header[511] = '\x0A';
+
+    fwrite(header, sizeof(char), headsz, file);
+    return TRUE;	
+}
+
+static gboolean filter_3d_edf_write_body(FILE * file, 
+		                         float * data, 
+					 const guint32 sz)
+{
+     fwrite(data, sizeof(float), sz, file);
+     return TRUE;
+}
+
+
+static void ufo_filter_3d_edf_writer_process(UfoFilter *self)
+{
+    g_return_if_fail(UFO_IS_FILTER(self));
+    UfoFilter3DEdfWriterPrivate *priv = UFO_FILTER_3D_EDF_WRITER_GET_PRIVATE(self);
+    UfoChannel *input_channel = ufo_filter_get_input_channel(self);
+    UfoBuffer *input = ufo_channel_get_input_buffer(input_channel);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(self);
+    g_message("ufo_filter_3d_edf_writer_process");
+
+    guint num_dims = 0;
+    guint *dim_size = NULL;
+    guint counter = 0;
+    GString *filename = g_string_new("");
+
+    while (input != NULL) 
+    {
+        ufo_buffer_get_dimensions(input, &num_dims, &dim_size);
+        g_assert(num_dims == 3);
+
+	const guint xs = dim_size[0];
+        const guint ys = dim_size[1];
+	const guint zs = dim_size[2];
+
+        float *data = ufo_buffer_get_host_array(input, command_queue);
+
+        g_string_printf(filename, "%s/%s%05i.edf", priv->path, priv->prefix, counter++);
+        //if (!filter_write_tiff(data, filename->str, width, height))
+	//				            g_message("something went wrong")
+	// write EDF header
+	// write raw data
+    // temporal saving unless 3d writer is not ready
+          FILE * idfile = fopen(filename->str,"wb");
+
+	  if(!filter_3d_edf_write_header(idfile, xs, ys, zs))
+             g_message("cannot write edf header");
+
+	  if(!filter_3d_edf_write_body(idfile, data, xs*ys*zs))
+              g_message("cannot write data");
+
+          fclose(idfile);
+    
+        ufo_channel_finalize_input_buffer(input_channel, input);
+        input = ufo_channel_get_input_buffer(input_channel);
+
+    }
+
+    g_string_free(filename, TRUE);
+    g_free(dim_size);
+
+}
+
+
+static void ufo_filter_3d_edf_writer_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilter3DEdfWriter *filter = UFO_FILTER_3D_EDF_WRITER(object);
+
+    switch (property_id) 
+    {
+        case PROP_PATH:
+            g_free(filter->priv->path);
+            filter->priv->path = g_strdup(g_value_get_string(value));
+            break;
+        case PROP_PREFIX:
+            g_free(filter->priv->prefix);
+            filter->priv->prefix = g_strdup(g_value_get_string(value));
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+
+static void ufo_filter_3d_edf_writer_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilter3DEdfWriter *filter = UFO_FILTER_3D_EDF_WRITER(object);
+
+   switch (property_id) 
+   {
+        case PROP_PATH:
+            g_value_set_string(value, filter->priv->path);
+            break;
+        case PROP_PREFIX:
+            g_value_set_string(value, filter->priv->prefix);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_3d_edf_writer_class_init(UfoFilter3DEdfWriterClass *klass)
+{
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_3d_edf_writer_set_property;
+    gobject_class->get_property = ufo_filter_3d_edf_writer_get_property;
+    filter_class->process = ufo_filter_3d_edf_writer_process;
+
+    edfwriter3d_properties[PROP_PREFIX] =
+        g_param_spec_string("prefix",
+		            "Filename prefix",
+  	                    "Prefix of output filename.",
+		            "",
+		            G_PARAM_READWRITE);
+
+   edfwriter3d_properties[PROP_PATH] =
+        g_param_spec_string("path",
+ 	                    "File path",
+		            "Path where to store files.",
+	                     ".",
+	 	             G_PARAM_READWRITE);
+
+   g_object_class_install_property(gobject_class, PROP_PATH, edfwriter3d_properties[PROP_PATH]);
+   g_object_class_install_property(gobject_class, PROP_PREFIX, edfwriter3d_properties[PROP_PREFIX]);
+
+   g_type_class_add_private(gobject_class, sizeof(UfoFilter3DEdfWriterPrivate));
+
+}
+
+static void ufo_filter_3d_edf_writer_init(UfoFilter3DEdfWriter *self)
+{
+    self->priv = UFO_FILTER_3D_EDF_WRITER_GET_PRIVATE(self);
+    self->priv->path = g_strdup(".");
+    self->priv->prefix = NULL;
+
+    ufo_filter_register_input(UFO_FILTER(self), "volume", 3);
+}
+
+G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
+{
+    return g_object_new(UFO_TYPE_FILTER_3D_EDF_WRITER, NULL);
+}
+
+
+
+

+ 36 - 0
src/ufo-filter-3d-edf-writer.h

@@ -0,0 +1,36 @@
+#ifndef __UFO_FILTER_3D_EDF_WRITER_H__
+#define __UFO_FILTER_3D_EDF_WRITER_H__ 
+
+#include <glib.h>
+#include <ufo/ufo-filter.h>
+
+#define UFO_TYPE_FILTER_3D_EDF_WRITER             (ufo_filter_3d_edf_writer_get_type())
+#define UFO_FILTER_3D_EDF_WRITER(obj)             (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_FILTER_3D_EDF_WRITER, UfoFilter3DEdfWriter))
+#define UFO_IS_FILTER_3D_EDF_WRITER(obj)          (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_FILTER_3D_EDF_WRITER))
+#define UFO_FILTER_3D_EDF_WRITER_CLASS(klass)     (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_FILTER_3D_EDF_WRITER, UfoFilter3DEdfWriterClass))
+#define UFO_IS_FILTER_3D_EDF_WRITER_CLASS(klass)  (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_FILTER_3D_EDF_WRITER))
+#define UFO_FILTER_3D_EDF_WRITER_GET_CLASS(obj)   (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_FILTER_3D_EDF_WRITER, UfoFilter3DEdfWriterClass))
+
+typedef struct _UfoFilter3DEdfWriter           UfoFilter3DEdfWriter;
+typedef struct _UfoFilter3DEdfWriterClass      UfoFilter3DEdfWriterClass;
+typedef struct _UfoFilter3DEdfWriterPrivate    UfoFilter3DEdfWriterPrivate;
+
+struct _UfoFilter3DEdfWriter {
+    /*< private >*/
+    UfoFilter parent_instance;
+    UfoFilter3DEdfWriterPrivate *priv;
+};
+
+/**
+  * UfoFilter3DEdfWriterClass:
+  *
+ * #UfoFilter3DEdfWriter class
+*/
+struct _UfoFilter3DEdfWriterClass {
+    /*< private >*/
+    UfoFilterClass parent_class;
+};
+
+GType ufo_filter_3d_edf_writer_get_type(void);
+
+#endif

+ 518 - 0
src/ufo-filter-lamino-bp-generic.c

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

+ 37 - 0
src/ufo-filter-lamino-bp-generic.h

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

+ 198 - 0
src/ufo-filter-scale.c

@@ -0,0 +1,198 @@
+#include <gmodule.h>
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <ufo/ufo-filter.h>
+// #include <ufo/ufo-element.h>
+#include <ufo/ufo-buffer.h>
+#include <ufo/ufo-resource-manager.h>
+
+#include "ufo-filter-scale.h"
+
+/**
+ * SECTION:ufo-filter-scale
+ * @Short_description: Scale image values 
+ * @Title: scale
+ *
+ * Scale input image values. The output
+ * is a new image.
+ * #UfoFilterScale: params.
+*/
+
+struct _UfoFilterScalePrivate {
+    float scale;
+    cl_kernel kernel;
+};
+
+GType ufo_filter_scale_get_type(void) G_GNUC_CONST;
+
+/* Inherit from UFO_TYPE_FILTER */
+G_DEFINE_TYPE(UfoFilterScale, ufo_filter_scale, UFO_TYPE_FILTER);
+
+#define UFO_FILTER_SCALE_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_SCALE, UfoFilterScalePrivate))
+
+enum {
+    PROP_0 =0,
+    PROP_SCALE,
+    N_PROPERTIES
+};
+
+static GParamSpec *scale_properties[N_PROPERTIES] = { NULL, };
+
+
+/* 
+ * virtual methods 
+ */
+static void ufo_filter_scale_initialize(UfoFilter *filter)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+    self->priv->kernel = NULL;
+
+    ufo_resource_manager_add_program(manager, "scale.cl", NULL,  &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+        return;
+    }
+
+    self->priv->kernel = ufo_resource_manager_get_kernel(manager, "scale", &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+}
+
+static void ufo_filter_scale_process(UfoFilter *filter)
+{
+    g_return_if_fail(UFO_IS_FILTER(filter));
+    UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+
+    UfoChannel *input_channel = ufo_filter_get_input_channel(filter);
+    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
+
+    UfoBuffer *input  = ufo_channel_get_input_buffer(input_channel);
+
+    guint num_dims = 0;
+    guint *dim_size = NULL;
+    ufo_buffer_get_dimensions(input, &num_dims, &dim_size);
+    ufo_channel_allocate_output_buffers(output_channel, num_dims, dim_size);
+    
+    size_t global_work_size[2] = { (size_t) dim_size[0], (size_t) dim_size[1] };
+    int width = dim_size[0];
+
+    float scale = (float) self->priv->scale;
+
+    cl_kernel kernel = self->priv->kernel;
+
+
+    while(input != NULL)
+    {
+         UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
+         cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array(input,  command_queue);
+         cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output, 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(int), &width));
+         CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(float), &scale));
+	
+         cl_event event;	 
+         CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
+					      2, NULL, global_work_size, NULL,
+					      0, NULL, &event));
+
+     // ufo_buffer_attach_event(output, event);
+
+         ufo_filter_account_gpu_time(filter, (void **) &event);
+
+         g_message("ufo-filter-scale: processing is completed");
+     
+         ufo_channel_finalize_input_buffer(input_channel, input);
+         ufo_channel_finalize_output_buffer(output_channel, output);
+         input = ufo_channel_get_input_buffer(input_channel);
+    
+    }
+    ufo_channel_finish(output_channel);
+    g_free(dim_size);
+
+}
+
+static void ufo_filter_scale_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            self->priv->scale = (float) g_value_get_double(value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            g_value_set_double(value, (double) self->priv->scale);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_class_init(UfoFilterScaleClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_scale_set_property;
+    gobject_class->get_property = ufo_filter_scale_get_property;
+    filter_class->initialize = ufo_filter_scale_initialize;
+    filter_class->process = ufo_filter_scale_process;
+
+    scale_properties[PROP_SCALE] = 
+      g_param_spec_double("scale",
+        "Scale",
+        "Scale for each pixel",
+        -5.0,   /* minimum */
+         10.0,   /* maximum */
+         1.0,   /* default */
+        G_PARAM_READWRITE);
+
+    g_object_class_install_property(gobject_class, PROP_SCALE, scale_properties[PROP_SCALE]);
+
+    /* install private data */
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterScalePrivate));
+}
+
+static void ufo_filter_scale_init(UfoFilterScale *self)
+{
+    UfoFilterScalePrivate *priv = self->priv = UFO_FILTER_SCALE_GET_PRIVATE(self);
+    priv->scale = 1.0;
+    priv->kernel = NULL;
+    
+    ufo_filter_register_input (UFO_FILTER(self), "image", 2);
+    ufo_filter_register_output(UFO_FILTER(self), "image", 2);
+}
+
+G_MODULE_EXPORT  *ufo_filter_plugin_new(void) 
+{
+    return g_object_new(UFO_TYPE_FILTER_SCALE, NULL);
+}

+ 39 - 0
src/ufo-filter-scale.h

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

+ 212 - 0
src/ufo-filter-scale.local

@@ -0,0 +1,212 @@
+#include <gmodule.h>
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <ufo/ufo-filter.h>
+// #include <ufo/ufo-element.h>
+#include <ufo/ufo-buffer.h>
+#include <ufo/ufo-resource-manager.h>
+
+#include "ufo-filter-scale.h"
+
+/**
+ * SECTION:ufo-filter-scale
+ * @Short_description: Scale image values 
+ * @Title: scale
+ *
+ * Scale input image values. The output
+ * is a new image.
+ * #UfoFilterScale: params.
+*/
+
+struct _UfoFilterScalePrivate {
+    float scale;
+    cl_kernel kernel;
+};
+
+GType ufo_filter_scale_get_type(void) G_GNUC_CONST;
+
+/* Inherit from UFO_TYPE_FILTER */
+G_DEFINE_TYPE(UfoFilterScale, ufo_filter_scale, UFO_TYPE_FILTER);
+
+#define UFO_FILTER_SCALE_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_SCALE, UfoFilterScalePrivate))
+
+enum {
+    PROP_0 =0,
+    PROP_SCALE,
+    N_PROPERTIES
+};
+
+static GParamSpec *scale_properties[N_PROPERTIES] = { NULL, };
+
+
+/* 
+ * virtual methods 
+ */
+static void ufo_filter_scale_initialize(UfoFilter *filter)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+    UfoResourceManager *manager = ufo_resource_manager();
+    GError *error = NULL;
+    self->priv->kernel = NULL;
+
+    ufo_resource_manager_add_program(manager, "scale.cl", NULL,  &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+        return;
+    }
+
+    self->priv->kernel = ufo_resource_manager_get_kernel(manager, "scale", &error);
+    if (error != NULL) {
+        g_warning("%s", error->message);
+        g_error_free(error);
+    }
+}
+
+static void ufo_filter_scale_process(UfoFilter *filter)
+{
+    g_return_if_fail(UFO_IS_FILTER(filter));
+   UfoFilterScale *self = UFO_FILTER_SCALE(filter);
+    UfoChannel *input_channel = ufo_filter_get_input_channel(filter);
+    UfoChannel *output_channel = ufo_filter_get_output_channel(filter);
+    cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(filter);
+
+     UfoBuffer *input  = ufo_channel_get_input_buffer(input_channel);
+     ufo_channel_allocate_output_buffers_like(output_channel, input);
+  
+     const gint32 num_elements = ufo_buffer_get_size(input) / sizeof(float);
+     float scale = (float) self->priv->scale;
+     while(input !=NULL)
+     {
+         float *in_data = ufo_buffer_get_host_array(input, command_queue);
+
+	 UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
+	  
+	 /* This avoids an unneccessary GPU-to-host transfer */
+	  ufo_buffer_invalidate_gpu_data(output);
+
+         float *out_data = ufo_buffer_get_host_array(output, command_queue);
+
+	 for (int i = 0; i < num_elements; i++) 
+		out_data[i] = scale * in_data[i];				         
+
+         ufo_channel_finalize_input_buffer(input_channel, input);
+         ufo_channel_finalize_output_buffer(output_channel, output);
+	 input = ufo_channel_get_input_buffer(input_channel); // read next
+	 g_message("ufo-filter-scale: processing is completed");     
+     }	     
+
+/*     UfoBuffer *oimage = NULL;
+
+
+    gint32 width, height;
+    UfoBuffer *buffer = (UfoBuffer *) g_async_queue_pop(input_queue);
+    while (!ufo_buffer_is_finished(buffer)) {
+        if (self->priv->kernel != NULL) {
+            float scale
+		    = (float) self->priv->scale;
+            size_t global_work_size[2];
+
+            ufo_buffer_get_dimensions(buffer, &width, &height);
+            global_work_size[0] = width;
+            global_work_size[1] = height;
+
+            cl_mem buffer_mem = (cl_mem) ufo_buffer_get_gpu_data(buffer, command_queue);
+            cl_int err = CL_SUCCESS;
+
+            err = clSetKernelArg(self->priv->kernel, 0, sizeof(float), &scale);
+            err = clSetKernelArg(self->priv->kernel, 1, sizeof(cl_mem), (void *) &buffer_mem);
+            err = clEnqueueNDRangeKernel(command_queue,
+                self->priv->kernel,
+                2, NULL, global_work_size, NULL,
+                0, NULL, &event);
+
+            ufo_filter_account_gpu_time(filter, (void **) &event);
+        }
+        g_async_queue_push(output_queue, buffer);
+        buffer = (UfoBuffer *) g_async_queue_pop(input_queue);
+    }
+   // g_message("ufo-filter-scale: 0s/%fs", ufo_filter_get_gpu_time(filter));
+   // g_async_queue_push(output_queue, buffer);
+*/
+    ufo_channel_finish(output_channel);
+
+}
+
+static void ufo_filter_scale_set_property(GObject *object,
+    guint           property_id,
+    const GValue    *value,
+    GParamSpec      *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            self->priv->scale = (float) g_value_get_double(value);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_get_property(GObject *object,
+    guint       property_id,
+    GValue      *value,
+    GParamSpec  *pspec)
+{
+    UfoFilterScale *self = UFO_FILTER_SCALE(object);
+
+    switch (property_id) {
+        case PROP_SCALE:
+            g_value_set_double(value, (double) self->priv->scale);
+            break;
+        default:
+            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
+            break;
+    }
+}
+
+static void ufo_filter_scale_class_init(UfoFilterScaleClass *klass)
+{
+    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
+    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
+
+    gobject_class->set_property = ufo_filter_scale_set_property;
+    gobject_class->get_property = ufo_filter_scale_get_property;
+    filter_class->initialize = ufo_filter_scale_initialize;
+    filter_class->process = ufo_filter_scale_process;
+
+    scale_properties[PROP_SCALE] = 
+      g_param_spec_double("scale",
+        "Scale",
+        "Scale for each pixel",
+        -5.0,   /* minimum */
+         10.0,   /* maximum */
+         1.0,   /* default */
+        G_PARAM_READWRITE);
+
+    g_object_class_install_property(gobject_class, PROP_SCALE, scale_properties[PROP_SCALE]);
+
+    /* install private data */
+    g_type_class_add_private(gobject_class, sizeof(UfoFilterScalePrivate));
+}
+
+static void ufo_filter_scale_init(UfoFilterScale *self)
+{
+    UfoFilterScalePrivate *priv = self->priv = UFO_FILTER_SCALE_GET_PRIVATE(self);
+    priv->scale = 1.0;
+    priv->kernel = NULL;
+    
+    ufo_filter_register_input (UFO_FILTER(self), "image", 2);
+    ufo_filter_register_output(UFO_FILTER(self), "image", 2);
+}
+
+G_MODULE_EXPORT  *ufo_filter_plugin_new(void) 
+{
+    return g_object_new(UFO_TYPE_FILTER_SCALE, NULL);
+}