diff options
| author | Picca Frédéric-Emmanuel <picca@debian.org> | 2016-11-26 20:11:17 +0100 |
|---|---|---|
| committer | Picca Frédéric-Emmanuel <picca@debian.org> | 2016-11-26 20:11:17 +0100 |
| commit | c695dd103e599f4718d6e2e0645ec5df1058ea81 (patch) | |
| tree | 0c66319e949d62d9094b38e5b82cf756882c2e74 | |
| parent | 95f0ec29fb3953afb30348bef9ed452da0376afa (diff) | |
| parent | b2f88bd4a239fba2a9016628c13645fb77305800 (diff) | |
Merge tag 'v0.12.0+dfsg1'
66 files changed, 1871 insertions, 851 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 7fdd023..15bd539 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -17,8 +17,8 @@ project(ufo C) set(TARNAME "ufo-filters") set(UFO_FILTERS_VERSION_MAJOR "0") -set(UFO_FILTERS_VERSION_MINOR "11") -set(UFO_FILTERS_VERSION_PATCH "1") +set(UFO_FILTERS_VERSION_MINOR "12") +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}") @@ -31,7 +31,7 @@ list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/common/cmake") include(GNUInstallDirs) include(PkgConfigVars) -set(PKG_UFO_CORE_MIN_REQUIRED "0.6") +set(PKG_UFO_CORE_MIN_REQUIRED "0.12") # Backprojection burst mode, must be one of 1, 2, 4, 8, 16 set(BP_BURST "16" CACHE STRING "Number of projections processed in one pass") @@ -1,3 +1,33 @@ +Version 0.12.0 +============== + +Enhancements: + +- Fortify source and enable large file support +- Re-arrange filter documentation + +Fixes: + +- Fix #127: use enums where possible +- Document the filter task +- Fix potential errors found with static analysis +- stdin: use gsize to avoid LFS problems +- dfi-sinc: do not call exit() +- raw/read: fix type translation for raw-offset + +Breaks: + +- metaballs: create filled balls rather than circles +- metaballs: remove run-infinitely and fps props +- filter: use enum instead of type-unsafe string +- loop: rename ::count to ::number + +New filters: + +- Add binarization filter +- Add basic segmentation filter + + Version 0.11.1 ============== diff --git a/docs/CMakeLists.txt b/docs/CMakeLists.txt index 1ccfca0..63dd582 100644 --- a/docs/CMakeLists.txt +++ b/docs/CMakeLists.txt @@ -1,5 +1,3 @@ -cmake_minimum_required(VERSION 2.6) - find_program(SPHINX sphinx-build PATHS /usr/local/bin /usr/bin) mark_as_advanced(SPHINX) diff --git a/docs/filters.rst b/docs/filters.rst index ff60e6a..f981ee2 100644 --- a/docs/filters.rst +++ b/docs/filters.rst @@ -4,30 +4,20 @@ Filters Filters transform data and have at least one input and one output. -Point-based filters -=================== +Point-based transformation +========================== -Transposition -------------- - -.. gobj:class:: transpose - - Transpose images from (x, y) to (y, x). +Binarization +------------ +.. gobj:class:: binarize -Flipping --------- + Binarizes an image. -.. gobj:class:: flip - - Flips images vertically or horizontally. - - .. gobj:prop:: direction:string - - Can be either `horizontal` or `vertical` and denotes the direction along - with the image is flipped. + .. gobj:prop:: threshold:float + Any values above the threshold are set to one all others to zero. Clipping -------- @@ -40,57 +30,10 @@ Clipping Minimum value, all values lower than `min` are set to `min`. - .. gobj:prop:: max:flaot + .. gobj:prop:: max:float Maximum value, all values higher than `max` are set to `max`. - -Binning -------- - -.. gobj:class:: bin - - Bin a square of pixels by summing their values. - - .. gobj:prop:: size:uint - - Number of pixels in one direction to bin to a single pixel value. - - -Averaging ---------- - -.. gobj:class:: average - - Read in full data stream and generate an averaged output. - - .. gobj:prop:: number:int - - Number of averaged images to output. By default one image is generated. - - -Flat-field correction ---------------------- - -.. gobj:class:: flat-field-correct - - Computes the flat field correction using three data streams: - - 1. Projection data on input 0 - 2. Dark field data on input 1 - 3. Flat field data on input 2 - - - .. gobj:prop:: absorption-correction:boolean - - If *TRUE*, compute the negative natural logarithm of the - flat-corrected data. - - .. gobj:prop:: fix-nan-and-inf:boolean - - If *TRUE*, replace all resulting NANs and INFs with zeros. - - Arithmetic expressions ---------------------- @@ -115,6 +58,7 @@ Arithmetic expressions .. _generic-opencl-ref: + Generic OpenCL -------------- @@ -142,221 +86,162 @@ Generic OpenCL Number of dimensions the kernel works on. Must be in [1, 3]. -Complex filters -=============== - -Median ------- - -.. gobj:class:: median-filter - - Filters input with a simple median. - - .. gobj:prop:: size:int - - Odd-numbered size of the neighbouring window. - - -Edge detection --------------- - -.. gobj:class:: detect-edge - - Detect edges by computing the power gradient image using different edge - filters. - - .. gobj:prop:: type:string - - Edge filter (or operator) which is one of ``sobel``, ``laplace`` and - ``prewitt``. By default, the ``sobel`` operator is used. - - -Sinogram transposition ----------------------- - -.. gobj:class:: transpose-projections - - Read a stream of two-dimensional projections and output a stream of - transposed sinograms. :gobj:prop:`num-projections` *must* be set to the - number of incoming projections to allocate enough memory. - - .. gobj:prop:: number:int - - Number of projections. - - .. Warning:: - - This is a memory intensive task and can easily exhaust your - system memory. Make sure you have enough memory, otherwise the process - will be killed. +Spatial transformation +====================== +Transposition +------------- -Tomographic backprojection --------------------------- +.. gobj:class:: transpose -.. gobj:class:: backproject + Transpose images from (x, y) to (y, x). - Computes the backprojection for a single sinogram. - .. gobj:prop:: axis-pos:float +Flipping +-------- - Position of the rotation axis in horizontal pixel dimension of a - sinogram or projection. If not given, the center of the sinogram is - assumed. +.. gobj:class:: flip - .. gobj:prop:: angle-step:float + Flips images vertically or horizontally. - Angle step increment in radians. If not given, pi divided by height - of input sinogram is assumed. + .. gobj:prop:: direction:string - .. gobj:prop:: angle-offset:float + Can be either `horizontal` or `vertical` and denotes the direction along + with the image is flipped. - Constant angle offset in radians. This determines effectively the - starting angle. - .. gobj:prop:: mode:enum +Binning +------- - Reconstruction mode which can be either ``nearest`` or ``texture``. +.. gobj:class:: bin - .. gobj:prop:: roi-x:int + Bin a square of pixels by summing their values. - Horizontal coordinate of the start of the ROI. By default 0. + .. gobj:prop:: size:uint - .. gobj:prop:: roi-y:int + Number of pixels in one direction to bin to a single pixel value. - Vertical coordinate of the start of the ROI. By default 0. - .. gobj:prop:: roi-width:int +Rescaling +--------- - Width of the region of interest. The default value of 0 denotes full - width. +.. gobj:class:: rescale - .. gobj:prop:: roi-height:int + Rescale input data by a fixed :gobj:prop:`factor`. - Height of the region of interest. The default value of 0 denotes full - height. + .. gobj:prop:: factor:int + Fixed factor for scaling the input in both directions. -Forward projection ------------------- + .. gobj:prop:: x-factor:int -.. gobj:class:: forwardproject + Fixed factor for scaling the input width. - Computes the forward projection of slices into sinograms. + .. gobj:prop:: y-factor:int - .. gobj:prop:: number:int + Fixed factor for scaling the input height. - Number of final 1D projections, that means height of the sinogram. + .. gobj:prop:: width:int - .. gobj:prop:: angle-step:float + Fixed width, disabling scalar rescaling. - Angular step between two adjacent projections. If not changed, it is - simply pi divided by :gobj:prop:`num-projections`. + .. gobj:prop:: height:int + Fixed height, disabling scalar rescaling. -Laminographic backprojection ----------------------------- + .. gobj:prop:: interpolation:string -.. gobj:class:: lamino-backproject + Interpolation method used for rescaling. - Backprojects parallel beam computed laminography projection-by-projection - into a 3D volume. +Padding +------- - .. gobj:prop:: x-region:GValueArray +.. gobj:class:: pad - X region for reconstruction as (from, to, step). + Pad an image to some extent with specific behavior for pixels falling + outside the original image. - .. gobj:prop:: y-region:GValueArray + .. gobj:prop:: x:int - Y region for reconstruction as (from, to, step). + Horizontal coordinate in the output image which will contain the first + input column. - .. gobj:prop:: z:float + .. gobj:prop:: y:int - Z coordinate of the reconstructed slice. + Vertical coordinate in the output image which will contain the first + input row. - .. gobj:prop:: region:GValueArray + .. gobj:prop:: width:int - Region for the parameter along z-axis as (from, to, step). + Width of the padded image. - .. gobj:prop:: projection-offset:GValueArray + .. gobj:prop:: height:int - Offset to projection data as (x, y) for the case input data is cropped - to the necessary range of interest. + Height of the padded image. - .. gobj:prop:: center:GValueArray + .. gobj:prop:: addressing-mode:string - Center of the volume with respect to projections (x, y), (rotation - axes). + Addressing mode specifies the behavior for pixels falling outside the + original image. See OpenCL sampler_t documentation for more information. - .. gobj:prop:: overall-angle:float - Angle covered by all projections (can be negative for negative steps in - case only num-projections is specified) +Cropping +-------- - .. gobj:prop:: num-projections:uint +.. gobj:class:: crop - Number of projections. + Crop a region of interest from two-dimensional input. If the region is + (partially) outside the input, only accessible data will be copied. - .. gobj:prop:: tomo-angle:float + .. gobj:prop:: x:int - Tomographic rotation angle in radians (used for acquiring projections). + Horizontal coordinate from where to start the ROI. - .. gobj:prop:: lamino-angle:float + .. gobj:prop:: y:int - Absolute laminogrpahic angle in radians determining the sample tilt. + Vertical coordinate from where to start the ROI. - .. gobj:prop:: roll-angle:float + .. gobj:prop:: width:int - Sample angular misalignment to the side (roll) in radians (CW is - positive). + Width of the region of interest. - .. gobj:prop:: parameter:string + .. gobj:prop:: height:int - Which paramter will be varied along the z-axis, from "z", "x-center", - "lamino-angle", "roll-angle". + Height of the region of interest. + .. gobj:prop:: from-center:boolean -Phase retrieval ---------------- + Start cropping from the center outwards. -.. gobj:class:: retrieve-phase - Computes and applies a fourier filter to correct phase-shifted data. - Expects frequencies as an input and produces frequencies as an output. +Filters +======= - .. gobj:prop:: method:string +Median +------ - Retrieval method which is one of ``tie``, ``ctf``, ``ctfhalfsin``, - ``qp``, ``qphalfsine`` or ``qp2``. +.. gobj:class:: median-filter - .. gobj:prop:: energy:float + Filters input with a simple median. - Energy in keV. + .. gobj:prop:: size:int - .. gobj:prop:: distance:float + Odd-numbered size of the neighbouring window. - Distance in meter. - .. gobj:prop:: pixel-size:float +Edge detection +-------------- - Pixel size in meter. +.. gobj:class:: detect-edge - .. gobj:prop:: regularization-rate:float + Detect edges by computing the power gradient image using different edge + filters. - Regularization parameter is log10 of the constant to be added to the - denominator to regularize the singularity at zero frequency: 1/sin(x) -> - 1/(sin(x)+10^-RegPar). - - Typical values [2, 3]. + .. gobj:prop:: type:string - .. gobj:prop:: thresholding-rate:float + Edge filter (or operator) which is one of ``sobel``, ``laplace`` and + ``prewitt``. By default, the ``sobel`` operator is used. - Parameter for Quasiparticle phase retrieval which defines the width of - the rings to be cropped around the zero crossing of the CTF denominator - in Fourier space. - - Typical values in [0.01, 0.1], ``qp`` retrieval is rather independent of - cropping width. Gaussian blur @@ -375,81 +260,41 @@ Gaussian blur Sigma of the kernel. -Padding -------- -.. gobj:class:: pad +Stream transformations +====================== - Pad an image to some extent with specific behavior for pixels falling - outside the original image. - - .. gobj:prop:: x:int - - Horizontal coordinate in the output image which will contain the first - input column. +Averaging +--------- - .. gobj:prop:: y:int +.. gobj:class:: average - Vertical coordinate in the output image which will contain the first - input row. +Read in full data stream and generate an averaged output. - .. gobj:prop:: width:int +.. gobj:prop:: number:int - Width of the padded image. +Number of averaged images to output. By default one image is generated. - .. gobj:prop:: height:int - Height of the padded image. +Slicing +------- - .. gobj:prop:: addressing-mode:string +.. gobj:class:: slice - Addressing mode specifies the behavior for pixels falling outside the - original image. See OpenCL sampler_t documentation for more information. + Slices a three-dimensional input buffer to two-dimensional slices. -Cropping +Stacking -------- -.. gobj:class:: crop - - Crop a region of interest from two-dimensional input. If the region is - (partially) outside the input, only accessible data will be copied. - - .. gobj:prop:: x:int - - Horizontal coordinate from where to start the ROI. - - .. gobj:prop:: y:int - - Vertical coordinate from where to start the ROI. - - .. gobj:prop:: width:int - - Width of the region of interest. - - .. gobj:prop:: height:int - - Height of the region of interest. - - .. gobj:prop:: from-center:boolean - - Start cropping from the center outwards. - - -Rescaling ---------- - -.. gobj:class:: rescale - - Rescale input data by a fixed :gobj:prop:`factor`. - - .. gobj:prop:: factor:int +.. gobj:class:: stack - Fixed factor for scaling the input. + Symmetrical to the slice filter, the stack filter stacks two-dimensional + input. - .. gobj:prop:: interpolation:string + .. gobj:prop:: number:int - Interpolation method used for rescaling. + Number of items, i.e. the length of the third dimension. Merging @@ -480,6 +325,9 @@ Slice mapping mapper, warnings are issued. +Fourier domain +============== + Fast Fourier transform ---------------------- @@ -544,6 +392,42 @@ Fast Fourier transform Height to crop output. +Frequency filtering +------------------- + +.. gobj:class:: filter + + Computes a frequency filter function and multiplies it with its input, + effectively attenuating certain frequencies. + + .. gobj:prop:: filter_:string + + Any of ``ramp``, ``ramp-fromreal``, ``butterworth``, ``faris-byer`` and + ``hamming``. The default filter is ``ramp-fromreal`` which computes a + correct ramp filter avoiding offset issues encountered with naive + implementations. + + .. gobj:prop:: scale:float + + Arbitrary scale that is multiplied to each frequency component. + + .. gobj:prop:: cutoff:float + + Cutoff frequency of the Butterworth filter. + + .. gobj:prop:: order:int + + Order of the Butterworth filter. + + .. gobj:prop:: tau:float + + Tau parameter of Faris-Byer filter. + + .. gobj:prop:: theta:float + + Theta parameter of Faris-Byer filter. + + 1D stripe filtering ------------------- @@ -560,8 +444,221 @@ Fast Fourier transform gaussian. +Reconstruction +============== + +Flat-field correction +--------------------- + +.. gobj:class:: flat-field-correct + + Computes the flat field correction using three data streams: + + 1. Projection data on input 0 + 2. Dark field data on input 1 + 3. Flat field data on input 2 + + + .. gobj:prop:: absorption-correction:boolean + + If *TRUE*, compute the negative natural logarithm of the + flat-corrected data. + + .. gobj:prop:: fix-nan-and-inf:boolean + + If *TRUE*, replace all resulting NANs and INFs with zeros. + + +Sinogram transposition +---------------------- + +.. gobj:class:: transpose-projections + + Read a stream of two-dimensional projections and output a stream of + transposed sinograms. :gobj:prop:`num-projections` *must* be set to the + number of incoming projections to allocate enough memory. + + .. gobj:prop:: number:int + + Number of projections. + + .. Warning:: + + This is a memory intensive task and can easily exhaust your + system memory. Make sure you have enough memory, otherwise the process + will be killed. + + +Tomographic backprojection +-------------------------- + +.. gobj:class:: backproject + + Computes the backprojection for a single sinogram. + + .. gobj:prop:: axis-pos:float + + Position of the rotation axis in horizontal pixel dimension of a + sinogram or projection. If not given, the center of the sinogram is + assumed. + + .. gobj:prop:: angle-step:float + + Angle step increment in radians. If not given, pi divided by height + of input sinogram is assumed. + + .. gobj:prop:: angle-offset:float + + Constant angle offset in radians. This determines effectively the + starting angle. + + .. gobj:prop:: mode:enum + + Reconstruction mode which can be either ``nearest`` or ``texture``. + + .. gobj:prop:: roi-x:int + + Horizontal coordinate of the start of the ROI. By default 0. + + .. gobj:prop:: roi-y:int + + Vertical coordinate of the start of the ROI. By default 0. + + .. gobj:prop:: roi-width:int + + Width of the region of interest. The default value of 0 denotes full + width. + + .. gobj:prop:: roi-height:int + + Height of the region of interest. The default value of 0 denotes full + height. + + +Forward projection +------------------ + +.. gobj:class:: forwardproject + + Computes the forward projection of slices into sinograms. + + .. gobj:prop:: number:int + + Number of final 1D projections, that means height of the sinogram. + + .. gobj:prop:: angle-step:float + + Angular step between two adjacent projections. If not changed, it is + simply pi divided by :gobj:prop:`num-projections`. + + +Laminographic backprojection +---------------------------- + +.. gobj:class:: lamino-backproject + + Backprojects parallel beam computed laminography projection-by-projection + into a 3D volume. + + .. gobj:prop:: x-region:GValueArray + + X region for reconstruction as (from, to, step). + + .. gobj:prop:: y-region:GValueArray + + Y region for reconstruction as (from, to, step). + + .. gobj:prop:: z:float + + Z coordinate of the reconstructed slice. + + .. gobj:prop:: region:GValueArray + + Region for the parameter along z-axis as (from, to, step). + + .. gobj:prop:: projection-offset:GValueArray + + Offset to projection data as (x, y) for the case input data is cropped + to the necessary range of interest. + + .. gobj:prop:: center:GValueArray + + Center of the volume with respect to projections (x, y), (rotation + axes). + + .. gobj:prop:: overall-angle:float + + Angle covered by all projections (can be negative for negative steps in + case only num-projections is specified) + + .. gobj:prop:: num-projections:uint + + Number of projections. + + .. gobj:prop:: tomo-angle:float + + Tomographic rotation angle in radians (used for acquiring projections). + + .. gobj:prop:: lamino-angle:float + + Absolute laminogrpahic angle in radians determining the sample tilt. + + .. gobj:prop:: roll-angle:float + + Sample angular misalignment to the side (roll) in radians (CW is + positive). + + .. gobj:prop:: parameter:string + + Which paramter will be varied along the z-axis, from "z", "x-center", + "lamino-angle", "roll-angle". + + +Phase retrieval +--------------- + +.. gobj:class:: retrieve-phase + + Computes and applies a fourier filter to correct phase-shifted data. + Expects frequencies as an input and produces frequencies as an output. + + .. gobj:prop:: method:string + + Retrieval method which is one of ``tie``, ``ctf``, ``ctfhalfsin``, + ``qp``, ``qphalfsine`` or ``qp2``. + + .. gobj:prop:: energy:float + + Energy in keV. + + .. gobj:prop:: distance:float + + Distance in meter. + + .. gobj:prop:: pixel-size:float + + Pixel size in meter. + + .. gobj:prop:: regularization-rate:float + + Regularization parameter is log10 of the constant to be added to the + denominator to regularize the singularity at zero frequency: 1/sin(x) -> + 1/(sin(x)+10^-RegPar). + + Typical values [2, 3]. + + .. gobj:prop:: thresholding-rate:float + + Parameter for Quasiparticle phase retrieval which defines the width of + the rings to be cropped around the zero crossing of the CTF denominator + in Fourier space. + + Typical values in [0.01, 0.1], ``qp`` retrieval is rather independent of + cropping width. + + General matrix-matrix multiplication ------------------------------------- +==================================== .. gobj:class:: gemm @@ -582,8 +679,26 @@ General matrix-matrix multiplication Scalar multiplied with C. -Auxiliary filters -================= +Segmentation +============ + +.. gobj:class:: segment + + Segments a stack of images given a field of labels using the random walk + algorithm described in [#]_. The first + input stream must contain three-dimensional image stacks, the second input + stream a label image with the same width and height as the images. Any pixel + value other than zero is treated as a label and used to determine segments + in all directions. + + .. [#] + Lösel and Heuveline, *Enhancing a Diffusion Algorithm for 4D Image + Segmentation Using Local Information* in Proc. SPIE 9784, Medical + Imaging 2016, https://doi.org/10.1117/12.2216202 + + +Auxiliary +========= Buffering --------- @@ -607,8 +722,8 @@ Loops avoid unnecessary copies. You can expect the data items to be on the device where the data originated. - .. gobj:prop:: count:int - + .. gobj:prop:: number:int + Number of iterations for each received data item. @@ -624,25 +739,3 @@ Monitoring If set print the given numbers of items on stdout as hexadecimally formatted numbers. - - -Slicing -------- - -.. gobj:class:: slice - - Slices a three-dimensional input buffer to two-dimensional slices. - - -Stacking --------- - -.. gobj:class:: stack - - Symmetrical to the slice filter, the stack filter stacks two-dimensional - input. - - .. gobj:prop:: number:int - - Number of items, i.e. the length of the third dimension. - diff --git a/docs/generators.rst b/docs/generators.rst index 73b29de..1e0bc53 100644 --- a/docs/generators.rst +++ b/docs/generators.rst @@ -5,11 +5,8 @@ Generators Generators produce data and have at least one output but no input. -Data readers -============ - File reader ------------ +=========== .. gobj:class:: read @@ -73,7 +70,7 @@ File reader Memory reader -------------- +============= .. gobj:class:: memory-in @@ -124,7 +121,7 @@ Memory reader UcaCamera reader ----------------- +================ .. gobj:class:: camera @@ -159,7 +156,7 @@ UcaCamera reader stdin reader ------------- +============ .. gobj:class:: stdin @@ -180,12 +177,8 @@ stdin reader Specifies the bit depth of input. - -Auxiliary generators -==================== - -Metaballs ---------- +Metaball simulation +=================== .. gobj:class:: metaballs @@ -208,14 +201,9 @@ Metaballs Length of data stream. - .. gobj:prop:: frames-per-second:int - - Simulate behaviour by restricting the number of output images per - second. - -Empty data ----------- +Data generation +=============== .. gobj:class:: dummy-data diff --git a/docs/sinks.rst b/docs/sinks.rst index 2c6975d..3e088a4 100644 --- a/docs/sinks.rst +++ b/docs/sinks.rst @@ -5,11 +5,8 @@ Sinks Sinks are endpoints and have at least one input but no output. -Data writers -============ - File writer ------------ +=========== .. gobj:class:: write @@ -46,7 +43,7 @@ File writer Memory writer -------------- +============= .. gobj:class:: memory-out @@ -87,7 +84,7 @@ Memory writer stdout writer -------------- +============= .. gobj:class:: stdout @@ -103,7 +100,7 @@ Auxiliary sink ============== Null ----- +==== .. gobj:class:: null diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 3370b90..3ae7230 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -5,6 +5,7 @@ set(ufofilter_SRCS ufo-average-task.c ufo-backproject-task.c ufo-bin-task.c + ufo-binarize-task.c ufo-blur-task.c ufo-buffer-task.c ufo-calculate-task.c @@ -58,6 +59,7 @@ set(ufofilter_SRCS ufo-ring-pattern-task.c ufo-ringwriter-task.c ufo-replicate-task.c + ufo-segment-task.c ufo-sleep-task.c ufo-slice-task.c ufo-stack-task.c @@ -112,7 +114,7 @@ set(ufofilter_LIBS set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -std=c99 -pedantic -Wall -Wextra -fPIC -Wno-unused-parameter -Wno-deprecated-declarations") -add_definitions ("-DBURST=${BP_BURST}") +add_definitions (-DBURST=${BP_BURST} -D_FORTIFY_SOURCE=1 -D_FILE_OFFSET_BITS=64 -D_LARGE_FILES) #}}} #{{{ Dependency checks find_package(TIFF) @@ -190,6 +192,18 @@ if (GSL_FOUND) list(APPEND ufofilter_SRCS ufo-measure-task.c) endif () +if (TARGET oclfft) + option(WITH_OCLFFT "Use Apple FFT" ON) + + if (WITH_OCLFFT) + list(APPEND fft_aux_LIBS oclfft) + list(APPEND ifft_aux_LIBS oclfft) + list(APPEND retrieve_phase_aux_LIBS oclfft) + list(APPEND filter_aux_LIBS oclfft) + set(HAVE_AMD OFF) + endif () +endif () + if (CLFFT_FOUND) option(WITH_CLFFT "Use AMD clFFT" ON) @@ -200,12 +214,6 @@ if (CLFFT_FOUND) list(APPEND retrieve_phase_aux_LIBS ${CLFFT_LIBRARIES}) list(APPEND filter_aux_LIBS ${CLFFT_LIBRARIES}) set(HAVE_AMD ON) - else () - list(APPEND fft_aux_LIBS oclfft) - list(APPEND ifft_aux_LIBS oclfft) - list(APPEND retrieve_phase_aux_LIBS oclfft) - list(APPEND filter_aux_LIBS oclfft) - set(HAVE_AMD OFF) endif () endif () diff --git a/src/kernels/binarize.cl b/src/kernels/binarize.cl new file mode 100644 index 0000000..9169d2f --- /dev/null +++ b/src/kernels/binarize.cl @@ -0,0 +1,27 @@ +/* + * Copyright (C) 2016 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 +binarize (global float *input, + global float *output, + const float threshold) +{ + size_t idx = get_global_id (0); + output[idx] = input[idx] < threshold ? 0.0f : 1.0f; +} diff --git a/src/kernels/metaballs.cl b/src/kernels/metaballs.cl index 49740a8..876e3a3 100644 --- a/src/kernels/metaballs.cl +++ b/src/kernels/metaballs.cl @@ -17,10 +17,17 @@ * License along with this library. If not, see <http://www.gnu.org/licenses/>. */ +typedef struct { + float x; + float y; + float vx; + float vy; + float size; +} Ball; + kernel void draw_metaballs (global float *output, - global float2 *positions, - constant float *sizes, + global Ball *balls, uint num_balls) { const int idx = get_global_id(0); @@ -28,12 +35,14 @@ draw_metaballs (global float *output, const float r = 1.0f; float sum = 0.0f; + for (int i = 0; i < num_balls; i++) { - float x = (positions[i].x - idx); - float y = (positions[i].y - idy); - sum += sizes[i] / sqrt(x*x + y*y); + float x = (balls[i].x - idx); + float y = (balls[i].y - idy); + sum += balls[i].size / sqrt(x*x + y*y); } - if ((sum > (r - 0.01f)) && (sum < (r + 0.01f))) + + if (sum > 1.0f) output[idy * get_global_size(0) + idx] = 1.0f; else output[idy * get_global_size(0) + idx] = 0.0f; diff --git a/src/kernels/segment.cl b/src/kernels/segment.cl new file mode 100644 index 0000000..9ff56ee --- /dev/null +++ b/src/kernels/segment.cl @@ -0,0 +1,216 @@ +/* + * Copyright (C) 2016 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/>. + */ + +typedef struct { + int x; + int y; +} Label; + +#define C_WEST 0 +#define C_NORTH 1 +#define C_EAST 2 +#define C_SOUTH 3 +#define C_UP 4 +#define C_DOWN 5 + +kernel void +walk (global float *slices, + global ushort *accumulator, + global Label *prelabeled, + const int width, + const int height, + const int num_slices, + global float *random) +{ + size_t idx; + int x, y, z; + int offset; + int i, j; + float current; + float c[6]; /* west, north, east, south, up, down */ + float d[6]; + float sample, sum, mean; + float sigma_sq; + + idx = get_global_id (0); + x = prelabeled[idx].x; + y = prelabeled[idx].y; + z = 0; + offset = width * height; + current = slices[y * width + x + z * offset]; + +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 2 +#endif + for (int depth = 0; depth < num_slices * 1000; depth++) { + /* FIXME: race condition */ + accumulator[y * width + x + z * offset] += 1; + + c[C_WEST] = slices[y * width + x - 1 + z * offset]; + c[C_EAST] = slices[y * width + x + 1 + z * offset]; + c[C_NORTH] = slices[(y - 1) * width + x + z * offset]; + c[C_SOUTH] = slices[(y + 1) * width + x + z * offset]; + + if (z > 0) + c[C_DOWN] = slices[y * width + x + (z - 1) * offset]; + else + c[C_DOWN] = 0.0f; + + if (z < num_slices - 1) + c[C_UP] = slices[y * width + x + (z + 1) * offset]; + else + c[C_UP] = 0.0f; + + /* compute mean of neighbourhood */ + sum = 0.0f; + + for (i = 0; i < 6; i++) + sum += c[i]; + + mean = sum / 6; + + /* compute std deviation */ + sum = 0.0f; + + for (i = 0; i < 6; i++) + sum += (c[i] - mean) * (c[i] - mean); + + sigma_sq = sqrt (sum / 5.0f); + + /* compute probabilities */ + for (i = 0; i < 6; i++) { + const float v = (c[i] - current); + c[i] = exp (- v * v / (2 * sigma_sq)) / (sqrt (2 * M_PI_F * sigma_sq)); + } + + /* normalize and initialize d array */ + sum = 0.0f; + + for (i = 0; i < 6; i++) + sum += c[i]; + + for (i = 0; i < 6; i++) { + c[i] /= sum; + d[i] = c[i]; + } + + /* insertion sort to get a cumulative distribution */ + for (i = 1; i < 6; i++) { + float v = d[i]; + j = i; + + while (j > 0 && d[j - 1] < v) { + d[j] = d[j - 1]; + j--; + } + + d[j] = v; + } + + /* draw sample according to our distribution using the inversion method */ + sample = random[(idx + depth) % 32768]; + sum = 0.0f; + i = 0; + + for (; i < 6; i++) { + sum += d[i]; + + if (sum >= sample) + break; + } + + /* find original position */ + for (j = 0; j < 6 && c[j] != d[i]; j++) + ; + + switch (j) { + case C_WEST: + x = max (x - 1, 0); + break; + case C_NORTH: + y = max (y - 1, 0); + break; + case C_EAST: + x = min (x + 1, width - 1); + break; + case C_SOUTH: + y = min (y + 1, height - 1); + break; + case C_UP: + z = min (z + 1, num_slices - 1); + break; + case C_DOWN: + z = max (z - 1, 0); + break; + } + } +} + +kernel void +threshold (global ushort *accumulator, + global uint *bitmap, + const int segment) +{ + const int x = get_global_id (0); + const int y = get_global_id (1); + const int z = get_global_id (2); + const int width = get_global_size (0); + const int height = get_global_size (1); + const int num_slices = get_global_size (2); + const int offset = y * width * 32 + z * width * 32 * height; + uint v = 0; + + for (int b = 0; b < 32; b++) { + ushort d = accumulator[offset + 32 * x + b]; + + if (d > 20) + v = v | (1 << b); + } + + bitmap[segment * width * height * num_slices + z * width * height + y * width + x] = (uint) v; +} + +kernel void +render (global uint *bitmap, + global float *output, + const global uint *label_map, + const int slice, + const int num_segments, + const int num_slices) +{ + int x = get_global_id (0); + int y = get_global_id (1); + int width = get_global_size (0); + int height = get_global_size (1); + int r[32] = {0}; + + for (int segment = 0; segment < num_segments; segment++) { + uint v; + + v = bitmap[segment * width * height * num_slices + slice * width * height + y * width + x]; + + for (int b = 0; b < 32; b++) { + if (v & (1 << b)) + r[b] = label_map[segment]; + } + } + + for (int i = 0; i < 32; i++) + output[y * width * 32 + 32 * x + i] = (float) r[i]; +} diff --git a/src/readers/ufo-edf-reader.c b/src/readers/ufo-edf-reader.c index 9cee5f7..1cf4f1f 100644 --- a/src/readers/ufo-edf-reader.c +++ b/src/readers/ufo-edf-reader.c @@ -120,6 +120,9 @@ ufo_edf_reader_read (UfoReader *reader, /* Read the full ROI at once if no stepping is specified */ num_bytes = width * roi_height; num_read = fread (data, 1, num_bytes, priv->fp); + + if (num_read != num_bytes) + return; } else { for (guint i = 0; i < num_rows - 1; i++) { diff --git a/src/readers/ufo-raw-reader.c b/src/readers/ufo-raw-reader.c index 665e19b..f7197f3 100644 --- a/src/readers/ufo-raw-reader.c +++ b/src/readers/ufo-raw-reader.c @@ -26,12 +26,12 @@ struct _UfoRawReaderPrivate { FILE *fp; - gssize total_size; - gssize frame_size; + gsize total_size; + gsize frame_size; gsize bytes_per_pixel; guint width; guint height; - glong offset; + gulong offset; UfoBufferDepth bitdepth; }; @@ -72,7 +72,7 @@ ufo_raw_reader_can_open (UfoReader *reader, if (!g_str_has_suffix (filename, ".raw")) return FALSE; - if (priv->width == 0 || priv->height == 0 || priv->bitdepth == 0) { + if (priv->width == 0 || priv->height == 0 || priv->bitdepth == UFO_BUFFER_DEPTH_INVALID) { g_warning ("`raw-width', `raw-height' or `raw-bitdepth' was not set"); return FALSE; } @@ -111,9 +111,11 @@ static gboolean ufo_raw_reader_data_available (UfoReader *reader) { UfoRawReaderPrivate *priv; + glong pos; priv = UFO_RAW_READER_GET_PRIVATE (reader); - return priv->fp != NULL && (ftell (priv->fp) + priv->offset + priv->frame_size) <= priv->total_size; + pos = ftell (priv->fp); + return priv->fp != NULL && pos >= 0 && (((gulong) pos) + priv->offset + priv->frame_size) <= priv->total_size; } static void @@ -185,7 +187,7 @@ ufo_raw_reader_set_property (GObject *object, } break; case PROP_OFFSET: - priv->offset = g_value_get_long (value); + priv->offset = g_value_get_ulong (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -212,7 +214,7 @@ ufo_raw_reader_get_property (GObject *object, g_value_set_uint (value, priv->bitdepth); break; case PROP_OFFSET: - g_value_set_long (value, priv->offset); + g_value_set_ulong (value, priv->offset); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -298,6 +300,6 @@ ufo_raw_reader_init (UfoRawReader *self) priv->fp = NULL; priv->width = 0; priv->height = 0; - priv->bitdepth = 0; + priv->bitdepth = UFO_BUFFER_DEPTH_INVALID; priv->offset = 0L; } diff --git a/src/ufo-backproject-task.c b/src/ufo-backproject-task.c index e4645aa..923b0f3 100644 --- a/src/ufo-backproject-task.c +++ b/src/ufo-backproject-task.c @@ -32,6 +32,12 @@ typedef enum { MODE_TEXTURE } Mode; +static GEnumValue mode_values[] = { + { MODE_NEAREST, "MODE_NEAREST", "nearest" }, + { MODE_TEXTURE, "MODE_TEXTURE", "texture" }, + { 0, NULL, NULL} +}; + struct _UfoBackprojectTaskPrivate { cl_context context; cl_kernel nearest_kernel; @@ -350,10 +356,7 @@ ufo_backproject_task_set_property (GObject *object, priv->luts_changed = TRUE; break; case PROP_MODE: - if (!g_strcmp0 (g_value_get_string (value), "nearest")) - priv->mode = MODE_NEAREST; - else if (!g_strcmp0 (g_value_get_string (value), "texture")) - priv->mode = MODE_TEXTURE; + priv->mode = g_value_get_enum (value); break; case PROP_ROI_X: priv->roi_x = g_value_get_uint (value); @@ -398,14 +401,7 @@ ufo_backproject_task_get_property (GObject *object, g_value_set_double (value, priv->angle_offset); break; case PROP_MODE: - switch (priv->mode) { - case MODE_NEAREST: - g_value_set_string (value, "nearest"); - break; - case MODE_TEXTURE: - g_value_set_string (value, "texture"); - break; - } + g_value_set_enum (value, priv->mode); break; case PROP_ROI_X: g_value_set_uint (value, priv->roi_x); @@ -475,11 +471,11 @@ ufo_backproject_task_class_init (UfoBackprojectTaskClass *klass) G_PARAM_READWRITE); properties[PROP_MODE] = - g_param_spec_string ("mode", - "Backprojection mode", - "Backprojection mode from: \"nearest\", \"texture\"", - "texture", - G_PARAM_READWRITE); + g_param_spec_enum ("mode", + "Backprojection mode (\"nearest\", \"texture\")", + "Backprojection mode (\"nearest\", \"texture\")", + g_enum_register_static ("mode", mode_values), + MODE_TEXTURE, G_PARAM_READWRITE); properties[PROP_ROI_X] = g_param_spec_uint ("roi-x", diff --git a/src/ufo-bin-task.h b/src/ufo-bin-task.h index 8cad6fc..a7ff60a 100644 --- a/src/ufo-bin-task.h +++ b/src/ufo-bin-task.h @@ -50,4 +50,5 @@ GType ufo_bin_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-binarize-task.c b/src/ufo-binarize-task.c new file mode 100644 index 0000000..082ff28 --- /dev/null +++ b/src/ufo-binarize-task.c @@ -0,0 +1,220 @@ +/* + * Copyright (C) 2016 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/>. + */ + +#ifdef __APPLE__ +#include <OpenCL/cl.h> +#else +#include <CL/cl.h> +#endif + +#include "ufo-binarize-task.h" + + +struct _UfoBinarizeTaskPrivate { + cl_kernel kernel; + gfloat threshold; +}; + +static void ufo_task_interface_init (UfoTaskIface *iface); + +G_DEFINE_TYPE_WITH_CODE (UfoBinarizeTask, ufo_binarize_task, UFO_TYPE_TASK_NODE, + G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK, + ufo_task_interface_init)) + +#define UFO_BINARIZE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_BINARIZE_TASK, UfoBinarizeTaskPrivate)) + +enum { + PROP_0, + PROP_THRESHOLD, + N_PROPERTIES +}; + +static GParamSpec *properties[N_PROPERTIES] = { NULL, }; + +UfoNode * +ufo_binarize_task_new (void) +{ + return UFO_NODE (g_object_new (UFO_TYPE_BINARIZE_TASK, NULL)); +} + +static void +ufo_binarize_task_setup (UfoTask *task, + UfoResources *resources, + GError **error) +{ + UfoBinarizeTaskPrivate *priv; + + priv = UFO_BINARIZE_TASK_GET_PRIVATE (task); + priv->kernel = ufo_resources_get_kernel (resources, "binarize.cl", "binarize", error); + + if (priv->kernel != NULL) + UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel)); +} + +static void +ufo_binarize_task_get_requisition (UfoTask *task, + UfoBuffer **inputs, + UfoRequisition *requisition) +{ + ufo_buffer_get_requisition (inputs[0], requisition); +} + +static guint +ufo_binarize_task_get_num_inputs (UfoTask *task) +{ + return 1; +} + +static guint +ufo_binarize_task_get_num_dimensions (UfoTask *task, + guint input) +{ + return 2; +} + +static UfoTaskMode +ufo_binarize_task_get_mode (UfoTask *task) +{ + return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU; +} + +static gboolean +ufo_binarize_task_process (UfoTask *task, + UfoBuffer **inputs, + UfoBuffer *output, + UfoRequisition *requisition) +{ + UfoBinarizeTaskPrivate *priv; + UfoGpuNode *node; + UfoProfiler *profiler; + cl_command_queue cmd_queue; + cl_mem in_mem; + cl_mem out_mem; + gsize size = 1; + + priv = UFO_BINARIZE_TASK_GET_PRIVATE (task); + node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task))); + cmd_queue = ufo_gpu_node_get_cmd_queue (node); + in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue); + out_mem = ufo_buffer_get_device_array (output, cmd_queue); + + for (guint i = 0; i < requisition->n_dims; i++) + size *= requisition->dims[i]; + + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_mem)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &out_mem)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (gfloat), &priv->threshold)); + + profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); + ufo_profiler_call (profiler, cmd_queue, priv->kernel, 1, &size, NULL); + + return TRUE; +} + +static void +ufo_binarize_task_set_property (GObject *object, + guint property_id, + const GValue *value, + GParamSpec *pspec) +{ + UfoBinarizeTaskPrivate *priv = UFO_BINARIZE_TASK_GET_PRIVATE (object); + + switch (property_id) { + case PROP_THRESHOLD: + priv->threshold = g_value_get_float (value); + break; + default: + G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); + break; + } +} + +static void +ufo_binarize_task_get_property (GObject *object, + guint property_id, + GValue *value, + GParamSpec *pspec) +{ + UfoBinarizeTaskPrivate *priv = UFO_BINARIZE_TASK_GET_PRIVATE (object); + + switch (property_id) { + case PROP_THRESHOLD: + g_value_set_float (value, priv->threshold); + break; + default: + G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); + break; + } +} + +static void +ufo_binarize_task_finalize (GObject *object) +{ + UfoBinarizeTaskPrivate *priv; + + priv = UFO_BINARIZE_TASK_GET_PRIVATE (object); + + if (priv->kernel != NULL) { + UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel)); + priv->kernel = NULL; + } + + G_OBJECT_CLASS (ufo_binarize_task_parent_class)->finalize (object); +} + +static void +ufo_task_interface_init (UfoTaskIface *iface) +{ + iface->setup = ufo_binarize_task_setup; + iface->get_num_inputs = ufo_binarize_task_get_num_inputs; + iface->get_num_dimensions = ufo_binarize_task_get_num_dimensions; + iface->get_mode = ufo_binarize_task_get_mode; + iface->get_requisition = ufo_binarize_task_get_requisition; + iface->process = ufo_binarize_task_process; +} + +static void +ufo_binarize_task_class_init (UfoBinarizeTaskClass *klass) +{ + GObjectClass *oclass = G_OBJECT_CLASS (klass); + + oclass->set_property = ufo_binarize_task_set_property; + oclass->get_property = ufo_binarize_task_get_property; + oclass->finalize = ufo_binarize_task_finalize; + + properties[PROP_THRESHOLD] = + g_param_spec_float ("threshold", + "Absolute threshold", + "Absolute threshold", + -G_MAXFLOAT, G_MAXFLOAT, 1.0f, + G_PARAM_READWRITE); + + for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) + g_object_class_install_property (oclass, i, properties[i]); + + g_type_class_add_private (oclass, sizeof(UfoBinarizeTaskPrivate)); +} + +static void +ufo_binarize_task_init(UfoBinarizeTask *self) +{ + self->priv = UFO_BINARIZE_TASK_GET_PRIVATE(self); + self->priv->threshold = 1.0f; + self->priv->kernel = NULL; +} diff --git a/src/ufo-binarize-task.h b/src/ufo-binarize-task.h new file mode 100644 index 0000000..bed3dc5 --- /dev/null +++ b/src/ufo-binarize-task.h @@ -0,0 +1,54 @@ +/* + * 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_BINARIZE_TASK_H +#define __UFO_BINARIZE_TASK_H + +#include <ufo/ufo.h> + +G_BEGIN_DECLS + +#define UFO_TYPE_BINARIZE_TASK (ufo_binarize_task_get_type()) +#define UFO_BINARIZE_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_BINARIZE_TASK, UfoBinarizeTask)) +#define UFO_IS_BINARIZE_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_BINARIZE_TASK)) +#define UFO_BINARIZE_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_BINARIZE_TASK, UfoBinarizeTaskClass)) +#define UFO_IS_BINARIZE_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_BINARIZE_TASK)) +#define UFO_BINARIZE_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_BINARIZE_TASK, UfoBinarizeTaskClass)) + +typedef struct _UfoBinarizeTask UfoBinarizeTask; +typedef struct _UfoBinarizeTaskClass UfoBinarizeTaskClass; +typedef struct _UfoBinarizeTaskPrivate UfoBinarizeTaskPrivate; + +struct _UfoBinarizeTask { + UfoTaskNode parent_instance; + + UfoBinarizeTaskPrivate *priv; +}; + +struct _UfoBinarizeTaskClass { + UfoTaskNodeClass parent_class; +}; + +UfoNode *ufo_binarize_task_new (void); +GType ufo_binarize_task_get_type (void); + +G_END_DECLS + +#endif + diff --git a/src/ufo-calculate-task.c b/src/ufo-calculate-task.c index 8eb386a..1b1d158 100644 --- a/src/ufo-calculate-task.c +++ b/src/ufo-calculate-task.c @@ -64,9 +64,12 @@ make_kernel (UfoCalculateTaskPrivate *priv, UfoResources *resources, GError **er expression = priv->expression == NULL ? default_expression : priv->expression; source = (gchar *) g_try_malloc (strlen (template) + strlen (expression)); - if (!source) { - g_warning ("Error allocating kernel string memory"); + if (source == NULL) { + g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP, + "Error allocating kernel string memory"); + return; } + if (priv->kernel) { UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel)); } diff --git a/src/ufo-calculate-task.h b/src/ufo-calculate-task.h index 6fc08e8..e8399b1 100644 --- a/src/ufo-calculate-task.h +++ b/src/ufo-calculate-task.h @@ -50,4 +50,5 @@ GType ufo_calculate_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-clip-task.h b/src/ufo-clip-task.h index aa04d68..b943a96 100644 --- a/src/ufo-clip-task.h +++ b/src/ufo-clip-task.h @@ -50,4 +50,5 @@ GType ufo_clip_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-concatenate-result-task.h b/src/ufo-concatenate-result-task.h index 6774743..e50b239 100644 --- a/src/ufo-concatenate-result-task.h +++ b/src/ufo-concatenate-result-task.h @@ -63,4 +63,5 @@ GType ufo_concatenate_result_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-detect-edge-task.c b/src/ufo-detect-edge-task.c index f2d85ad..a206019 100644 --- a/src/ufo-detect-edge-task.c +++ b/src/ufo-detect-edge-task.c @@ -26,20 +26,26 @@ #include "ufo-detect-edge-task.h" -struct _UfoDetectEdgeTaskPrivate { - enum { - FILTER_SOBEL = 0, - FILTER_LAPLACE, - FILTER_PREWITT, - } type; +typedef enum { + FILTER_SOBEL = 0, + FILTER_LAPLACE, + FILTER_PREWITT, +} Filter; + +static GEnumValue filter_values[] = { + { FILTER_SOBEL, "FILTER_SOBEL", "sobel" }, + { FILTER_LAPLACE, "FILTER_LAPLACE", "laplace" }, + { FILTER_PREWITT, "FILTER_PREWITT", "prewitt" }, + { 0, NULL, NULL} +}; +struct _UfoDetectEdgeTaskPrivate { + Filter type; cl_context context; cl_kernel kernel; cl_mem mask_mem; }; -static const gchar* FILTER_NAMES[] = { "sobel", "laplace", "prewitt", NULL }; - static gfloat FILTER_MASKS[][9] = { { 1.0f, 0.0f, -1.0f, /* Sobel in one direction */ 2.0f, 0.0f, -2.0f, @@ -52,7 +58,6 @@ static gfloat FILTER_MASKS[][9] = { -1.0f, 0.0f, 1.0f } }; - static void ufo_task_interface_init (UfoTaskIface *iface); G_DEFINE_TYPE_WITH_CODE (UfoDetectEdgeTask, ufo_detect_edge_task, UFO_TYPE_TASK_NODE, @@ -166,17 +171,8 @@ ufo_detect_edge_task_set_property (GObject *object, switch (property_id) { case PROP_TYPE: - { - const gchar *type; - - type = g_value_get_string (value); - - for (guint i = 0; FILTER_NAMES[i] != NULL; i++) - if (!g_strcmp0 (type, FILTER_NAMES[i])) - priv->type = i; - } + priv->type = g_value_get_enum (value); break; - default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); break; @@ -195,9 +191,8 @@ ufo_detect_edge_task_get_property (GObject *object, switch (property_id) { case PROP_TYPE: - g_value_set_string (value, FILTER_NAMES[priv->type]); + g_value_set_enum (value, priv->type); break; - default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); break; @@ -245,11 +240,12 @@ ufo_detect_edge_task_class_init (UfoDetectEdgeTaskClass *klass) oclass->finalize = ufo_detect_edge_task_finalize; properties[PROP_TYPE] = - g_param_spec_string ("type", - "Filter type (\"sobel\", \"laplace\", \"prewitt\")", - "Filter type (\"sobel\", \"laplace\", \"prewitt\")", - "sobel", - G_PARAM_READWRITE); + g_param_spec_enum ("type", + "Filter type (\"sobel\", \"laplace\", \"prewitt\")", + "Filter type (\"sobel\", \"laplace\", \"prewitt\")", + g_enum_register_static ("type", filter_values), + FILTER_SOBEL, + G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (oclass, i, properties[i]); diff --git a/src/ufo-detect-edge-task.h b/src/ufo-detect-edge-task.h index 029bc4e..7a9d623 100644 --- a/src/ufo-detect-edge-task.h +++ b/src/ufo-detect-edge-task.h @@ -50,4 +50,5 @@ GType ufo_detect_edge_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-dfi-sinc-task.c b/src/ufo-dfi-sinc-task.c index 7301a6d..9f2c8ee 100644 --- a/src/ufo-dfi-sinc-task.c +++ b/src/ufo-dfi-sinc-task.c @@ -24,7 +24,6 @@ #endif #include <stdio.h> -#include <stdlib.h> #include <string.h> #include <math.h> @@ -123,15 +122,10 @@ ufo_dfi_sinc_task_sinc(gfloat x) * Returns: an array. */ static gfloat * -ufo_dfi_sinc_task_get_ktbl(size_t length) +ufo_dfi_sinc_task_get_ktbl (size_t length) { gfloat *ktbl = (gfloat *) g_malloc0 (length * sizeof (gfloat)); - if (!length % 2) { - g_print("Error: Length %zu of ktbl cannot be even!\n", length); - exit(1); - } - size_t ktbl_len2 = (length - 1) / 2; gfloat step = (gfloat) G_PI / (gfloat) ktbl_len2; gfloat value = -(gfloat) ktbl_len2 * step; @@ -259,8 +253,8 @@ ufo_dfi_sinc_task_process (UfoTask *task, if (priv->roi_size >= 1 && priv->roi_size <= raster_size) { interp_grid_rows = (int) ceil ((float) priv->roi_size/2.0 / (float)BLOCK_SIZE); - } - + } + gint spectrum_offset = (raster_size - (interp_grid_cols * BLOCK_SIZE)) / 2; gfloat max_radius = (gfloat) (interp_grid_cols * BLOCK_SIZE) / 2.0f; @@ -321,7 +315,7 @@ ufo_dfi_sinc_task_process (UfoTask *task, profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); ufo_profiler_call (profiler, cmd_queue, priv->dfi_sinc_kernel, requisition->n_dims, working_size, local_work_size); - + return TRUE; } @@ -338,7 +332,16 @@ ufo_dfi_sinc_task_set_property (GObject *object, priv->L = g_value_get_uint (value); break; case PROP_NUM_PRESAMPLED_VLS: - priv->number_presampled_values = g_value_get_uint (value); + { + guint v; + + v = g_value_get_uint (value); + + if (v % 2) + g_warning ("::number-presampled-values cannot be even"); + else + priv->number_presampled_values = v; + } break; case PROP_ROI_SIZE: priv->roi_size = g_value_get_int (value); @@ -383,7 +386,7 @@ static void ufo_dfi_sinc_task_finalize (GObject *object) { UfoDfiSincTaskPrivate *priv = UFO_DFI_SINC_TASK_GET_PRIVATE (object); - + if (priv->in_tex) UFO_RESOURCES_CHECK_CLERR(clReleaseMemObject (priv->in_tex)); G_OBJECT_CLASS (ufo_dfi_sinc_task_parent_class)->finalize (object); diff --git a/src/ufo-dump-ring-task.h b/src/ufo-dump-ring-task.h index e1c203d..7eaaeba 100644 --- a/src/ufo-dump-ring-task.h +++ b/src/ufo-dump-ring-task.h @@ -63,4 +63,5 @@ GType ufo_dump_ring_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-duplicate-task.h b/src/ufo-duplicate-task.h index f40b5fa..bfb530f 100644 --- a/src/ufo-duplicate-task.h +++ b/src/ufo-duplicate-task.h @@ -63,4 +63,5 @@ GType ufo_duplicate_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-fftmult-task.c b/src/ufo-fftmult-task.c index 5f3ad2f..ff088ce 100644 --- a/src/ufo-fftmult-task.c +++ b/src/ufo-fftmult-task.c @@ -174,7 +174,7 @@ ufo_fftmult_task_process (UfoTask *task, UfoRequisition *requisition) { UfoFftmultTaskPrivate *priv; - priv = UFO_FFTMULT_TASK_GET_PRIVATE (task); + /* Forwarding ring radius metada to next plugin */ unsigned radius, number_ones; get_ring_metadata(inputs[1], &number_ones, &radius); diff --git a/src/ufo-filter-particle-task.c b/src/ufo-filter-particle-task.c index e054fc5..0964b41 100644 --- a/src/ufo-filter-particle-task.c +++ b/src/ufo-filter-particle-task.c @@ -352,7 +352,7 @@ pfind (float *input, unsigned *out, unsigned nc, unsigned nr, unsigned *pcount, /* dupindex gives the new index for each index, when two index are duplicates * dupindex will affect those two indexes a same value */ for(k=0;k<ndup;k++) { - k1=dupx[k]; k2=dupy[k]; + k1=dupx[k]; for(j=k;j<ndup;j++) { if(dupy[j]== dupy[k]||dupx[j]== dupy[k]) { dupindex[ dupx[j] ]=dupindex[k1]; @@ -498,6 +498,10 @@ ufo_filter_particle_task_process (UfoTask *task, res[i].intensity = input[y * req.dims[0] + x]; } + free (cx); + free (cy); + free (rr); + return TRUE; } diff --git a/src/ufo-filter-particle-task.h b/src/ufo-filter-particle-task.h index 6220a0d..52870d8 100644 --- a/src/ufo-filter-particle-task.h +++ b/src/ufo-filter-particle-task.h @@ -63,4 +63,5 @@ GType ufo_filter_particle_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-filter-stripes1d-task.c b/src/ufo-filter-stripes1d-task.c index 88c3d03..574aac2 100644 --- a/src/ufo-filter-stripes1d-task.c +++ b/src/ufo-filter-stripes1d-task.c @@ -72,9 +72,11 @@ create_coefficients (UfoFilterStripes1dTaskPrivate *priv, const guint width) const guint real_width = width / 2 + 1; cl_int cl_err; - if (!coefficients) { + if (coefficients == NULL) { g_warning ("Could not allocate memory for coeefficients"); + return; } + if (width % 2) { g_warning ("Width must be an even number"); } diff --git a/src/ufo-filter-stripes1d-task.h b/src/ufo-filter-stripes1d-task.h index 1a9042d..c829b3d 100644 --- a/src/ufo-filter-stripes1d-task.h +++ b/src/ufo-filter-stripes1d-task.h @@ -50,4 +50,5 @@ GType ufo_filter_stripes1d_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-filter-task.c b/src/ufo-filter-task.c index c986731..2952bce 100644 --- a/src/ufo-filter-task.c +++ b/src/ufo-filter-task.c @@ -52,15 +52,15 @@ typedef enum { FILTER_BUTTERWORTH, FILTER_FARIS_BYER, FILTER_HAMMING, - N_FILTERS, } Filter; -static const gchar *filter_names[] = { - "ramp", - "ramp-fromreal", - "butterworth", - "faris-byer", - "hamming", +static GEnumValue filter_values[] = { + { FILTER_RAMP, "FILTER_RAMP", "ramp" }, + { FILTER_RAMP_FROMREAL, "FILTER_RAMP_FROMREAL", "ramp-fromreal" }, + { FILTER_BUTTERWORTH, "FILTER_BUTTERWORTH", "butterworth"}, + { FILTER_FARIS_BYER, "FILTER_FARIS_BYER", "faris-byer"}, + { FILTER_HAMMING, "FILTER_HAMMING", "hamming"}, + { 0, NULL, NULL} }; static SetupFunc filter_funcs[] = { @@ -381,12 +381,7 @@ ufo_filter_task_set_property (GObject *object, switch (property_id) { case PROP_FILTER: - for (guint i = 0; i < N_FILTERS; i++) { - if (!g_strcmp0 (g_value_get_string (value), filter_names[i])) { - priv->filter = (Filter) i; - break; - } - } + priv->filter = g_value_get_enum (value); break; case PROP_CUTOFF: priv->cutoff = g_value_get_float (value); @@ -419,7 +414,7 @@ ufo_filter_task_get_property (GObject *object, switch (property_id) { case PROP_FILTER: - g_value_set_string (value, filter_names[priv->filter]); + g_value_set_enum (value, priv->filter); break; case PROP_CUTOFF: g_value_set_float (value, priv->cutoff); @@ -456,11 +451,11 @@ ufo_filter_task_class_init (UfoFilterTaskClass *klass) oclass->get_property = ufo_filter_task_get_property; properties[PROP_FILTER] = - g_param_spec_string ("filter", + g_param_spec_enum ("filter", "Type of filter (\"ramp\", \"ramp-fromreal\", \"butterworth\", \"faris-byer\", \"hamming\")", "Type of filter (\"ramp\", \"ramp-fromreal\", \"butterworth\", \"faris-byer\", \"hamming\")", - "ramp", - G_PARAM_READWRITE); + g_enum_register_static ("filter", filter_values), + 0, G_PARAM_READWRITE); properties[PROP_CUTOFF] = g_param_spec_float ("cutoff", diff --git a/src/ufo-flatten-inplace-task.c b/src/ufo-flatten-inplace-task.c index 3117f41..a89a8f4 100644 --- a/src/ufo-flatten-inplace-task.c +++ b/src/ufo-flatten-inplace-task.c @@ -20,14 +20,17 @@ #include "ufo-flatten-inplace-task.h" typedef enum { - M_0, - M_SUM, - M_MIN, - M_MAX, - M_LAST + MODE_SUM, + MODE_MIN, + MODE_MAX, } Mode; -static const gchar *modes[] = {"sum", "min", "max"}; +static GEnumValue mode_values[] = { + { MODE_SUM, "MODE_SUM", "sum" }, + { MODE_MIN, "MODE_MIN", "min" }, + { MODE_MAX, "MODE_MAX", "max" }, + { 0, NULL, NULL} +}; struct _UfoFlattenInplaceTaskPrivate { Mode mode; @@ -56,17 +59,6 @@ ufo_flatten_inplace_task_new (void) return UFO_NODE (g_object_new (UFO_TYPE_FLATTEN_INPLACE_TASK, NULL)); } -static Mode -string_to_mode (const gchar *s) -{ - for (Mode i = M_0 + 1; i < M_LAST; i++) { - if (!g_strcmp0 (s, modes[i - 1])) - return i; - } - - return M_0; -} - static void ufo_flatten_inplace_task_setup (UfoTask *task, UfoResources *resources, @@ -118,19 +110,19 @@ ufo_flatten_inplace_task_process (UfoTask *task, out_array = ufo_buffer_get_host_array (output, NULL); switch (priv->mode) { - case M_SUM: + case MODE_SUM: for (gsize i = 0; i < n_pixels; i++) out_array[i] += in_array[i]; break; - case M_MIN: + case MODE_MIN: for (gsize i = 0; i < n_pixels; i++) { if (in_array[i] < out_array[i]) out_array[i] = in_array[i]; } break; - case M_MAX: + case MODE_MAX: for (gsize i = 0; i < n_pixels; i++) { if (in_array[i] > out_array[i]) out_array[i] = in_array[i]; @@ -169,12 +161,7 @@ ufo_flatten_inplace_task_set_property (GObject *object, switch (property_id) { case PROP_MODE: - { - Mode mode = string_to_mode (g_value_get_string (value)); - - if (mode != M_0) - priv->mode = mode; - } + priv->mode = g_value_get_enum (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -192,7 +179,7 @@ ufo_flatten_inplace_task_get_property (GObject *object, switch (property_id) { case PROP_MODE: - g_value_set_string (value, modes[priv->mode]); + g_value_set_enum (value, priv->mode); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -228,11 +215,11 @@ ufo_flatten_inplace_task_class_init (UfoFlattenInplaceTaskClass *klass) oclass->finalize = ufo_flatten_inplace_task_finalize; properties[PROP_MODE] = - g_param_spec_string ("mode", - "Mode (min, max, sum)", - "Mode (min, max, sum)", - "", - G_PARAM_READWRITE); + g_param_spec_enum ("mode", + "Mode (min, max, sum)", + "Mode (min, max, sum)", + g_enum_register_static ("mode", mode_values), + MODE_SUM, G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (oclass, i, properties[i]); @@ -244,6 +231,6 @@ static void ufo_flatten_inplace_task_init(UfoFlattenInplaceTask *self) { self->priv = UFO_FLATTEN_INPLACE_TASK_GET_PRIVATE(self); - self->priv->mode = M_SUM; + self->priv->mode = MODE_SUM; self->priv->generated = FALSE; } diff --git a/src/ufo-flip-task.c b/src/ufo-flip-task.c index 588a0d5..a81a339 100644 --- a/src/ufo-flip-task.c +++ b/src/ufo-flip-task.c @@ -26,10 +26,16 @@ #include "ufo-flip-task.h" typedef enum { - HORIZONTAL = 0, - VERTICAL + DIRECTION_HORIZONTAL = 0, + DIRECTION_VERTICAL } Direction; +static GEnumValue direction_values[] = { + { DIRECTION_HORIZONTAL, "DIRECTION_HORIZONTAL", "horizontal" }, + { DIRECTION_VERTICAL, "DIRECTION_VERTICAL", "vertical" }, + { 0, NULL, NULL} +}; + struct _UfoFlipTaskPrivate { Direction direction; cl_kernel kernels[2]; @@ -65,8 +71,8 @@ ufo_flip_task_setup (UfoTask *task, UfoFlipTaskPrivate *priv; priv = UFO_FLIP_TASK_GET_PRIVATE (task); - priv->kernels[HORIZONTAL] = ufo_resources_get_kernel (resources, "flip.cl", "flip_horizontal", error); - priv->kernels[VERTICAL] = ufo_resources_get_kernel (resources, "flip.cl", "flip_vertical", error); + priv->kernels[DIRECTION_HORIZONTAL] = ufo_resources_get_kernel (resources, "flip.cl", "flip_horizontal", error); + priv->kernels[DIRECTION_VERTICAL] = ufo_resources_get_kernel (resources, "flip.cl", "flip_vertical", error); } static void @@ -136,12 +142,7 @@ ufo_flip_task_set_property (GObject *object, switch (property_id) { case PROP_DIRECTION: - if (g_strcmp0 (g_value_get_string (value), "horizontal") == 0) - priv->direction = HORIZONTAL; - else if (g_strcmp0 (g_value_get_string (value), "vertical") == 0) - priv->direction = VERTICAL; - else - g_warning ("`%s' cannot be set", g_value_get_string (value)); + priv->direction = g_value_get_enum (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -159,10 +160,7 @@ ufo_flip_task_get_property (GObject *object, switch (property_id) { case PROP_DIRECTION: - if (priv->direction == HORIZONTAL) - g_value_set_string (value, "horizontal"); - else if (priv->direction == VERTICAL) - g_value_set_string (value, "vertical"); + g_value_set_enum (value, priv->direction); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -197,10 +195,11 @@ ufo_flip_task_class_init (UfoFlipTaskClass *klass) oclass->finalize = ufo_flip_task_finalize; properties[PROP_DIRECTION] = - g_param_spec_string ("direction", - "Flip direction (either `horizontal' or `vertical')", - "Flip direction (either `horizontal' or `vertical')", - "horizontal", + g_param_spec_enum ("direction", + "Flip direction (`horizontal' or `vertical')", + "Flip direction (`horizontal' or `vertical')", + g_enum_register_static ("direction", direction_values), + DIRECTION_HORIZONTAL, G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) @@ -213,7 +212,7 @@ static void ufo_flip_task_init(UfoFlipTask *self) { self->priv = UFO_FLIP_TASK_GET_PRIVATE(self); - self->priv->direction = HORIZONTAL; - self->priv->kernels[HORIZONTAL] = NULL; - self->priv->kernels[VERTICAL] = NULL; + self->priv->direction = DIRECTION_HORIZONTAL; + self->priv->kernels[DIRECTION_HORIZONTAL] = NULL; + self->priv->kernels[DIRECTION_VERTICAL] = NULL; } diff --git a/src/ufo-forwardproject-task.h b/src/ufo-forwardproject-task.h index ea240d2..924f129 100644 --- a/src/ufo-forwardproject-task.h +++ b/src/ufo-forwardproject-task.h @@ -63,4 +63,5 @@ GType ufo_forwardproject_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-get-dup-circ-task.h b/src/ufo-get-dup-circ-task.h index 138fec5..396789d 100644 --- a/src/ufo-get-dup-circ-task.h +++ b/src/ufo-get-dup-circ-task.h @@ -63,4 +63,5 @@ GType ufo_get_dup_circ_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-lamino-backproject-task.c b/src/ufo-lamino-backproject-task.c index 2010296..65348ae 100644 --- a/src/ufo-lamino-backproject-task.c +++ b/src/ufo-lamino-backproject-task.c @@ -43,11 +43,19 @@ typedef enum { - PARAM_Z, - PARAM_CENTER, - PARAM_LAMINO, - PARAM_ROLL -} Param; + PARAMETER_Z, + PARAMETER_X_CENTER, + PARAMETER_LAMINO_ANGLE, + PARAMETER_ROLL_ANGLE +} Parameter; + +static GEnumValue parameter_values[] = { + { PARAMETER_Z, "PARAMETER_Z", "z" }, + { PARAMETER_X_CENTER, "X_CENTER", "x-center" }, + { PARAMETER_LAMINO_ANGLE, "LAMINO_ANGLE", "lamino-angle" }, + { PARAMETER_ROLL_ANGLE, "ROLL_ANGLE", "roll-angle" }, + { 0, NULL, NULL} +}; struct _UfoLaminoBackprojectTaskPrivate { /* private */ @@ -79,7 +87,7 @@ struct _UfoLaminoBackprojectTaskPrivate { gfloat lamino_angle; gfloat z; gfloat roll_angle; - Param parameter; + Parameter parameter; }; static void ufo_task_interface_init (UfoTaskIface *iface); @@ -182,16 +190,16 @@ ufo_lamino_backproject_task_setup (UfoTask *task, priv->context = ufo_resources_get_context (resources); switch (priv->parameter) { - case PARAM_Z: + case PARAMETER_Z: kernel_filename = g_strdup ("z_kernel.cl"); break; - case PARAM_CENTER: + case PARAMETER_X_CENTER: kernel_filename = g_strdup ("center_kernel.cl"); break; - case PARAM_LAMINO: + case PARAMETER_LAMINO_ANGLE: kernel_filename = g_strdup ("lamino_kernel.cl"); break; - case PARAM_ROLL: + case PARAMETER_ROLL_ANGLE: kernel_filename = g_strdup ("roll_kernel.cl"); break; default: @@ -337,7 +345,8 @@ ufo_lamino_backproject_task_process (UfoTask *task, x_region[1] = (gfloat) EXTRACT_INT (priv->x_region, 2); y_region[0] = (gfloat) EXTRACT_INT (priv->y_region, 0); y_region[1] = (gfloat) EXTRACT_INT (priv->y_region, 2); - if (priv->parameter == PARAM_Z) { + + if (priv->parameter == PARAMETER_Z) { z_ends[0] = z_region[0] = EXTRACT_FLOAT (priv->region, 0); z_region[1] = EXTRACT_FLOAT (priv->region, 2); z_ends[1] = EXTRACT_FLOAT (priv->region, 1); @@ -345,19 +354,22 @@ ufo_lamino_backproject_task_process (UfoTask *task, z_ends[0] = z_region[0] = priv->z; z_ends[1] = priv->z + 1.0f; } - if (priv->parameter == PARAM_CENTER) { + + if (priv->parameter == PARAMETER_X_CENTER) { x_center[0] = EXTRACT_FLOAT (priv->region, 0) - EXTRACT_INT (priv->projection_offset, 0); x_center[1] = EXTRACT_FLOAT (priv->region, 2); } else { x_center[0] = x_center[1] = EXTRACT_FLOAT (priv->center, 0) - EXTRACT_INT (priv->projection_offset, 0); } - if (priv->parameter == PARAM_LAMINO) { + + if (priv->parameter == PARAMETER_LAMINO_ANGLE) { lamino_angles[0] = EXTRACT_FLOAT (priv->region, 0); lamino_angles[1] = EXTRACT_FLOAT (priv->region, 2); } else { lamino_angles[0] = lamino_angles[1] = priv->lamino_angle; } - if (priv->parameter == PARAM_ROLL) { + + if (priv->parameter == PARAMETER_ROLL_ANGLE) { roll_angles[0] = EXTRACT_FLOAT (priv->region, 0); roll_angles[1] = EXTRACT_FLOAT (priv->region, 2); } else { @@ -374,8 +386,8 @@ ufo_lamino_backproject_task_process (UfoTask *task, /* If COPY_PROJECTION_REGION is True we copy only the part necessary */ /* for a given tomographic and laminographic angle */ - /* TODO: Extend the region determination to be able to handle PARAM_LAMINO */ - if (COPY_PROJECTION_REGION && priv->parameter != PARAM_LAMINO) { + /* TODO: Extend the region determination to be able to handle PARAMETER_LAMINO_ANGLE */ + if (COPY_PROJECTION_REGION && priv->parameter != PARAMETER_LAMINO_ANGLE) { determine_x_region (x_copy_region, priv->x_region, priv->y_region, tomo_angle, EXTRACT_FLOAT (priv->center, 0), in_req.dims[0]); determine_y_region (y_copy_region, priv->x_region, priv->y_region, z_ends, @@ -532,15 +544,7 @@ ufo_lamino_backproject_task_set_property (GObject *object, priv->roll_angle = g_value_get_float (value); break; case PROP_PARAMETER: - if (!g_strcmp0 (g_value_get_string (value), "z")) { - priv->parameter = PARAM_Z; - } else if (!g_strcmp0 (g_value_get_string (value), "x-center")) { - priv->parameter = PARAM_CENTER; - } else if (!g_strcmp0 (g_value_get_string (value), "lamino-angle")) { - priv->parameter = PARAM_LAMINO; - } else if (!g_strcmp0 (g_value_get_string (value), "roll-angle")) { - priv->parameter = PARAM_ROLL; - } + priv->parameter = g_value_get_enum (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -591,20 +595,7 @@ ufo_lamino_backproject_task_get_property (GObject *object, g_value_set_float (value, priv->roll_angle); break; case PROP_PARAMETER: - switch (priv->parameter) { - case PARAM_Z: - g_value_set_string (value, "z"); - break; - case PARAM_CENTER: - g_value_set_string (value, "x-center"); - break; - case PARAM_LAMINO: - g_value_set_string (value, "lamino-angle"); - break; - case PARAM_ROLL: - g_value_set_string (value, "roll-angle"); - break; - } + g_value_set_enum (value, priv->parameter); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -781,12 +772,13 @@ ufo_lamino_backproject_task_class_init (UfoLaminoBackprojectTaskClass *klass) G_PARAM_READWRITE); properties[PROP_PARAMETER] = - g_param_spec_string ("parameter", - "Which paramter will be varied along the z-axis", - "Which paramter will be varied along the z-axis, from \"z\", \"x-center\", \"lamino-angle\",\ - \"roll-angle\"", - "z", - G_PARAM_READWRITE); + g_param_spec_enum ("parameter", + "Which parameter will be varied along the z-axis", + "Which parameter will be varied along the z-axis (\"z\", \"x-center\", \"lamino-angle\",\ + \"roll-angle\")", + g_enum_register_static ("parameter", parameter_values), + PARAMETER_Z, + G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (oclass, i, properties[i]); @@ -828,7 +820,7 @@ ufo_lamino_backproject_task_init(UfoLaminoBackprojectTask *self) self->priv->tomo_angle = -G_MAXFLOAT; self->priv->lamino_angle = 0.0f; self->priv->roll_angle = 0.0f; - self->priv->parameter = PARAM_Z; + self->priv->parameter = PARAMETER_Z; self->priv->count = 0; self->priv->generated = FALSE; } diff --git a/src/ufo-lamino-backproject-task.h b/src/ufo-lamino-backproject-task.h index 989feaa..8308f60 100644 --- a/src/ufo-lamino-backproject-task.h +++ b/src/ufo-lamino-backproject-task.h @@ -50,4 +50,5 @@ GType ufo_lamino_backproject_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-loop-task.c b/src/ufo-loop-task.c index 4423931..f72b7c3 100644 --- a/src/ufo-loop-task.c +++ b/src/ufo-loop-task.c @@ -21,7 +21,7 @@ struct _UfoLoopTaskPrivate { - guint count; + guint number; guint current; UfoBuffer *temporary; GList *filled; @@ -37,7 +37,7 @@ G_DEFINE_TYPE_WITH_CODE (UfoLoopTask, ufo_loop_task, UFO_TYPE_TASK_NODE, enum { PROP_0, - PROP_COUNT, + PROP_NUMBER, N_PROPERTIES }; @@ -120,7 +120,7 @@ ufo_loop_task_generate (UfoTask *task, priv = UFO_LOOP_TASK_GET_PRIVATE (task); - if (priv->current == priv->count) + if (priv->current == priv->number) return FALSE; /* @@ -146,8 +146,8 @@ ufo_loop_task_set_property (GObject *object, UfoLoopTaskPrivate *priv = UFO_LOOP_TASK_GET_PRIVATE (object); switch (property_id) { - case PROP_COUNT: - priv->count = g_value_get_uint (value); + case PROP_NUMBER: + priv->number = g_value_get_uint (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -164,8 +164,8 @@ ufo_loop_task_get_property (GObject *object, UfoLoopTaskPrivate *priv = UFO_LOOP_TASK_GET_PRIVATE (object); switch (property_id) { - case PROP_COUNT: - g_value_set_uint (value, priv->count); + case PROP_NUMBER: + g_value_set_uint (value, priv->number); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -205,8 +205,8 @@ ufo_loop_task_class_init (UfoLoopTaskClass *klass) oclass->get_property = ufo_loop_task_get_property; oclass->finalize = ufo_loop_task_finalize; - properties[PROP_COUNT] = - g_param_spec_uint ("count", + properties[PROP_NUMBER] = + g_param_spec_uint ("number", "Number of loop iterations", "Number of loop iterations", 0, G_MAXUINT, 1, @@ -222,6 +222,6 @@ static void ufo_loop_task_init(UfoLoopTask *self) { self->priv = UFO_LOOP_TASK_GET_PRIVATE(self); - self->priv->count = 1; + self->priv->number = 1; self->priv->filled = NULL; } diff --git a/src/ufo-loop-task.h b/src/ufo-loop-task.h index 5b0b2e8..09918be 100644 --- a/src/ufo-loop-task.h +++ b/src/ufo-loop-task.h @@ -50,4 +50,5 @@ GType ufo_loop_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-map-slice-task.h b/src/ufo-map-slice-task.h index d51744f..38487a1 100644 --- a/src/ufo-map-slice-task.h +++ b/src/ufo-map-slice-task.h @@ -63,4 +63,5 @@ GType ufo_map_slice_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-median-filter-task.h b/src/ufo-median-filter-task.h index 1a038bf..b61fb68 100644 --- a/src/ufo-median-filter-task.h +++ b/src/ufo-median-filter-task.h @@ -63,4 +63,5 @@ GType ufo_median_filter_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-memory-in-task.h b/src/ufo-memory-in-task.h index 28fbe14..c5f5e25 100644 --- a/src/ufo-memory-in-task.h +++ b/src/ufo-memory-in-task.h @@ -50,4 +50,5 @@ GType ufo_memory_in_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-memory-out-task.h b/src/ufo-memory-out-task.h index 5821ea8..df15462 100644 --- a/src/ufo-memory-out-task.h +++ b/src/ufo-memory-out-task.h @@ -50,4 +50,5 @@ GType ufo_memory_out_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-merge-task.h b/src/ufo-merge-task.h index 4b016cd..9c45640 100644 --- a/src/ufo-merge-task.h +++ b/src/ufo-merge-task.h @@ -50,4 +50,5 @@ GType ufo_merge_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-metaballs-task.c b/src/ufo-metaballs-task.c index 47ac4e7..f4c6cbe 100644 --- a/src/ufo-metaballs-task.c +++ b/src/ufo-metaballs-task.c @@ -1,5 +1,5 @@ /* - * Copyright (C) 2011-2013 Karlsruhe Institute of Technology + * Copyright (C) 2011-2016 Karlsruhe Institute of Technology * * This file is part of Ufo. * @@ -25,28 +25,25 @@ #include "ufo-metaballs-task.h" +typedef struct { + gfloat x; + gfloat y; + gfloat vx; + gfloat vy; + gfloat size; +} Ball; struct _UfoMetaballsTaskPrivate { - cl_context context; cl_kernel kernel; - cl_mem positions_mem; - cl_mem sizes_mem; + cl_mem balls_mem; guint width; guint height; guint num_balls; guint num_iterations; guint current_iteration; - gboolean run_infinitely; - guint frames_per_second; - gsize num_position_bytes; - GTimer *timer; - gdouble seconds_per_frame; - - gfloat *positions; - gfloat *velocities; - gfloat *sizes; + Ball *balls; }; static void ufo_task_interface_init (UfoTaskIface *iface); @@ -63,8 +60,6 @@ enum { PROP_HEIGHT, PROP_NUM_BALLS, PROP_NUM_ITERATIONS, - PROP_RUN_INFINITELY, - PROP_FRAMES_PER_SECOND, N_PROPERTIES }; @@ -78,69 +73,46 @@ ufo_metaballs_task_new (void) static void ufo_metaballs_task_setup (UfoTask *task, - UfoResources *resources, - GError **error) + UfoResources *resources, + GError **error) { UfoMetaballsTaskPrivate *priv; gfloat f_width; gfloat f_height; - gsize size; + cl_context context; cl_int err = CL_SUCCESS; priv = UFO_METABALLS_TASK_GET_PRIVATE (task); - priv->context = ufo_resources_get_context (resources); - - priv->kernel = ufo_resources_get_kernel (resources, - "metaballs.cl", - "draw_metaballs", - error); + context = ufo_resources_get_context (resources); - UFO_RESOURCES_CHECK_CLERR (clRetainContext (priv->context)); + priv->kernel = ufo_resources_get_kernel (resources, "metaballs.cl", "draw_metaballs", error); if (priv->kernel != NULL) UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel)); - size = priv->num_balls * sizeof (gfloat); - priv->current_iteration = 0; - priv->seconds_per_frame = 1.0 / ((gdouble) priv->frames_per_second); - priv->num_position_bytes = 2 * size; - priv->positions = g_malloc0 (priv->num_position_bytes); - priv->velocities = g_malloc0 (priv->num_position_bytes); - priv->sizes = g_malloc0 (size); - priv->timer = g_timer_new (); + priv->balls = g_malloc0 (priv->num_balls * sizeof (Ball)); f_width = (gfloat) priv->width; f_height = (gfloat) priv->height; for (guint i = 0; i < priv->num_balls; i++) { - const guint x = 2*i, y = 2*i + 1; - priv->sizes[i] = (gfloat) g_random_double_range (f_width / 50.0f, f_width / 10.0f); - priv->positions[x] = (gfloat) g_random_double_range (0.0, (double) f_width); - priv->positions[y] = (gfloat) g_random_double_range (0.0, (double) f_height); - priv->velocities[x] = (gfloat) g_random_double_range (-4.0, 4.0); - priv->velocities[y] = (gfloat) g_random_double_range (-4.0, 4.0); + priv->balls[i].size = (gfloat) g_random_double_range (0.01 * f_width, 0.05 * f_width); + priv->balls[i].x = (gfloat) g_random_double_range (0.0, (double) f_width); + priv->balls[i].y = (gfloat) g_random_double_range (0.0, (double) f_height); + priv->balls[i].vx = (gfloat) g_random_double_range (-4.0, 4.0); + priv->balls[i].vy = (gfloat) g_random_double_range (-4.0, 4.0); }; - priv->positions_mem = clCreateBuffer(priv->context, - CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - priv->num_position_bytes, priv->positions, &err); + priv->balls_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + priv->num_balls * sizeof (Ball), priv->balls, &err); UFO_RESOURCES_CHECK_CLERR (err); - - priv->sizes_mem = clCreateBuffer(priv->context, - CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - size, priv->sizes, &err); - UFO_RESOURCES_CHECK_CLERR (err); - - UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &priv->positions_mem)); - UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_mem), &priv->sizes_mem)); - UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 3, sizeof (cl_uint), &priv->num_balls)); } static void ufo_metaballs_task_get_requisition (UfoTask *task, - UfoBuffer **inputs, - UfoRequisition *requisition) + UfoBuffer **inputs, + UfoRequisition *requisition) { UfoMetaballsTaskPrivate *priv; @@ -171,17 +143,18 @@ ufo_metaballs_task_get_mode (UfoTask *task) static gboolean ufo_metaballs_task_generate (UfoTask *task, - UfoBuffer *output, - UfoRequisition *requisition) + UfoBuffer *output, + UfoRequisition *requisition) { UfoMetaballsTaskPrivate *priv; UfoGpuNode *node; + UfoProfiler *profiler; cl_command_queue cmd_queue; cl_mem out_mem; priv = UFO_METABALLS_TASK_GET_PRIVATE (task); - if (!priv->run_infinitely && (priv->current_iteration++) >= priv->num_iterations) + if (priv->current_iteration++ >= priv->num_iterations) return FALSE; node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task))); @@ -189,49 +162,37 @@ ufo_metaballs_task_generate (UfoTask *task, out_mem = ufo_buffer_get_device_array (output, cmd_queue); UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof(cl_mem), (cl_mem) &out_mem)); - UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, - priv->kernel, - 2, NULL, requisition->dims, NULL, - 0, NULL, NULL)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &priv->balls_mem)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_uint), &priv->num_balls)); - /* Update positions and velocities */ - for (guint j = 0; j < priv->num_balls; j++) { - const guint x = 2*j, y = 2*j + 1; + profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); + ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL); - priv->positions[x] += priv->velocities[x]; - priv->positions[y] += priv->velocities[y]; + /* Update positions and velocities */ + for (guint i = 0; i < priv->num_balls; i++) { + priv->balls[i].x += priv->balls[i].vx; + priv->balls[i].y += priv->balls[i].vy; - if (priv->positions[x] < 0 || priv->positions[x] > priv->width) - priv->velocities[x] = -priv->velocities[x]; + if (priv->balls[i].x < 0 || priv->balls[i].x > priv->width) + priv->balls[i].vx = -priv->balls[i].vx; - if (priv->positions[y] < 0 || priv->positions[y] > priv->height) - priv->velocities[y] = -priv->velocities[y]; + if (priv->balls[i].y < 0 || priv->balls[i].y > priv->height) + priv->balls[i].vy = -priv->balls[i].vy; } - UFO_RESOURCES_CHECK_CLERR (clEnqueueWriteBuffer (cmd_queue, - priv->positions_mem, + UFO_RESOURCES_CHECK_CLERR (clEnqueueWriteBuffer (cmd_queue, priv->balls_mem, CL_FALSE, - 0, priv->num_position_bytes, priv->positions, + 0, priv->num_balls * sizeof (Ball), priv->balls, 0, NULL, NULL)); - g_timer_stop(priv->timer); - - if (priv->frames_per_second > 0) { - const gdouble elapsed = g_timer_elapsed (priv->timer, NULL); - const gdouble delta = priv->seconds_per_frame - elapsed; - - if (delta > 0.0) - g_usleep (G_USEC_PER_SEC * ((gulong) delta)); - } - return TRUE; } static void ufo_metaballs_task_set_property (GObject *object, - guint property_id, - const GValue *value, - GParamSpec *pspec) + guint property_id, + const GValue *value, + GParamSpec *pspec) { UfoMetaballsTaskPrivate *priv = UFO_METABALLS_TASK_GET_PRIVATE (object); @@ -248,12 +209,6 @@ ufo_metaballs_task_set_property (GObject *object, case PROP_NUM_ITERATIONS: priv->num_iterations = g_value_get_uint(value); break; - case PROP_RUN_INFINITELY: - priv->run_infinitely = g_value_get_boolean(value); - break; - case PROP_FRAMES_PER_SECOND: - priv->frames_per_second = g_value_get_uint(value); - break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); break; @@ -262,9 +217,9 @@ ufo_metaballs_task_set_property (GObject *object, static void ufo_metaballs_task_get_property (GObject *object, - guint property_id, - GValue *value, - GParamSpec *pspec) + guint property_id, + GValue *value, + GParamSpec *pspec) { UfoMetaballsTaskPrivate *priv = UFO_METABALLS_TASK_GET_PRIVATE (object); @@ -281,12 +236,6 @@ ufo_metaballs_task_get_property (GObject *object, case PROP_NUM_ITERATIONS: g_value_set_uint(value, priv->num_iterations); break; - case PROP_RUN_INFINITELY: - g_value_set_boolean(value, priv->run_infinitely); - break; - case PROP_FRAMES_PER_SECOND: - g_value_set_uint(value, priv->frames_per_second); - break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); break; @@ -305,25 +254,12 @@ ufo_metaballs_task_finalize (GObject *object) priv->kernel = NULL; } - if (priv->positions_mem) { - UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->positions_mem)); - priv->positions_mem= NULL; - } - - if (priv->sizes_mem) { - UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->sizes_mem)); - priv->sizes_mem= NULL; + if (priv->balls_mem) { + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->balls_mem)); + priv->balls_mem = NULL; } - if (priv->context) { - UFO_RESOURCES_CHECK_CLERR (clReleaseContext (priv->context)); - priv->context = NULL; - } - - g_free (priv->sizes); - g_free (priv->positions); - g_free (priv->velocities); - g_timer_destroy (priv->timer); + g_free (priv->balls); G_OBJECT_CLASS (ufo_metaballs_task_parent_class)->finalize (object); } @@ -376,20 +312,6 @@ ufo_metaballs_task_class_init (UfoMetaballsTaskClass *klass) 1, G_MAXUINT, 1, G_PARAM_READWRITE); - properties[PROP_RUN_INFINITELY] = - g_param_spec_boolean("run-infinitely", - "Run infinitely", - "Run infinitely", - FALSE, - G_PARAM_READWRITE); - - properties[PROP_FRAMES_PER_SECOND] = - g_param_spec_uint("frames-per-second", - "Number of frames per second (0 for maximum possible rate)", - "Number of frames per second (0 for maximum possible rate)", - 0, G_MAXINT, 0, - G_PARAM_READWRITE); - for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (gobject_class, i, properties[i]); @@ -406,6 +328,4 @@ ufo_metaballs_task_init(UfoMetaballsTask *self) priv->height = 512; priv->num_balls = 1; priv->num_iterations = 1; - priv->run_infinitely = FALSE; - priv->frames_per_second = 0; } diff --git a/src/ufo-monitor-task.h b/src/ufo-monitor-task.h index 8a883e2..c03e6ce 100644 --- a/src/ufo-monitor-task.h +++ b/src/ufo-monitor-task.h @@ -51,3 +51,4 @@ GType ufo_monitor_task_get_type (void); G_END_DECLS #endif + diff --git a/src/ufo-ordfilt-task.c b/src/ufo-ordfilt-task.c index 944f640..2d2e69a 100644 --- a/src/ufo-ordfilt-task.c +++ b/src/ufo-ordfilt-task.c @@ -288,8 +288,6 @@ ufo_ordfilt_task_process (UfoTask *task, priv = UFO_ORDFILT_TASK_GET_PRIVATE (task); node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task))); cmd_queue = ufo_gpu_node_get_cmd_queue (node); - - priv = UFO_ORDFILT_TASK_GET_PRIVATE (task); compute_ordfilt (priv, inputs[0], inputs[1], output, cmd_queue); return TRUE; diff --git a/src/ufo-ordfilt-task.h b/src/ufo-ordfilt-task.h index 3b3d61f..b9bf7d0 100644 --- a/src/ufo-ordfilt-task.h +++ b/src/ufo-ordfilt-task.h @@ -63,4 +63,5 @@ GType ufo_ordfilt_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-pad-task.c b/src/ufo-pad-task.c index 389b20b..25ee831 100644 --- a/src/ufo-pad-task.c +++ b/src/ufo-pad-task.c @@ -24,12 +24,22 @@ #include "ufo-pad-task.h" -/** - * SECTION:ufo-pad-task - * @Short_description: Pad images to some extent - * @Title: pad - * - */ +typedef enum { + ADDRESS_NONE = CL_ADDRESS_NONE, + ADDRESS_CLAMP_TO_EDGE = CL_ADDRESS_CLAMP_TO_EDGE, + ADDRESS_CLAMP = CL_ADDRESS_CLAMP, + ADDRESS_REPEAT = CL_ADDRESS_REPEAT, + ADDRESS_MIRRORED_REPEAT = CL_ADDRESS_MIRRORED_REPEAT +} AddressingMode; + +static GEnumValue addressing_values[] = { + { ADDRESS_NONE, "ADDRESS_NONE", "none" }, + { ADDRESS_CLAMP_TO_EDGE, "ADDRESS_CLAMP_TO_EDGE", "clamp_to_edge" }, + { ADDRESS_CLAMP, "ADDRESS_CLAMP", "clamp" }, + { ADDRESS_REPEAT, "ADDRESS_REPEAT", "repeat" }, + { ADDRESS_MIRRORED_REPEAT, "ADDRESS_MIRRORED_REPEAT", "mirrored_repeat" }, + { 0, NULL, NULL} +}; struct _UfoPadTaskPrivate { /* OpenCL */ @@ -40,7 +50,7 @@ struct _UfoPadTaskPrivate { /* properties */ guint width, height; gint x, y; - cl_addressing_mode addressing_mode; + AddressingMode addressing_mode; }; static void ufo_task_interface_init (UfoTaskIface *iface); @@ -234,23 +244,7 @@ ufo_pad_task_set_property (GObject *object, priv->y = g_value_get_int (value); break; case PROP_ADDRESSING_MODE: - if (!g_strcmp0 (g_value_get_string (value), "none")) { - priv->addressing_mode = CL_ADDRESS_NONE; - } - else if (!g_strcmp0 (g_value_get_string (value), "clamp")) { - priv->addressing_mode = CL_ADDRESS_CLAMP; - } - else if (!g_strcmp0 (g_value_get_string (value), "clamp_to_edge")) { - priv->addressing_mode = CL_ADDRESS_CLAMP_TO_EDGE; - } - else if (!g_strcmp0 (g_value_get_string (value), "repeat")) { - priv->addressing_mode = CL_ADDRESS_REPEAT; - } else { - g_warning ("Invalid addressing mode \"%s\", "\ - "it has to be one of [\"none\", \"clamp\", "\ - "clamp_to_edge\", \"repeat\"]", - g_value_get_string (value)); - } + priv->addressing_mode = g_value_get_enum (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -280,20 +274,7 @@ ufo_pad_task_get_property (GObject *object, g_value_set_int (value, priv->y); break; case PROP_ADDRESSING_MODE: - switch (priv->addressing_mode) { - case CL_ADDRESS_NONE: - g_value_set_string (value, "none"); - break; - case CL_ADDRESS_CLAMP: - g_value_set_string (value, "clamp"); - break; - case CL_ADDRESS_CLAMP_TO_EDGE: - g_value_set_string (value, "clamp_to_edge"); - break; - case CL_ADDRESS_REPEAT: - g_value_set_string (value, "repeat"); - break; - } + g_value_set_enum (value, priv->addressing_mode); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -374,11 +355,12 @@ ufo_pad_task_class_init (UfoPadTaskClass *klass) G_PARAM_READWRITE); properties[PROP_ADDRESSING_MODE] = - g_param_spec_string ("addressing-mode", - "Outlier treatment", - "Outlier treatment from : \"none\", \"clamp\", \"clamp_to_edge\", \"repeat\"", - "clamp", - G_PARAM_READWRITE); + g_param_spec_enum ("addressing-mode", + "Outlier treatment (\"none\", \"clamp\", \"clamp_to_edge\", \"repeat\")", + "Outlier treatment (\"none\", \"clamp\", \"clamp_to_edge\", \"repeat\")", + g_enum_register_static ("addressing_mode", addressing_values), + CL_ADDRESS_CLAMP, + G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (gobject_class, i, properties[i]); diff --git a/src/ufo-read-task.c b/src/ufo-read-task.c index c5acd37..4d6d609 100644 --- a/src/ufo-read-task.c +++ b/src/ufo-read-task.c @@ -37,7 +37,7 @@ #include "readers/ufo-hdf5-reader.h" #endif -/* XXX: keep enum and types array in sync! */ +/* XXX: keep enum and values array in sync! */ typedef enum { TYPE_EDF, TYPE_RAW, @@ -50,16 +50,17 @@ typedef enum { TYPE_UNSPECIFIED } FileType; -static const gchar *types[] = { - "edf", - "raw", +static GEnumValue type_values[] = { + { TYPE_EDF, "TYPE_EDF", "edf" }, + { TYPE_RAW, "TYPE_RAW", "raw" }, #ifdef HAVE_TIFF - "tiff", + { TYPE_TIFF, "TYPE_TIFF", "tiff" }, #endif #ifdef WITH_HDF5 - "hdf5", + { TYPE_HDF5, "TYPE_HDF5", "hdf5" }, #endif - NULL + { TYPE_UNSPECIFIED, "TYPE_UNSPECIFIED", "unspecified" }, + { 0, NULL, NULL} }; struct _UfoReadTaskPrivate { @@ -350,7 +351,6 @@ ufo_read_task_set_property (GObject *object, case PROP_RAW_WIDTH: case PROP_RAW_HEIGHT: case PROP_RAW_BITDEPTH: - case PROP_RAW_OFFSET: { const gchar *prop_name; @@ -360,21 +360,11 @@ ufo_read_task_set_property (GObject *object, g_object_set (priv->raw_reader, prop_name, g_value_get_uint (value), NULL); } break; + case PROP_RAW_OFFSET: + g_object_set (priv->raw_reader, "offset", g_value_get_ulong (value), NULL); + break; case PROP_TYPE: - { - const gchar *type; - - type = g_value_get_string (value); - - for (guint i = 0; types[i] != NULL; i++) { - if (g_strcmp0 (type, types[i]) == 0) { - priv->type = (FileType) i; - return; - } - } - - g_warning ("File type `%s' not recognized", type); - } + priv->type = g_value_get_enum (value); break; default: @@ -431,10 +421,7 @@ ufo_read_task_get_property (GObject *object, } break; case PROP_TYPE: - if (priv->type != TYPE_UNSPECIFIED) - g_value_set_string (value, types[priv->type]); - else - g_value_set_string (value, "unspecified"); + g_value_set_enum (value, priv->type); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -503,94 +490,95 @@ ufo_read_task_class_init(UfoReadTaskClass *klass) gobject_class->finalize = ufo_read_task_finalize; properties[PROP_PATH] = - g_param_spec_string("path", + g_param_spec_string ("path", "Glob-style pattern.", "Glob-style pattern that describes the file path.", "*.tif", G_PARAM_READWRITE); properties[PROP_STEP] = - g_param_spec_uint("step", + g_param_spec_uint ("step", "Read every \"step\" file", "Read every \"step\" file", 1, G_MAXUINT, 1, G_PARAM_READWRITE); properties[PROP_ROI_Y] = - g_param_spec_uint("y", + g_param_spec_uint ("y", "Vertical coordinate", "Vertical coordinate from where to start reading the image", 0, G_MAXUINT, 0, G_PARAM_READWRITE); properties[PROP_ROI_HEIGHT] = - g_param_spec_uint("height", + g_param_spec_uint ("height", "Height", "Height of the region of interest to read", 0, G_MAXUINT, 0, G_PARAM_READWRITE); properties[PROP_ROI_STEP] = - g_param_spec_uint("y-step", + g_param_spec_uint ("y-step", "Read every \"step\" row", "Read every \"step\" row", 1, G_MAXUINT, 1, G_PARAM_READWRITE); properties[PROP_CONVERT] = - g_param_spec_boolean("convert", + g_param_spec_boolean ("convert", "Enable automatic conversion", "Enable automatic conversion of input data types to float", TRUE, G_PARAM_READWRITE); properties[PROP_START] = - g_param_spec_uint("start", + g_param_spec_uint ("start", "Offset to the first read file", "Offset to the first read file", 0, G_MAXUINT, 0, G_PARAM_READWRITE); properties[PROP_NUMBER] = - g_param_spec_uint("number", + g_param_spec_uint ("number", "Number of files that will be read at most", "Number of files that will be read at most", 0, G_MAXUINT, G_MAXUINT, G_PARAM_READWRITE); properties[PROP_RAW_WIDTH] = - g_param_spec_uint("raw-width", + g_param_spec_uint ("raw-width", "Width of raw image", "Width of raw image", 0, G_MAXUINT, G_MAXUINT, G_PARAM_READWRITE); properties[PROP_RAW_HEIGHT] = - g_param_spec_uint("raw-height", + g_param_spec_uint ("raw-height", "Height of raw image", "Height of raw image", 0, G_MAXUINT, G_MAXUINT, G_PARAM_READWRITE); properties[PROP_RAW_BITDEPTH] = - g_param_spec_uint("raw-bitdepth", + g_param_spec_uint ("raw-bitdepth", "Bitdepth of raw image", "Bitdepth of raw image", 0, G_MAXUINT, G_MAXUINT, G_PARAM_READWRITE); properties[PROP_RAW_OFFSET] = - g_param_spec_uint("raw-offset", + g_param_spec_ulong ("raw-offset", "Offset to the beginning of raw image in bytes", "Offset to the beginning of raw image in bytes", - 0, G_MAXUINT, G_MAXUINT, + 0, G_MAXULONG, 0, G_PARAM_READWRITE); properties[PROP_TYPE] = - g_param_spec_string ("type", + g_param_spec_enum ("type", "Override type detection based on extension", "Override type detection based on extension", - "unspecified", + g_enum_register_static ("type", type_values), + TYPE_UNSPECIFIED, G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) diff --git a/src/ufo-reduce-task.h b/src/ufo-reduce-task.h index bb55f6b..6753844 100644 --- a/src/ufo-reduce-task.h +++ b/src/ufo-reduce-task.h @@ -63,4 +63,5 @@ GType ufo_reduce_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-refeed-task.h b/src/ufo-refeed-task.h index ea04780..0d883f9 100644 --- a/src/ufo-refeed-task.h +++ b/src/ufo-refeed-task.h @@ -50,4 +50,5 @@ GType ufo_refeed_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-remove-circle-task.c b/src/ufo-remove-circle-task.c index ff29bff..bfd09c3 100644 --- a/src/ufo-remove-circle-task.c +++ b/src/ufo-remove-circle-task.c @@ -223,7 +223,7 @@ void remove_inner_circle(URCS *src, URCS *dst) } dst->nb_elt = (float) counter; if (counter >= src->nb_elt && min_i != -1) - g_print("Remove Circle : Some memory corruption occured. This is a bug " + g_print("Remove Circle: Some memory corruption occurred. This is a bug " " in the software that needs fixing\n"); } diff --git a/src/ufo-remove-circle-task.h b/src/ufo-remove-circle-task.h index c7e5dcf..b2c2d7f 100644 --- a/src/ufo-remove-circle-task.h +++ b/src/ufo-remove-circle-task.h @@ -63,4 +63,5 @@ GType ufo_remove_circle_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-replicate-task.h b/src/ufo-replicate-task.h index af39d88..d0ec663 100644 --- a/src/ufo-replicate-task.h +++ b/src/ufo-replicate-task.h @@ -63,4 +63,5 @@ GType ufo_replicate_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-rescale-task.c b/src/ufo-rescale-task.c index 020df7d..b92dda2 100644 --- a/src/ufo-rescale-task.c +++ b/src/ufo-rescale-task.c @@ -26,10 +26,21 @@ #include "ufo-rescale-task.h" +typedef enum { + INTERPOLATION_NEAREST = CL_FILTER_NEAREST, + INTERPOLATION_LINEAR = CL_FILTER_LINEAR, +} Interpolation; + +static GEnumValue interpolation_values[] = { + { INTERPOLATION_NEAREST, "INTERPOLATION_NEAREST", "nearest" }, + { INTERPOLATION_LINEAR, "INTERPOLATION_LINEAR", "linear" }, + { 0, NULL, NULL} +}; + struct _UfoRescaleTaskPrivate { cl_context context; cl_kernel kernel; - cl_filter_mode filter_mode; + Interpolation interpolation; gfloat x_factor; gfloat y_factor; guint width; @@ -80,7 +91,7 @@ ufo_rescale_task_setup (UfoTask *task, priv->sampler = clCreateSampler (priv->context, (cl_bool) FALSE, CL_ADDRESS_NONE, - priv->filter_mode, + (cl_filter_mode) priv->interpolation, &err); UFO_RESOURCES_CHECK_CLERR (err); @@ -198,11 +209,7 @@ ufo_rescale_task_set_property (GObject *object, priv->height = g_value_get_uint (value); break; case PROP_INTERPOLATION: - if (!g_strcmp0 (g_value_get_string (value), "nearest")) { - priv->filter_mode = CL_FILTER_NEAREST; - } else if (!g_strcmp0 (g_value_get_string (value), "linear")) { - priv->filter_mode = CL_FILTER_LINEAR; - } + priv->interpolation = g_value_get_enum (value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -238,14 +245,7 @@ ufo_rescale_task_get_property (GObject *object, g_value_set_uint (value, priv->height); break; case PROP_INTERPOLATION: - switch (priv->filter_mode) { - case CL_FILTER_NEAREST: - g_value_set_string (value, "nearest"); - break; - case CL_FILTER_LINEAR: - g_value_set_string (value, "linear"); - break; - } + g_value_set_enum (value, priv->interpolation); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); @@ -318,11 +318,11 @@ ufo_rescale_task_class_init (UfoRescaleTaskClass *klass) G_PARAM_READWRITE); properties[PROP_INTERPOLATION] = - g_param_spec_string ("interpolation", - "Interpolation mode", - "Interpolation mode from: \"nearest\", \"linear\"", - "linear", - G_PARAM_READWRITE); + g_param_spec_enum ("interpolation", + "Interpolation mode (\"nearest\", \"linear\")", + "Interpolation mode (\"nearest\", \"linear\")", + g_enum_register_static ("interpolation", interpolation_values), + INTERPOLATION_LINEAR, G_PARAM_READWRITE); properties[PROP_WIDTH] = g_param_spec_uint ("width", @@ -352,5 +352,5 @@ ufo_rescale_task_init(UfoRescaleTask *self) self->priv->y_factor = 2.0f; self->priv->width = 0; self->priv->height = 0; - self->priv->filter_mode = CL_FILTER_LINEAR; + self->priv->interpolation = INTERPOLATION_LINEAR; } diff --git a/src/ufo-retrieve-phase-task.c b/src/ufo-retrieve-phase-task.c index a75fe47..e681b6c 100644 --- a/src/ufo-retrieve-phase-task.c +++ b/src/ufo-retrieve-phase-task.c @@ -30,7 +30,7 @@ #define IS_POW_OF_2(x) !(x & (x - 1)) typedef enum { - METHOD_TIE, + METHOD_TIE = 0, METHOD_CTF, METHOD_CTFHALFSINE, METHOD_QP, @@ -39,6 +39,16 @@ typedef enum { N_METHODS } Method; +static GEnumValue method_values[] = { + { METHOD_TIE, "METHOD_TIE", "tie" }, + { METHOD_CTF, "METHOD_CTF", "ctf" }, + { METHOD_CTFHALFSINE, "METHOD_CTFHALFSINE", "ctfhalfsine" }, + { METHOD_QP, "METHOD_QP", "qp" }, + { METHOD_QPHALFSINE, "METHOD_QPHALFSINE", "qphalfsine" }, + { METHOD_QP2, "METHOD_QP2", "qp2" }, + { 0, NULL, NULL} +}; + struct _UfoRetrievePhaseTaskPrivate { Method method; gfloat energy; @@ -85,8 +95,8 @@ ufo_retrieve_phase_task_new (void) static void ufo_retrieve_phase_task_setup (UfoTask *task, - UfoResources *resources, - GError **error) + UfoResources *resources, + GError **error) { UfoRetrievePhaseTaskPrivate *priv; gfloat lambda; @@ -161,9 +171,9 @@ ufo_filter_task_get_mode (UfoTask *task) static gboolean ufo_retrieve_phase_task_process (UfoTask *task, - UfoBuffer **inputs, - UfoBuffer *output, - UfoRequisition *requisition) + UfoBuffer **inputs, + UfoBuffer *output, + UfoRequisition *requisition) { UfoRetrievePhaseTaskPrivate *priv; UfoGpuNode *node; @@ -208,26 +218,15 @@ ufo_retrieve_phase_task_process (UfoTask *task, static void ufo_retrieve_phase_task_get_property (GObject *object, - guint property_id, - GValue *value, - GParamSpec *pspec) + guint property_id, + GValue *value, + GParamSpec *pspec) { UfoRetrievePhaseTaskPrivate *priv = UFO_RETRIEVE_PHASE_TASK_GET_PRIVATE (object); switch (property_id) { case PROP_METHOD: - if (priv->method == METHOD_TIE) - g_value_set_string (value, "tie"); - else if (priv->method == METHOD_CTF) - g_value_set_string (value, "ctf"); - else if (priv->method == METHOD_CTFHALFSINE) - g_value_set_string (value, "ctfhalfsine"); - else if (priv->method == METHOD_QP) - g_value_set_string (value, "qp"); - else if (priv->method == METHOD_QPHALFSINE) - g_value_set_string (value, "qphalfsine"); - else if (priv->method == METHOD_QP2) - g_value_set_string (value, "qp2"); + g_value_set_enum (value, priv->method); break; case PROP_ENERGY: g_value_set_float (value, priv->energy); @@ -252,26 +251,15 @@ ufo_retrieve_phase_task_get_property (GObject *object, static void ufo_retrieve_phase_task_set_property (GObject *object, - guint property_id, - const GValue *value, - GParamSpec *pspec) + guint property_id, + const GValue *value, + GParamSpec *pspec) { UfoRetrievePhaseTaskPrivate *priv = UFO_RETRIEVE_PHASE_TASK_GET_PRIVATE (object); switch (property_id) { case PROP_METHOD: - if (!g_strcmp0 (g_value_get_string (value), "tie")) - priv->method = METHOD_TIE; - else if (!g_strcmp0 (g_value_get_string (value), "ctf")) - priv->method = METHOD_CTF; - else if (!g_strcmp0 (g_value_get_string (value), "ctfhalfsine")) - priv->method = METHOD_CTFHALFSINE; - else if (!g_strcmp0 (g_value_get_string (value), "qp")) - priv->method = METHOD_QP; - else if (!g_strcmp0 (g_value_get_string (value), "qphalfsine")) - priv->method = METHOD_QPHALFSINE; - else if (!g_strcmp0 (g_value_get_string (value), "qp2")) - priv->method = METHOD_QP2; + priv->method = g_value_get_enum (value); break; case PROP_ENERGY: priv->energy = g_value_get_float (value); @@ -349,10 +337,11 @@ ufo_retrieve_phase_task_class_init (UfoRetrievePhaseTaskClass *klass) gobject_class->finalize = ufo_retrieve_phase_task_finalize; properties[PROP_METHOD] = - g_param_spec_string ("method", - "Name of method", - "Method.", - "tie", + g_param_spec_enum ("method", + "Method name", + "Method name", + g_enum_register_static ("method", method_values), + METHOD_TIE, G_PARAM_READWRITE); properties[PROP_ENERGY] = diff --git a/src/ufo-ring-pattern-task.h b/src/ufo-ring-pattern-task.h index 131fbb6..824062f 100644 --- a/src/ufo-ring-pattern-task.h +++ b/src/ufo-ring-pattern-task.h @@ -63,4 +63,5 @@ GType ufo_ring_pattern_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-ringwriter-task.h b/src/ufo-ringwriter-task.h index bacf3ef..4748e22 100644 --- a/src/ufo-ringwriter-task.h +++ b/src/ufo-ringwriter-task.h @@ -63,4 +63,5 @@ GType ufo_ringwriter_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-segment-task.c b/src/ufo-segment-task.c new file mode 100644 index 0000000..a113cfd --- /dev/null +++ b/src/ufo-segment-task.c @@ -0,0 +1,437 @@ +/* + * Copyright (C) 2016 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/>. + */ + +#ifdef __APPLE__ +#include <OpenCL/cl.h> +#else +#include <CL/cl.h> +#endif + +#include "ufo-segment-task.h" + +#define MAX_SEGMENTS 16 +#define MAX_LABELS 32768 + +/* + * Data structures + * --------------- + * + * accumulator: stores the number of hits for the current segment. + * + * bitmap: stores if the pixel belongs to a segment. + * + * + * Algorithm + * --------- + * + * 1. For each segment from the label field we generate arrays of labels, + * essentially the position. + * + * 2. For each segment + * + * 2.1. we use the label array to call the `walk` kernel which starts random + * walks from each label position. Going to a new position increases the + * accumulator array at that position. + * + * 2.2. we use the accumulator to fill out the bitmap. One entry contains 32 + * successive bits (MSB) denoting if the corresponding pixel belongs to that + * particular segment. To avoid too many buffers we use a linearized array, + * i.e. it has four dimensions: width x height x slices x segment. + * + * 3. For the current slice, we look up the segment bitmaps and store the + * segment in the output. + * + */ + +typedef struct { + int x; + int y; +} Label; + +typedef struct { + Label *labels; + guint num_labels; +} SegmentLabels; + +struct _UfoSegmentTaskPrivate { + cl_context context; + cl_kernel walk; + cl_kernel render; + cl_kernel threshold; + cl_mem bitmap; + cl_mem label_map; + guint num_slices; + guint num_segments; + guint current; +}; + +static void ufo_task_interface_init (UfoTaskIface *iface); + +G_DEFINE_TYPE_WITH_CODE (UfoSegmentTask, ufo_segment_task, UFO_TYPE_TASK_NODE, + G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK, + ufo_task_interface_init)) + +#define UFO_SEGMENT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_SEGMENT_TASK, UfoSegmentTaskPrivate)) + +UfoNode * +ufo_segment_task_new (void) +{ + return UFO_NODE (g_object_new (UFO_TYPE_SEGMENT_TASK, NULL)); +} + +static void +ufo_segment_task_setup (UfoTask *task, + UfoResources *resources, + GError **error) +{ + UfoSegmentTaskPrivate *priv; + + priv = UFO_SEGMENT_TASK_GET_PRIVATE (task); + + priv->context = ufo_resources_get_context (resources); + priv->walk = ufo_resources_get_kernel (resources, "segment.cl", "walk", error); + priv->render = ufo_resources_get_kernel (resources, "segment.cl", "render", error); + priv->threshold = ufo_resources_get_kernel (resources, "segment.cl", "threshold", error); + + UFO_RESOURCES_CHECK_CLERR (clRetainContext (priv->context)); + + if (priv->walk != NULL) + UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->walk)); + + if (priv->render != NULL) + UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->render)); + + if (priv->threshold != NULL) + UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->threshold)); +} + +static void +ufo_segment_task_get_requisition (UfoTask *task, + UfoBuffer **inputs, + UfoRequisition *requisition) +{ + UfoRequisition label_req; + + ufo_buffer_get_requisition (inputs[0], requisition); + requisition->n_dims = 2; + + /* ensure inputs match */ + ufo_buffer_get_requisition (inputs[1], &label_req); + + if ((label_req.dims[0] != requisition->dims[0]) || + (label_req.dims[1] != requisition->dims[1])) { + g_warning ("Label field and input dimensions do not match ([%zu, %zu] != [%zu, %zu])", + label_req.dims[0], label_req.dims[1], requisition->dims[0], requisition->dims[1]); + } +} + +static guint +ufo_segment_task_get_num_inputs (UfoTask *task) +{ + return 2; +} + +static guint +ufo_segment_task_get_num_dimensions (UfoTask *task, + guint input) +{ + if (input == 0) + return 3; + + return 2; +} + +static UfoTaskMode +ufo_segment_task_get_mode (UfoTask *task) +{ + return UFO_TASK_MODE_REDUCTOR | UFO_TASK_MODE_GPU; +} + +static SegmentLabels * +extract_labels (UfoBuffer *buffer) +{ + UfoRequisition requisition; + gfloat *data; + SegmentLabels *result; + guint width; + guint height; + + ufo_buffer_get_requisition (buffer, &requisition); + data = ufo_buffer_get_host_array (buffer, NULL); + result = g_malloc0 (sizeof (SegmentLabels) * MAX_SEGMENTS); + width = requisition.dims[0]; + height = requisition.dims[1]; + + for (guint x = 0; x < width; x++) { + for (guint y = 0; y < height; y++) { + gint v = (gint) data[y * width + x] - 1; + + if (v >= 0 && v < MAX_SEGMENTS) { + guint n; + + if (result[v].labels == NULL) { + result[v].labels = g_malloc0 (sizeof (Label) * MAX_LABELS); + result[v].num_labels = 0; + } + + n = result[v].num_labels; + + if (n < MAX_LABELS) { + result[v].labels[n].x = x; + result[v].labels[n].y = y; + result[v].num_labels++; + } + } + } + } + + return result; +} + +static gboolean +ufo_segment_task_process (UfoTask *task, + UfoBuffer **inputs, + UfoBuffer *output, + UfoRequisition *requisition) +{ + UfoSegmentTaskPrivate *priv; + UfoRequisition in_req; + UfoGpuNode *node; + UfoProfiler *profiler; + cl_command_queue cmd_queue; + SegmentLabels *prelabeled; + gfloat *random_host; + cl_mem random_device; + cl_mem prelabeled_device; + cl_mem accumulator; + cl_mem slices; + cl_int error; + guint width; + guint height; + guint segment; + guint label_map_host[MAX_LABELS]; + gsize work_size[3]; + guint16 fill_pattern = 0; + + priv = UFO_SEGMENT_TASK_GET_PRIVATE (task); + node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task))); + cmd_queue = ufo_gpu_node_get_cmd_queue (node); + + /* extract labels */ + prelabeled = extract_labels (inputs[1]); + priv->num_segments = 0; + + for (guint s = 0; s < MAX_SEGMENTS; s++) { + if (prelabeled[s].num_labels > 0) + priv->num_segments++; + } + + /* create uniformly distributed data */ + random_host = g_malloc0 (32768 * sizeof (gfloat)); + + for (guint i = 0; i < 32768; i++) + random_host[i] = (gfloat) g_random_double (); + + random_device = clCreateBuffer (priv->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + 32768 * sizeof (gfloat), random_host, &error); + + UFO_RESOURCES_CHECK_CLERR (error); + + ufo_buffer_get_requisition (inputs[0], &in_req); + + width = in_req.dims[0]; + height = in_req.dims[1]; + priv->current = in_req.dims[2]; + priv->num_slices = in_req.dims[2]; + + profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); + slices = ufo_buffer_get_device_array (inputs[0], cmd_queue); + + /* create bitmap memory */ + priv->bitmap = clCreateBuffer (priv->context, CL_MEM_READ_WRITE, + sizeof (guint32) * width / 32 * height * priv->num_slices * priv->num_segments, + NULL, &error); + + UFO_RESOURCES_CHECK_CLERR (error); + + /* create and initialize accumulator memory */ + accumulator = clCreateBuffer (priv->context, CL_MEM_READ_WRITE, + sizeof (guint16) * width * height * priv->num_slices, + NULL, &error); + UFO_RESOURCES_CHECK_CLERR (error); + + prelabeled_device = clCreateBuffer (priv->context, CL_MEM_READ_ONLY, + MAX_LABELS * sizeof (Label), + NULL, &error); + UFO_RESOURCES_CHECK_CLERR (error); + + segment = 0; + + for (guint s = 0; s < MAX_SEGMENTS; s++) { + if (prelabeled[s].num_labels == 0) + continue; + + /* upload labels for this segment */ + UFO_RESOURCES_CHECK_CLERR (clEnqueueWriteBuffer (cmd_queue, prelabeled_device, CL_TRUE, + 0, prelabeled[s].num_labels * sizeof (Label), + prelabeled[s].labels, + 0, NULL, NULL)); + + /* reset accumulator to zero */ + UFO_RESOURCES_CHECK_CLERR (clEnqueueFillBuffer (cmd_queue, accumulator, + &fill_pattern, sizeof (fill_pattern), + 0, width * height * priv->num_slices * sizeof (guint16), + 0, NULL, NULL)); + + /* start random walk */ + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 0, sizeof (cl_mem), &slices)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 1, sizeof (cl_mem), &accumulator)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 2, sizeof (cl_mem), &prelabeled_device)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 3, sizeof (guint), &width)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 4, sizeof (guint), &height)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 5, sizeof (guint), &priv->num_slices)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->walk, 6, sizeof (cl_mem), &random_device)); + + work_size[0] = prelabeled[s].num_labels; + ufo_profiler_call (profiler, cmd_queue, priv->walk, 1, work_size, NULL); + + /* threshold and record hit in bitmap */ + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->threshold, 0, sizeof (cl_mem), &accumulator)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->threshold, 1, sizeof (cl_mem), &priv->bitmap)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->threshold, 2, sizeof (guint), &segment)); + + work_size[0] = width / 32; + work_size[1] = height; + work_size[2] = priv->num_slices; + ufo_profiler_call (profiler, cmd_queue, priv->threshold, 3, work_size, NULL); + + label_map_host[segment] = s + 1; + segment++; + } + + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (accumulator)); + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (prelabeled_device)); + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (random_device)); + + priv->label_map = clCreateBuffer (priv->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + MAX_LABELS * sizeof (guint), + label_map_host, &error); + UFO_RESOURCES_CHECK_CLERR (error); + + for (guint s = 0; s < MAX_SEGMENTS; s++) + g_free (prelabeled[s].labels); + + g_free (prelabeled); + g_free (random_host); + + return TRUE; +} + +static gboolean +ufo_segment_task_generate (UfoTask *task, + UfoBuffer *output, + UfoRequisition *requisition) +{ + UfoSegmentTaskPrivate *priv; + UfoProfiler *profiler; + UfoGpuNode *node; + cl_command_queue cmd_queue; + cl_mem out_mem; + guint slice; + gsize work_size[3]; + + priv = UFO_SEGMENT_TASK_GET_PRIVATE (task); + + if (priv->current == 0) { + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->bitmap)); + UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject (priv->label_map)); + return FALSE; + } + + 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); + slice = priv->num_slices - priv->current; + + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 0, sizeof (cl_mem), &priv->bitmap)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 1, sizeof (cl_mem), &out_mem)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 2, sizeof (cl_mem), &priv->label_map)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 3, sizeof (guint), &slice)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 4, sizeof (guint), &priv->num_segments)); + UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->render, 5, sizeof (guint), &priv->num_slices)); + + work_size[0] = requisition->dims[0] / 32; + work_size[1] = requisition->dims[1]; + + profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task)); + ufo_profiler_call (profiler, cmd_queue, priv->render, 2, work_size, NULL); + + priv->current--; + return TRUE; +} + +static void +ufo_segment_task_finalize (GObject *object) +{ + UfoSegmentTaskPrivate *priv; + + priv = UFO_SEGMENT_TASK_GET_PRIVATE (object); + + if (priv->walk != NULL) + UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->walk)); + + if (priv->render != NULL) + UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->render)); + + if (priv->threshold != NULL) + UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->threshold)); + + UFO_RESOURCES_CHECK_CLERR (clReleaseContext (priv->context)); + + G_OBJECT_CLASS (ufo_segment_task_parent_class)->finalize (object); +} + +static void +ufo_task_interface_init (UfoTaskIface *iface) +{ + iface->setup = ufo_segment_task_setup; + iface->get_num_inputs = ufo_segment_task_get_num_inputs; + iface->get_num_dimensions = ufo_segment_task_get_num_dimensions; + iface->get_mode = ufo_segment_task_get_mode; + iface->get_requisition = ufo_segment_task_get_requisition; + iface->process = ufo_segment_task_process; + iface->generate = ufo_segment_task_generate; +} + +static void +ufo_segment_task_class_init (UfoSegmentTaskClass *klass) +{ + GObjectClass *oclass = G_OBJECT_CLASS (klass); + + oclass->finalize = ufo_segment_task_finalize; + + g_type_class_add_private (oclass, sizeof(UfoSegmentTaskPrivate)); +} + +static void +ufo_segment_task_init(UfoSegmentTask *self) +{ + self->priv = UFO_SEGMENT_TASK_GET_PRIVATE(self); +} diff --git a/src/ufo-segment-task.h b/src/ufo-segment-task.h new file mode 100644 index 0000000..8c97997 --- /dev/null +++ b/src/ufo-segment-task.h @@ -0,0 +1,54 @@ +/* + * 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_SEGMENT_TASK_H +#define __UFO_SEGMENT_TASK_H + +#include <ufo/ufo.h> + +G_BEGIN_DECLS + +#define UFO_TYPE_SEGMENT_TASK (ufo_segment_task_get_type()) +#define UFO_SEGMENT_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_SEGMENT_TASK, UfoSegmentTask)) +#define UFO_IS_SEGMENT_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_SEGMENT_TASK)) +#define UFO_SEGMENT_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_SEGMENT_TASK, UfoSegmentTaskClass)) +#define UFO_IS_SEGMENT_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_SEGMENT_TASK)) +#define UFO_SEGMENT_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_SEGMENT_TASK, UfoSegmentTaskClass)) + +typedef struct _UfoSegmentTask UfoSegmentTask; +typedef struct _UfoSegmentTaskClass UfoSegmentTaskClass; +typedef struct _UfoSegmentTaskPrivate UfoSegmentTaskPrivate; + +struct _UfoSegmentTask { + UfoTaskNode parent_instance; + + UfoSegmentTaskPrivate *priv; +}; + +struct _UfoSegmentTaskClass { + UfoTaskNodeClass parent_class; +}; + +UfoNode *ufo_segment_task_new (void); +GType ufo_segment_task_get_type (void); + +G_END_DECLS + +#endif + diff --git a/src/ufo-slice-task.h b/src/ufo-slice-task.h index e81213a..8b7f54d 100644 --- a/src/ufo-slice-task.h +++ b/src/ufo-slice-task.h @@ -63,4 +63,5 @@ GType ufo_slice_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-stack-task.h b/src/ufo-stack-task.h index 17b1bc3..b0ed154 100644 --- a/src/ufo-stack-task.h +++ b/src/ufo-stack-task.h @@ -63,4 +63,5 @@ GType ufo_stack_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-stdin-task.c b/src/ufo-stdin-task.c index 8101c9b..b6e3dfd 100644 --- a/src/ufo-stdin-task.c +++ b/src/ufo-stdin-task.c @@ -22,9 +22,9 @@ struct _UfoStdinTaskPrivate { - guint width; - guint height; - guint bytes_per_pixel; + gsize width; + gsize height; + gsize bytes_per_pixel; UfoBufferDepth bitdepth; gboolean convert; }; @@ -122,10 +122,10 @@ ufo_stdin_task_set_property (GObject *object, switch (property_id) { case PROP_WIDTH: - priv->width = g_value_get_uint (value); + priv->width = (gsize) g_value_get_uint (value); break; case PROP_HEIGHT: - priv->height = g_value_get_uint (value); + priv->height = (gsize) g_value_get_uint (value); break; case PROP_BITDEPTH: switch (g_value_get_uint (value)) { @@ -164,10 +164,10 @@ ufo_stdin_task_get_property (GObject *object, switch (property_id) { case PROP_WIDTH: - g_value_set_uint (value, priv->width); + g_value_set_uint (value, (guint) priv->width); break; case PROP_HEIGHT: - g_value_set_uint (value, priv->height); + g_value_set_uint (value, (guint) priv->height); break; case PROP_BITDEPTH: g_value_set_uint (value, priv->bitdepth); diff --git a/src/ufo-stdin-task.h b/src/ufo-stdin-task.h index 947dad5..77eb87c 100644 --- a/src/ufo-stdin-task.h +++ b/src/ufo-stdin-task.h @@ -50,4 +50,5 @@ GType ufo_stdin_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + diff --git a/src/ufo-stdout-task.h b/src/ufo-stdout-task.h index ca89103..7224065 100644 --- a/src/ufo-stdout-task.h +++ b/src/ufo-stdout-task.h @@ -50,4 +50,5 @@ GType ufo_stdout_task_get_type (void); G_END_DECLS -#endif
\ No newline at end of file +#endif + |