diff options
Diffstat (limited to 'src/kernels')
| -rw-r--r-- | src/kernels/CMakeFiles/CMakeDirectoryInformation.cmake | 16 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/DependInfo.cmake | 19 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/build.make | 105 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/cmake_clean.cmake | 12 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/depend.internal | 3 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/depend.make | 3 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/burst.dir/progress.make | 2 | ||||
| -rw-r--r-- | src/kernels/CMakeFiles/progress.marks | 1 | ||||
| -rw-r--r-- | src/kernels/CTestTestfile.cmake | 6 | ||||
| -rw-r--r-- | src/kernels/Makefile | 212 | ||||
| -rw-r--r-- | src/kernels/backproject.cl | 6 | ||||
| -rw-r--r-- | src/kernels/center_kernel.cl | 478 | ||||
| -rw-r--r-- | src/kernels/cmake_install.cmake | 684 | ||||
| -rw-r--r-- | src/kernels/forwardproject.cl | 6 | ||||
| -rw-r--r-- | src/kernels/lamino_kernel.cl | 478 | ||||
| -rw-r--r-- | src/kernels/roll_kernel.cl | 478 | ||||
| -rw-r--r-- | src/kernels/stacked-backproject.cl | 287 | ||||
| -rw-r--r-- | src/kernels/stacked-forwardproject.cl | 274 | ||||
| -rw-r--r-- | src/kernels/z_kernel.cl | 478 |
19 files changed, 3510 insertions, 38 deletions
diff --git a/src/kernels/CMakeFiles/CMakeDirectoryInformation.cmake b/src/kernels/CMakeFiles/CMakeDirectoryInformation.cmake new file mode 100644 index 0000000..b70123c --- /dev/null +++ b/src/kernels/CMakeFiles/CMakeDirectoryInformation.cmake @@ -0,0 +1,16 @@ +# CMAKE generated file: DO NOT EDIT! +# Generated by "Unix Makefiles" Generator, CMake Version 3.16 + +# Relative path conversion top directories. +set(CMAKE_RELATIVE_PATH_TOP_SOURCE "/ccpi/repos/ufo-filters") +set(CMAKE_RELATIVE_PATH_TOP_BINARY "/ccpi/repos/ufo-filters") + +# Force unix paths in dependencies. +set(CMAKE_FORCE_UNIX_PATHS 1) + + +# The C and CXX include file regular expressions for this directory. +set(CMAKE_C_INCLUDE_REGEX_SCAN "^.*$") +set(CMAKE_C_INCLUDE_REGEX_COMPLAIN "^$") +set(CMAKE_CXX_INCLUDE_REGEX_SCAN ${CMAKE_C_INCLUDE_REGEX_SCAN}) +set(CMAKE_CXX_INCLUDE_REGEX_COMPLAIN ${CMAKE_C_INCLUDE_REGEX_COMPLAIN}) diff --git a/src/kernels/CMakeFiles/burst.dir/DependInfo.cmake b/src/kernels/CMakeFiles/burst.dir/DependInfo.cmake new file mode 100644 index 0000000..edc37fc --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/DependInfo.cmake @@ -0,0 +1,19 @@ +# The set of languages for which implicit dependencies are needed: +set(CMAKE_DEPENDS_LANGUAGES + ) +# The set of files for implicit dependencies of each language: + +# Pairs of files generated by the same build rule. +set(CMAKE_MULTIPLE_OUTPUT_PAIRS + "/ccpi/repos/ufo-filters/src/kernels/center_kernel.cl" "/ccpi/repos/ufo-filters/src/kernels/z_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/lamino_kernel.cl" "/ccpi/repos/ufo-filters/src/kernels/z_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/roll_kernel.cl" "/ccpi/repos/ufo-filters/src/kernels/z_kernel.cl" + ) + + +# Targets to which this target links. +set(CMAKE_TARGET_LINKED_INFO_FILES + ) + +# Fortran module output directory. +set(CMAKE_Fortran_TARGET_MODULE_DIR "") diff --git a/src/kernels/CMakeFiles/burst.dir/build.make b/src/kernels/CMakeFiles/burst.dir/build.make new file mode 100644 index 0000000..19d9eef --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/build.make @@ -0,0 +1,105 @@ +# CMAKE generated file: DO NOT EDIT! +# Generated by "Unix Makefiles" Generator, CMake Version 3.16 + +# Delete rule output on recipe failure. +.DELETE_ON_ERROR: + + +#============================================================================= +# Special targets provided by cmake. + +# Disable implicit rules so canonical targets will work. +.SUFFIXES: + + +# Remove some rules from gmake that .SUFFIXES does not remove. +SUFFIXES = + +.SUFFIXES: .hpux_make_needs_suffix_list + + +# Suppress display of executed commands. +$(VERBOSE).SILENT: + + +# A target that is always out of date. +cmake_force: + +.PHONY : cmake_force + +#============================================================================= +# Set environment variables for the build. + +# The shell in which to execute make rules. +SHELL = /bin/sh + +# The CMake executable. +CMAKE_COMMAND = /usr/bin/cmake + +# The command to remove a file. +RM = /usr/bin/cmake -E remove -f + +# Escaping for special characters. +EQUALS = = + +# The top-level source directory on which CMake was run. +CMAKE_SOURCE_DIR = /ccpi/repos/ufo-filters + +# The top-level build directory on which CMake was run. +CMAKE_BINARY_DIR = /ccpi/repos/ufo-filters + +# Utility rule file for burst. + +# Include the progress variables for this target. +include src/kernels/CMakeFiles/burst.dir/progress.make + +src/kernels/CMakeFiles/burst: src/kernels/z_kernel.cl +src/kernels/CMakeFiles/burst: src/kernels/center_kernel.cl +src/kernels/CMakeFiles/burst: src/kernels/lamino_kernel.cl +src/kernels/CMakeFiles/burst: src/kernels/roll_kernel.cl + + +src/kernels/z_kernel.cl: src/kernels/tools/make_burst_kernels.py +src/kernels/z_kernel.cl: src/kernels/templates/common.in +src/kernels/z_kernel.cl: src/kernels/templates/definitions.in +src/kernels/z_kernel.cl: src/kernels/templates/z_template.in +src/kernels/z_kernel.cl: src/kernels/templates/center_template.in +src/kernels/z_kernel.cl: src/kernels/templates/lamino_template.in +src/kernels/z_kernel.cl: src/kernels/templates/roll_template.in + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --blue --bold --progress-dir=/ccpi/repos/ufo-filters/CMakeFiles --progress-num=$(CMAKE_PROGRESS_1) "Generating burst backprojection kernels" + cd /ccpi/repos/ufo-filters/src/kernels && /usr/bin/python /ccpi/repos/ufo-filters/src/kernels/tools/make_burst_kernels.py /ccpi/repos/ufo-filters/src/kernels/templates/z_template.in 1 2 4 8 16 > z_kernel.cl + cd /ccpi/repos/ufo-filters/src/kernels && /usr/bin/python /ccpi/repos/ufo-filters/src/kernels/tools/make_burst_kernels.py /ccpi/repos/ufo-filters/src/kernels/templates/center_template.in 1 2 4 8 16 > center_kernel.cl + cd /ccpi/repos/ufo-filters/src/kernels && /usr/bin/python /ccpi/repos/ufo-filters/src/kernels/tools/make_burst_kernels.py /ccpi/repos/ufo-filters/src/kernels/templates/lamino_template.in 1 2 4 8 16 > lamino_kernel.cl + cd /ccpi/repos/ufo-filters/src/kernels && /usr/bin/python /ccpi/repos/ufo-filters/src/kernels/tools/make_burst_kernels.py /ccpi/repos/ufo-filters/src/kernels/templates/roll_template.in 1 2 4 8 16 > roll_kernel.cl + +src/kernels/center_kernel.cl: src/kernels/z_kernel.cl + @$(CMAKE_COMMAND) -E touch_nocreate src/kernels/center_kernel.cl + +src/kernels/lamino_kernel.cl: src/kernels/z_kernel.cl + @$(CMAKE_COMMAND) -E touch_nocreate src/kernels/lamino_kernel.cl + +src/kernels/roll_kernel.cl: src/kernels/z_kernel.cl + @$(CMAKE_COMMAND) -E touch_nocreate src/kernels/roll_kernel.cl + +burst: src/kernels/CMakeFiles/burst +burst: src/kernels/z_kernel.cl +burst: src/kernels/center_kernel.cl +burst: src/kernels/lamino_kernel.cl +burst: src/kernels/roll_kernel.cl +burst: src/kernels/CMakeFiles/burst.dir/build.make + +.PHONY : burst + +# Rule to build all files generated by this target. +src/kernels/CMakeFiles/burst.dir/build: burst + +.PHONY : src/kernels/CMakeFiles/burst.dir/build + +src/kernels/CMakeFiles/burst.dir/clean: + cd /ccpi/repos/ufo-filters/src/kernels && $(CMAKE_COMMAND) -P CMakeFiles/burst.dir/cmake_clean.cmake +.PHONY : src/kernels/CMakeFiles/burst.dir/clean + +src/kernels/CMakeFiles/burst.dir/depend: + cd /ccpi/repos/ufo-filters && $(CMAKE_COMMAND) -E cmake_depends "Unix Makefiles" /ccpi/repos/ufo-filters /ccpi/repos/ufo-filters/src/kernels /ccpi/repos/ufo-filters /ccpi/repos/ufo-filters/src/kernels /ccpi/repos/ufo-filters/src/kernels/CMakeFiles/burst.dir/DependInfo.cmake --color=$(COLOR) +.PHONY : src/kernels/CMakeFiles/burst.dir/depend + diff --git a/src/kernels/CMakeFiles/burst.dir/cmake_clean.cmake b/src/kernels/CMakeFiles/burst.dir/cmake_clean.cmake new file mode 100644 index 0000000..fd937bc --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/cmake_clean.cmake @@ -0,0 +1,12 @@ +file(REMOVE_RECURSE + "CMakeFiles/burst" + "center_kernel.cl" + "lamino_kernel.cl" + "roll_kernel.cl" + "z_kernel.cl" +) + +# Per-language clean rules from dependency scanning. +foreach(lang ) + include(CMakeFiles/burst.dir/cmake_clean_${lang}.cmake OPTIONAL) +endforeach() diff --git a/src/kernels/CMakeFiles/burst.dir/depend.internal b/src/kernels/CMakeFiles/burst.dir/depend.internal new file mode 100644 index 0000000..f647855 --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/depend.internal @@ -0,0 +1,3 @@ +# CMAKE generated file: DO NOT EDIT! +# Generated by "Unix Makefiles" Generator, CMake Version 3.16 + diff --git a/src/kernels/CMakeFiles/burst.dir/depend.make b/src/kernels/CMakeFiles/burst.dir/depend.make new file mode 100644 index 0000000..f647855 --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/depend.make @@ -0,0 +1,3 @@ +# CMAKE generated file: DO NOT EDIT! +# Generated by "Unix Makefiles" Generator, CMake Version 3.16 + diff --git a/src/kernels/CMakeFiles/burst.dir/progress.make b/src/kernels/CMakeFiles/burst.dir/progress.make new file mode 100644 index 0000000..225de34 --- /dev/null +++ b/src/kernels/CMakeFiles/burst.dir/progress.make @@ -0,0 +1,2 @@ +CMAKE_PROGRESS_1 = + diff --git a/src/kernels/CMakeFiles/progress.marks b/src/kernels/CMakeFiles/progress.marks new file mode 100644 index 0000000..573541a --- /dev/null +++ b/src/kernels/CMakeFiles/progress.marks @@ -0,0 +1 @@ +0 diff --git a/src/kernels/CTestTestfile.cmake b/src/kernels/CTestTestfile.cmake new file mode 100644 index 0000000..4528b8b --- /dev/null +++ b/src/kernels/CTestTestfile.cmake @@ -0,0 +1,6 @@ +# CMake generated Testfile for +# Source directory: /ccpi/repos/ufo-filters/src/kernels +# Build directory: /ccpi/repos/ufo-filters/src/kernels +# +# This file includes the relevant testing commands required for +# testing this directory and lists subdirectories to be tested as well. diff --git a/src/kernels/Makefile b/src/kernels/Makefile new file mode 100644 index 0000000..a14a6a7 --- /dev/null +++ b/src/kernels/Makefile @@ -0,0 +1,212 @@ +# CMAKE generated file: DO NOT EDIT! +# Generated by "Unix Makefiles" Generator, CMake Version 3.16 + +# Default target executed when no arguments are given to make. +default_target: all + +.PHONY : default_target + +# Allow only one "make -f Makefile2" at a time, but pass parallelism. +.NOTPARALLEL: + + +#============================================================================= +# Special targets provided by cmake. + +# Disable implicit rules so canonical targets will work. +.SUFFIXES: + + +# Remove some rules from gmake that .SUFFIXES does not remove. +SUFFIXES = + +.SUFFIXES: .hpux_make_needs_suffix_list + + +# Suppress display of executed commands. +$(VERBOSE).SILENT: + + +# A target that is always out of date. +cmake_force: + +.PHONY : cmake_force + +#============================================================================= +# Set environment variables for the build. + +# The shell in which to execute make rules. +SHELL = /bin/sh + +# The CMake executable. +CMAKE_COMMAND = /usr/bin/cmake + +# The command to remove a file. +RM = /usr/bin/cmake -E remove -f + +# Escaping for special characters. +EQUALS = = + +# The top-level source directory on which CMake was run. +CMAKE_SOURCE_DIR = /ccpi/repos/ufo-filters + +# The top-level build directory on which CMake was run. +CMAKE_BINARY_DIR = /ccpi/repos/ufo-filters + +#============================================================================= +# Targets provided globally by CMake. + +# Special rule for the target install/strip +install/strip: preinstall + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Installing the project stripped..." + /usr/bin/cmake -DCMAKE_INSTALL_DO_STRIP=1 -P cmake_install.cmake +.PHONY : install/strip + +# Special rule for the target install/strip +install/strip/fast: preinstall/fast + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Installing the project stripped..." + /usr/bin/cmake -DCMAKE_INSTALL_DO_STRIP=1 -P cmake_install.cmake +.PHONY : install/strip/fast + +# Special rule for the target install/local +install/local: preinstall + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Installing only the local directory..." + /usr/bin/cmake -DCMAKE_INSTALL_LOCAL_ONLY=1 -P cmake_install.cmake +.PHONY : install/local + +# Special rule for the target install/local +install/local/fast: preinstall/fast + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Installing only the local directory..." + /usr/bin/cmake -DCMAKE_INSTALL_LOCAL_ONLY=1 -P cmake_install.cmake +.PHONY : install/local/fast + +# Special rule for the target install +install: preinstall + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Install the project..." + /usr/bin/cmake -P cmake_install.cmake +.PHONY : install + +# Special rule for the target install +install/fast: preinstall/fast + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Install the project..." + /usr/bin/cmake -P cmake_install.cmake +.PHONY : install/fast + +# Special rule for the target list_install_components +list_install_components: + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Available install components are: \"Unspecified\"" +.PHONY : list_install_components + +# Special rule for the target list_install_components +list_install_components/fast: list_install_components + +.PHONY : list_install_components/fast + +# Special rule for the target rebuild_cache +rebuild_cache: + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Running CMake to regenerate build system..." + /usr/bin/cmake -S$(CMAKE_SOURCE_DIR) -B$(CMAKE_BINARY_DIR) +.PHONY : rebuild_cache + +# Special rule for the target rebuild_cache +rebuild_cache/fast: rebuild_cache + +.PHONY : rebuild_cache/fast + +# Special rule for the target edit_cache +edit_cache: + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "No interactive CMake dialog available..." + /usr/bin/cmake -E echo No\ interactive\ CMake\ dialog\ available. +.PHONY : edit_cache + +# Special rule for the target edit_cache +edit_cache/fast: edit_cache + +.PHONY : edit_cache/fast + +# Special rule for the target test +test: + @$(CMAKE_COMMAND) -E cmake_echo_color --switch=$(COLOR) --cyan "Running tests..." + /usr/bin/ctest --force-new-ctest-process $(ARGS) +.PHONY : test + +# Special rule for the target test +test/fast: test + +.PHONY : test/fast + +# The main all target +all: cmake_check_build_system + cd /ccpi/repos/ufo-filters && $(CMAKE_COMMAND) -E cmake_progress_start /ccpi/repos/ufo-filters/CMakeFiles /ccpi/repos/ufo-filters/src/kernels/CMakeFiles/progress.marks + cd /ccpi/repos/ufo-filters && $(MAKE) -f CMakeFiles/Makefile2 src/kernels/all + $(CMAKE_COMMAND) -E cmake_progress_start /ccpi/repos/ufo-filters/CMakeFiles 0 +.PHONY : all + +# The main clean target +clean: + cd /ccpi/repos/ufo-filters && $(MAKE) -f CMakeFiles/Makefile2 src/kernels/clean +.PHONY : clean + +# The main clean target +clean/fast: clean + +.PHONY : clean/fast + +# Prepare targets for installation. +preinstall: all + cd /ccpi/repos/ufo-filters && $(MAKE) -f CMakeFiles/Makefile2 src/kernels/preinstall +.PHONY : preinstall + +# Prepare targets for installation. +preinstall/fast: + cd /ccpi/repos/ufo-filters && $(MAKE) -f CMakeFiles/Makefile2 src/kernels/preinstall +.PHONY : preinstall/fast + +# clear depends +depend: + cd /ccpi/repos/ufo-filters && $(CMAKE_COMMAND) -S$(CMAKE_SOURCE_DIR) -B$(CMAKE_BINARY_DIR) --check-build-system CMakeFiles/Makefile.cmake 1 +.PHONY : depend + +# Convenience name for target. +src/kernels/CMakeFiles/burst.dir/rule: + cd /ccpi/repos/ufo-filters && $(MAKE) -f CMakeFiles/Makefile2 src/kernels/CMakeFiles/burst.dir/rule +.PHONY : src/kernels/CMakeFiles/burst.dir/rule + +# Convenience name for target. +burst: src/kernels/CMakeFiles/burst.dir/rule + +.PHONY : burst + +# fast build rule for target. +burst/fast: + cd /ccpi/repos/ufo-filters && $(MAKE) -f src/kernels/CMakeFiles/burst.dir/build.make src/kernels/CMakeFiles/burst.dir/build +.PHONY : burst/fast + +# Help Target +help: + @echo "The following are some of the valid targets for this Makefile:" + @echo "... all (the default if no target is provided)" + @echo "... clean" + @echo "... depend" + @echo "... install/strip" + @echo "... install/local" + @echo "... install" + @echo "... list_install_components" + @echo "... rebuild_cache" + @echo "... edit_cache" + @echo "... test" + @echo "... burst" +.PHONY : help + + + +#============================================================================= +# Special targets to cleanup operation of make. + +# Special rule to run CMake to check the build system integrity. +# No rule that depends on this can have commands that come from listfiles +# because they might be regenerated. +cmake_check_build_system: + cd /ccpi/repos/ufo-filters && $(CMAKE_COMMAND) -S$(CMAKE_SOURCE_DIR) -B$(CMAKE_BINARY_DIR) --check-build-system CMakeFiles/Makefile.cmake 0 +.PHONY : cmake_check_build_system + diff --git a/src/kernels/backproject.cl b/src/kernels/backproject.cl index 8a81790..88467de 100644 --- a/src/kernels/backproject.cl +++ b/src/kernels/backproject.cl @@ -19,7 +19,7 @@ constant sampler_t volumeSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | - CLK_FILTER_LINEAR; + CLK_FILTER_NEAREST; kernel void backproject_nearest (global float *sinogram, @@ -86,10 +86,10 @@ backproject_tex (read_only image2d_t sinogram, #pragma unroll 4 #endif for(int proj = 0; proj < n_projections; proj++) { - float h = by * sin_lut[angle_offset + proj] + bx * cos_lut[angle_offset + proj] + axis_pos; + float h = -by * sin_lut[angle_offset + proj] + bx * cos_lut[angle_offset + proj] + axis_pos; sum += read_imagef (sinogram, volumeSampler, (float2)(h, proj + 0.5f)).x; } - slice[idy * get_global_size(0) + idx] = sum * M_PI_F / n_projections; + slice[idy * get_global_size(0) + idx] = sum; } diff --git a/src/kernels/center_kernel.cl b/src/kernels/center_kernel.cl new file mode 100644 index 0000000..ff9108b --- /dev/null +++ b/src/kernels/center_kernel.cl @@ -0,0 +1,478 @@ +#define rotate() pixel.x -= x_center.x; \ + pixel.y -= y_center; \ + pixel.x = pixel.x * cos_roll + pixel.y * sin_roll; \ + pixel.y = -pixel.x * sin_roll + pixel.y * cos_roll; \ + pixel.x += x_center.x; \ + pixel.y += y_center; + +/* + * Copyright (C) 2015-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 backproject_burst_1 ( +read_only image2d_t projection_0, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float sines, + const float cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + float x_center_current = mad((float) idz, x_center.y, x_center.x); + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines, mad(voxel.y, sines, x_center_current)); + pixel.y = mad(tmp_x, sines, mad(tmp_y, cosines, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_2 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float2 sines, + const float2 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + float x_center_current = mad((float) idz, x_center.y, x_center.x); + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center_current)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center_current)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_4 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float4 sines, + const float4 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + float x_center_current = mad((float) idz, x_center.y, x_center.x); + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center_current)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center_current)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center_current)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center_current)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_8 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float8 sines, + const float8 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + float x_center_current = mad((float) idz, x_center.y, x_center.x); + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center_current)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center_current)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center_current)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center_current)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center_current)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center_current)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center_current)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center_current)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_16 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, +read_only image2d_t projection_8, +read_only image2d_t projection_9, +read_only image2d_t projection_10, +read_only image2d_t projection_11, +read_only image2d_t projection_12, +read_only image2d_t projection_13, +read_only image2d_t projection_14, +read_only image2d_t projection_15, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float16 sines, + const float16 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + float x_center_current = mad((float) idz, x_center.y, x_center.x); + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center_current)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center_current)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center_current)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center_current)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center_current)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center_current)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center_current)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center_current)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s8, mad(voxel.y, sines.s8, x_center_current)); + pixel.y = mad(tmp_x, sines.s8, mad(tmp_y, cosines.s8, tmp)); + rotate (); + result += read_imagef (projection_8, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s9, mad(voxel.y, sines.s9, x_center_current)); + pixel.y = mad(tmp_x, sines.s9, mad(tmp_y, cosines.s9, tmp)); + rotate (); + result += read_imagef (projection_9, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sa, mad(voxel.y, sines.sa, x_center_current)); + pixel.y = mad(tmp_x, sines.sa, mad(tmp_y, cosines.sa, tmp)); + rotate (); + result += read_imagef (projection_10, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sb, mad(voxel.y, sines.sb, x_center_current)); + pixel.y = mad(tmp_x, sines.sb, mad(tmp_y, cosines.sb, tmp)); + rotate (); + result += read_imagef (projection_11, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sc, mad(voxel.y, sines.sc, x_center_current)); + pixel.y = mad(tmp_x, sines.sc, mad(tmp_y, cosines.sc, tmp)); + rotate (); + result += read_imagef (projection_12, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sd, mad(voxel.y, sines.sd, x_center_current)); + pixel.y = mad(tmp_x, sines.sd, mad(tmp_y, cosines.sd, tmp)); + rotate (); + result += read_imagef (projection_13, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.se, mad(voxel.y, sines.se, x_center_current)); + pixel.y = mad(tmp_x, sines.se, mad(tmp_y, cosines.se, tmp)); + rotate (); + result += read_imagef (projection_14, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sf, mad(voxel.y, sines.sf, x_center_current)); + pixel.y = mad(tmp_x, sines.sf, mad(tmp_y, cosines.sf, tmp)); + rotate (); + result += read_imagef (projection_15, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} + diff --git a/src/kernels/cmake_install.cmake b/src/kernels/cmake_install.cmake new file mode 100644 index 0000000..c3a2e2f --- /dev/null +++ b/src/kernels/cmake_install.cmake @@ -0,0 +1,684 @@ +# Install script for directory: /ccpi/repos/ufo-filters/src/kernels + +# Set the install prefix +if(NOT DEFINED CMAKE_INSTALL_PREFIX) + set(CMAKE_INSTALL_PREFIX "/usr/local") +endif() +string(REGEX REPLACE "/$" "" CMAKE_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}") + +# Set the install configuration name. +if(NOT DEFINED CMAKE_INSTALL_CONFIG_NAME) + if(BUILD_TYPE) + string(REGEX REPLACE "^[^A-Za-z0-9_]+" "" + CMAKE_INSTALL_CONFIG_NAME "${BUILD_TYPE}") + else() + set(CMAKE_INSTALL_CONFIG_NAME "") + endif() + message(STATUS "Install configuration: \"${CMAKE_INSTALL_CONFIG_NAME}\"") +endif() + +# Set the component getting installed. +if(NOT CMAKE_INSTALL_COMPONENT) + if(COMPONENT) + message(STATUS "Install component: \"${COMPONENT}\"") + set(CMAKE_INSTALL_COMPONENT "${COMPONENT}") + else() + set(CMAKE_INSTALL_COMPONENT) + endif() +endif() + +# Install shared libraries without execute permission? +if(NOT DEFINED CMAKE_INSTALL_SO_NO_EXE) + set(CMAKE_INSTALL_SO_NO_EXE "1") +endif() + +# Is this installation the result of a crosscompile? +if(NOT DEFINED CMAKE_CROSSCOMPILING) + set(CMAKE_CROSSCOMPILING "FALSE") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/z_kernel.cl;/usr/share/ufo/center_kernel.cl;/usr/share/ufo/lamino_kernel.cl;/usr/share/ufo/roll_kernel.cl;/usr/share/ufo/general_bp_definitions.in;/usr/share/ufo/general_bp_header_scalar.in;/usr/share/ufo/general_bp_header_vector.in;/usr/share/ufo/general_bp_body.in") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES + "/ccpi/repos/ufo-filters/src/kernels/z_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/center_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/lamino_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/roll_kernel.cl" + "/ccpi/repos/ufo-filters/src/kernels/templates/general_bp_definitions.in" + "/ccpi/repos/ufo-filters/src/kernels/templates/general_bp_header_scalar.in" + "/ccpi/repos/ufo-filters/src/kernels/templates/general_bp_header_vector.in" + "/ccpi/repos/ufo-filters/src/kernels/templates/general_bp_body.in" + ) +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/arithmetics.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/arithmetics.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/backproject.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/backproject.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/bin.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/bin.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/binarize.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/binarize.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/center_kernel.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/center_kernel.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/clip.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/clip.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/complex.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/complex.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/conebeam.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/conebeam.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/correlate.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/correlate.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/cumsum.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/cumsum.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/cut-sinogram.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/cut-sinogram.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/cut.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/cut.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/denoise.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/denoise.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/dfi.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/dfi.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/edge.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/edge.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/estimate-noise.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/estimate-noise.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/ffc.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/ffc.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/fft.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/fft.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/fftmult.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/fftmult.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/filter.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/filter.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/flip.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/flip.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/forwardproject.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/forwardproject.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/gaussian.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/gaussian.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/gradient.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/gradient.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/histthreshold.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/histthreshold.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/interpolator.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/interpolator.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/lamino_kernel.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/lamino_kernel.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/mask.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/mask.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/median.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/median.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/metaballs.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/metaballs.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/morphology.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/morphology.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/nlm.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/nlm.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/opencl-reduce.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/opencl-reduce.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/opencl.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/opencl.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/ordfilt.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/ordfilt.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/pad.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/pad.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/phase-retrieval.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/phase-retrieval.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/piv.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/piv.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/polar.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/polar.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/reductor.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/reductor.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/rescale.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/rescale.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/rm-outliers.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/rm-outliers.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/roll_kernel.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/roll_kernel.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/rotate.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/rotate.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/segment.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/segment.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/split.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/split.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/stacked-backproject.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/stacked-backproject.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/stacked-forwardproject.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/stacked-forwardproject.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/swap-quadrants.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/swap-quadrants.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/transpose.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/transpose.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/z_kernel.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/z_kernel.cl") +endif() + +if("x${CMAKE_INSTALL_COMPONENT}x" STREQUAL "xUnspecifiedx" OR NOT CMAKE_INSTALL_COMPONENT) + list(APPEND CMAKE_ABSOLUTE_DESTINATION_FILES + "/usr/share/ufo/zeropad.cl") + if(CMAKE_WARN_ON_ABSOLUTE_INSTALL_DESTINATION) + message(WARNING "ABSOLUTE path INSTALL DESTINATION : ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() + if(CMAKE_ERROR_ON_ABSOLUTE_INSTALL_DESTINATION) + message(FATAL_ERROR "ABSOLUTE path INSTALL DESTINATION forbidden (by caller): ${CMAKE_ABSOLUTE_DESTINATION_FILES}") + endif() +file(INSTALL DESTINATION "/usr/share/ufo" TYPE FILE FILES "/ccpi/repos/ufo-filters/src/kernels/zeropad.cl") +endif() + diff --git a/src/kernels/forwardproject.cl b/src/kernels/forwardproject.cl index 18a8019..89faabd 100644 --- a/src/kernels/forwardproject.cl +++ b/src/kernels/forwardproject.cl @@ -19,7 +19,7 @@ constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | - CLK_FILTER_LINEAR; + CLK_FILTER_NEAREST; kernel void forwardproject(read_only image2d_t slice, @@ -37,10 +37,10 @@ forwardproject(read_only image2d_t slice, /* positive/negative distance from detector center */ const float d = idx - axis_pos + 0.5f; /* length of the cut through the circle */ - const float l = sqrt(4.0f*r*r - 4.0f*d*d); + const float l = sqrt(8.0f*r*r - 4.0f*d*d); /* vector in detector direction */ - float2 D = (float2) (cos(angle), sin(angle)); + float2 D = (float2) (cos(angle), -sin(angle)); D = normalize(D); /* vector perpendicular to the detector */ diff --git a/src/kernels/lamino_kernel.cl b/src/kernels/lamino_kernel.cl new file mode 100644 index 0000000..e573910 --- /dev/null +++ b/src/kernels/lamino_kernel.cl @@ -0,0 +1,478 @@ +#define rotate() pixel.x -= x_center.x; \ + pixel.y -= y_center; \ + pixel.x = pixel.x * cos_roll + pixel.y * sin_roll; \ + pixel.y = -pixel.x * sin_roll + pixel.y * cos_roll; \ + pixel.x += x_center.x; \ + pixel.y += y_center; + +/* + * Copyright (C) 2015-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 backproject_burst_1 ( +read_only image2d_t projection_0, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float sines, + const float cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_lamino = sincos (mad((float) idz, lamino_region.y, lamino_region.x), &cos_lamino); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines, mad(voxel.y, sines, x_center.x)); + pixel.y = mad(tmp_x, sines, mad(tmp_y, cosines, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_2 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float2 sines, + const float2 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_lamino = sincos (mad((float) idz, lamino_region.y, lamino_region.x), &cos_lamino); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_4 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float4 sines, + const float4 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_lamino = sincos (mad((float) idz, lamino_region.y, lamino_region.x), &cos_lamino); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_8 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float8 sines, + const float8 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_lamino = sincos (mad((float) idz, lamino_region.y, lamino_region.x), &cos_lamino); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_16 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, +read_only image2d_t projection_8, +read_only image2d_t projection_9, +read_only image2d_t projection_10, +read_only image2d_t projection_11, +read_only image2d_t projection_12, +read_only image2d_t projection_13, +read_only image2d_t projection_14, +read_only image2d_t projection_15, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float16 sines, + const float16 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_lamino = sincos (mad((float) idz, lamino_region.y, lamino_region.x), &cos_lamino); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s8, mad(voxel.y, sines.s8, x_center.x)); + pixel.y = mad(tmp_x, sines.s8, mad(tmp_y, cosines.s8, tmp)); + rotate (); + result += read_imagef (projection_8, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s9, mad(voxel.y, sines.s9, x_center.x)); + pixel.y = mad(tmp_x, sines.s9, mad(tmp_y, cosines.s9, tmp)); + rotate (); + result += read_imagef (projection_9, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sa, mad(voxel.y, sines.sa, x_center.x)); + pixel.y = mad(tmp_x, sines.sa, mad(tmp_y, cosines.sa, tmp)); + rotate (); + result += read_imagef (projection_10, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sb, mad(voxel.y, sines.sb, x_center.x)); + pixel.y = mad(tmp_x, sines.sb, mad(tmp_y, cosines.sb, tmp)); + rotate (); + result += read_imagef (projection_11, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sc, mad(voxel.y, sines.sc, x_center.x)); + pixel.y = mad(tmp_x, sines.sc, mad(tmp_y, cosines.sc, tmp)); + rotate (); + result += read_imagef (projection_12, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sd, mad(voxel.y, sines.sd, x_center.x)); + pixel.y = mad(tmp_x, sines.sd, mad(tmp_y, cosines.sd, tmp)); + rotate (); + result += read_imagef (projection_13, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.se, mad(voxel.y, sines.se, x_center.x)); + pixel.y = mad(tmp_x, sines.se, mad(tmp_y, cosines.se, tmp)); + rotate (); + result += read_imagef (projection_14, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sf, mad(voxel.y, sines.sf, x_center.x)); + pixel.y = mad(tmp_x, sines.sf, mad(tmp_y, cosines.sf, tmp)); + rotate (); + result += read_imagef (projection_15, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} + diff --git a/src/kernels/roll_kernel.cl b/src/kernels/roll_kernel.cl new file mode 100644 index 0000000..129e14f --- /dev/null +++ b/src/kernels/roll_kernel.cl @@ -0,0 +1,478 @@ +#define rotate() pixel.x -= x_center.x; \ + pixel.y -= y_center; \ + pixel.x = pixel.x * cos_roll + pixel.y * sin_roll; \ + pixel.y = -pixel.x * sin_roll + pixel.y * cos_roll; \ + pixel.x += x_center.x; \ + pixel.y += y_center; + +/* + * Copyright (C) 2015-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 backproject_burst_1 ( +read_only image2d_t projection_0, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float sines, + const float cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_roll = sincos (mad((float) idz, -roll_region.y, -roll_region.x), &cos_roll); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines, mad(voxel.y, sines, x_center.x)); + pixel.y = mad(tmp_x, sines, mad(tmp_y, cosines, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_2 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float2 sines, + const float2 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_roll = sincos (mad((float) idz, -roll_region.y, -roll_region.x), &cos_roll); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_4 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float4 sines, + const float4 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_roll = sincos (mad((float) idz, -roll_region.y, -roll_region.x), &cos_roll); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_8 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float8 sines, + const float8 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_roll = sincos (mad((float) idz, -roll_region.y, -roll_region.x), &cos_roll); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_16 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, +read_only image2d_t projection_8, +read_only image2d_t projection_9, +read_only image2d_t projection_10, +read_only image2d_t projection_11, +read_only image2d_t projection_12, +read_only image2d_t projection_13, +read_only image2d_t projection_14, +read_only image2d_t projection_15, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float16 sines, + const float16 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + sin_roll = sincos (mad((float) idz, -roll_region.y, -roll_region.x), &cos_roll); + tmp = mad(z_region.x, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s8, mad(voxel.y, sines.s8, x_center.x)); + pixel.y = mad(tmp_x, sines.s8, mad(tmp_y, cosines.s8, tmp)); + rotate (); + result += read_imagef (projection_8, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s9, mad(voxel.y, sines.s9, x_center.x)); + pixel.y = mad(tmp_x, sines.s9, mad(tmp_y, cosines.s9, tmp)); + rotate (); + result += read_imagef (projection_9, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sa, mad(voxel.y, sines.sa, x_center.x)); + pixel.y = mad(tmp_x, sines.sa, mad(tmp_y, cosines.sa, tmp)); + rotate (); + result += read_imagef (projection_10, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sb, mad(voxel.y, sines.sb, x_center.x)); + pixel.y = mad(tmp_x, sines.sb, mad(tmp_y, cosines.sb, tmp)); + rotate (); + result += read_imagef (projection_11, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sc, mad(voxel.y, sines.sc, x_center.x)); + pixel.y = mad(tmp_x, sines.sc, mad(tmp_y, cosines.sc, tmp)); + rotate (); + result += read_imagef (projection_12, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sd, mad(voxel.y, sines.sd, x_center.x)); + pixel.y = mad(tmp_x, sines.sd, mad(tmp_y, cosines.sd, tmp)); + rotate (); + result += read_imagef (projection_13, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.se, mad(voxel.y, sines.se, x_center.x)); + pixel.y = mad(tmp_x, sines.se, mad(tmp_y, cosines.se, tmp)); + rotate (); + result += read_imagef (projection_14, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sf, mad(voxel.y, sines.sf, x_center.x)); + pixel.y = mad(tmp_x, sines.sf, mad(tmp_y, cosines.sf, tmp)); + rotate (); + result += read_imagef (projection_15, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} + diff --git a/src/kernels/stacked-backproject.cl b/src/kernels/stacked-backproject.cl index d6fff5f..c9938cb 100644 --- a/src/kernels/stacked-backproject.cl +++ b/src/kernels/stacked-backproject.cl @@ -23,7 +23,7 @@ constant sampler_t volumeSampler_single = CLK_NORMALIZED_COORDS_FALSE | constant sampler_t volumeSampler_half = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | - CLK_FILTER_NEAREST; + CLK_FILTER_LINEAR; constant sampler_t volumeSampler_int8 = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | @@ -47,6 +47,34 @@ interleave_single ( global float *sinogram, write_imagef(interleaved_sinograms, (int4)(idx, idy, idz, 0),(float4)(x,y,0.0f,0.0f)); } +/*kernel void +texture_single (read_only image2d_array_t sinogram, + global float2 *reconstructed_buffer, + constant float *sin_lut, + constant float *cos_lut, + const unsigned int x_offset, + const unsigned int y_offset, + const unsigned int angle_offset, + const unsigned int n_projections, + const float axis_pos, + unsigned long size) +{ + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + const float bx = idx - axis_pos + x_offset + 0.5f; + const float by = idy - axis_pos + y_offset + 0.5f; + float2 sum = {0.0f, 0.0f}; + + for(int proj = 0; proj < n_projections; proj++) { + float h = -by * sin_lut[angle_offset + proj] + bx * cos_lut[angle_offset + proj] + axis_pos; + sum += read_imagef (sinogram, volumeSampler_single, (float4)(h, proj + 0.5f,idz,0.0f)).xy; + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; +}*/ + + kernel void texture_single ( read_only image2d_array_t sinogram, @@ -73,12 +101,12 @@ texture_single ( int global_sizex = get_global_size(0); int global_sizey = get_global_size(1); - /* Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant */ + // Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant int square = local_idy%4; int quadrant = local_idx/4; int pixel = local_idx%4; - /* Computing projection and pixel offsets */ + // Computing projection and pixel offsets int projection_index = local_idy/4; int2 remapped_index_local = {(4*square + 2*(quadrant%2) + (pixel%2)), @@ -94,6 +122,27 @@ texture_single ( __local float2 shared_mem[64][4]; __local float2 reconstructed_cache[16][16]; +/*#ifdef DEVICE_TESLA_K20XM +#pragma unroll 4 +#endif +#ifdef DEVICE_TESLA_P100_PCIE_16GB +#pragma unroll 2 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 8 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN +#pragma unroll 14 +#endif +#ifdef DEVICE_GEFORCE_GTX_1080_TI +#pragma unroll 10 +#endif +#ifdef DEVICE_QUADRO_M6000 +#pragma unroll 2 +#endif +#ifdef DEVICE_GFX1010 +#pragma unroll 4 +#endif*/ for(int proj = projection_index; proj < n_projections; proj+=4) { float sine_value = sin_lut[angle_offset + proj]; @@ -106,7 +155,7 @@ texture_single ( int2 remapped_index = {(local_idx%4), (4*local_idy + (local_idx/4))}; for(int q=0; q<4;q+=1){ - /* Moving partial sums to shared memory */ + // Moving partial sums to shared memory shared_mem[(local_sizex*remapped_index_local.y + remapped_index_local.x)][projection_index] = sum[q]; barrier(CLK_LOCAL_MEM_FENCE); // syncthreads @@ -124,7 +173,7 @@ texture_single ( barrier(CLK_LOCAL_MEM_FENCE); // syncthreads } - reconstructed_buffer[global_idx + global_idy*size + idz*size*size] = reconstructed_cache[local_idy][local_idx] * M_PI_F / n_projections; + reconstructed_buffer[global_idx + global_idy*size + idz*size*size] = reconstructed_cache[local_idy][local_idx]; } kernel void @@ -163,6 +212,54 @@ interleave_half (global float *sinogram, write_imagef(interleaved_sinograms, (int4)(idx, idy, idz, 0),(float4)(b)); } +/*kernel void +texture_half (read_only image2d_array_t sinogram, + global float4 *reconstructed_buffer, + constant float *sin_lut, + constant float *cos_lut, + const unsigned int x_offset, + const unsigned int y_offset, + const unsigned int angle_offset, + const unsigned int n_projections, + const float axis_pos, + unsigned long size) +{ + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + const float bx = idx - axis_pos + x_offset + 0.5f; + const float by = idy - axis_pos + y_offset + 0.5f; + float4 sum = {0.0f, 0.0f, 0.0f, 0.0f}; + +#ifdef DEVICE_TESLA_K20XM +#pragma unroll 4 +#endif +#ifdef DEVICE_TESLA_P100_PCIE_16GB +#pragma unroll 2 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 8 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN +#pragma unroll 14 +#endif +#ifdef DEVICE_GEFORCE_GTX_1080_TI +#pragma unroll 10 +#endif +#ifdef DEVICE_QUADRO_M6000 +#pragma unroll 2 +#endif +#ifdef DEVICE_GFX1010 +#pragma unroll 4 +#endif + for(int proj = 0; proj < n_projections; proj++) { + float h = -by * sin_lut[angle_offset + proj] + bx * cos_lut[angle_offset + proj] + axis_pos; + sum += read_imagef (sinogram, volumeSampler_half, (float4)(h, proj + 0.5f,idz,0.0f)); + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; +}*/ + kernel void texture_half ( read_only image2d_array_t sinogram, @@ -189,12 +286,12 @@ texture_half ( int global_sizex = get_global_size(0); int global_sizey = get_global_size(1); - /* Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant */ + // Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant int square = local_idy%4; int quadrant = local_idx/4; int pixel = local_idx%4; - /* Computing projection and pixel offsets */ + // Computing projection and pixel offsets int projection_index = local_idy/4; int2 remapped_index_local = {(4*square + 2*(quadrant%2) + (pixel%2)),(2* (quadrant/2) + (pixel/2))}; int2 remapped_index_global = {(get_group_id(0)*get_local_size(0)+remapped_index_local.x), @@ -206,6 +303,27 @@ texture_half ( __local float4 shared_mem[64][4]; __local float4 reconstructed_cache[16][16]; +#ifdef DEVICE_TESLA_K20XM +#pragma unroll 4 +#endif +#ifdef DEVICE_TESLA_P100_PCIE_16GB +#pragma unroll 2 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 8 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN +#pragma unroll 14 +#endif +#ifdef DEVICE_GEFORCE_GTX_1080_TI +#pragma unroll 10 +#endif +#ifdef DEVICE_QUADRO_M6000 +#pragma unroll 2 +#endif +#ifdef DEVICE_GFX1010 +#pragma unroll 4 +#endif for(int proj = projection_index; proj < n_projections; proj+=4) { float sine_value = sin_lut[angle_offset + proj]; @@ -218,7 +336,7 @@ texture_half ( int2 remapped_index = {(local_idx%4), (4*local_idy + (local_idx/4))}; for(int q=0; q<4;q+=1){ - /* Moving partial sums to shared memory */ + // Moving partial sums to shared memory shared_mem[(local_sizex*remapped_index_local.y + remapped_index_local.x)][projection_index] = sum[q]; barrier(CLK_LOCAL_MEM_FENCE); // syncthreads @@ -235,7 +353,7 @@ texture_half ( } barrier(CLK_LOCAL_MEM_FENCE); // syncthreads } - reconstructed_buffer[global_idx + global_idy*size + idz*size*size] = reconstructed_cache[local_idy][local_idx] * M_PI_F / n_projections; + reconstructed_buffer[global_idx + global_idy*size + idz*size*size] = reconstructed_cache[local_idy][local_idx]; } kernel void @@ -258,6 +376,11 @@ uninterleave_half (global float4 *reconstructed_buffer, output[idx + idy*sizex + (output_offset+3)*sizex*sizey] = b.w; } +union converter { + uint2 storage; + uchar8 a; +}; + kernel void interleave_uint (global float *sinogram, write_only image2d_array_t interleaved_sinograms, @@ -270,22 +393,85 @@ interleave_uint (global float *sinogram, const int sizex = get_global_size(0); const int sizey = get_global_size(1); - int sinogram_offset = idz*4; + int sinogram_offset = idz*8; const float scale = 255.0f / (max - min); - uint4 b = {(sinogram[idx + idy * sizex + (sinogram_offset) * sizex * sizey] - min)*scale, - (sinogram[idx + idy * sizex + (sinogram_offset+1) * sizex * sizey] - min)*scale, - (sinogram[idx + idy * sizex + (sinogram_offset+2) * sizex * sizey] - min)*scale, - (sinogram[idx + idy * sizex + (sinogram_offset+3) * sizex * sizey] - min)*scale}; - - write_imageui(interleaved_sinograms, (int4)(idx, idy, idz, 0),(uint4)(b)); + union converter il; + il.a.s0 = (sinogram[idx + idy * sizex + (sinogram_offset) * sizex * sizey] - min)*scale; + il.a.s1 = (sinogram[idx + idy * sizex + (sinogram_offset+1) * sizex * sizey] - min)*scale; + il.a.s2 = (sinogram[idx + idy * sizex + (sinogram_offset+2) * sizex * sizey] - min)*scale; + il.a.s3 = (sinogram[idx + idy * sizex + (sinogram_offset+3) * sizex * sizey] - min)*scale; + il.a.s4 = (sinogram[idx + idy * sizex + (sinogram_offset+4) * sizex * sizey] - min)*scale; + il.a.s5 = (sinogram[idx + idy * sizex + (sinogram_offset+5) * sizex * sizey] - min)*scale; + il.a.s6 = (sinogram[idx + idy * sizex + (sinogram_offset+6) * sizex * sizey] - min)*scale; + il.a.s7 = (sinogram[idx + idy * sizex + (sinogram_offset+7) * sizex * sizey] - min)*scale; + + write_imageui(interleaved_sinograms, (int4)(idx, idy, idz, 0),(uint4)((uint)il.storage.x,(uint)il.storage.y,0,0)); } +/*kernel void +texture_uint (read_only image2d_array_t sinogram, + global uint8 *reconstructed_buffer, + constant float *sin_lut, + constant float *cos_lut, + const unsigned int x_offset, + const unsigned int y_offset, + const unsigned int angle_offset, + const unsigned int n_projections, + const float axis_pos, + unsigned long size) +{ + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + const float bx = idx - axis_pos + x_offset + 0.5f; + const float by = idy - axis_pos + y_offset + 0.5f; + uint8 sum = {0,0,0,0,0,0,0,0}; + +#ifdef DEVICE_TESLA_K20XM +#pragma unroll 4 +#endif +#ifdef DEVICE_TESLA_P100_PCIE_16GB +#pragma unroll 2 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 8 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN +#pragma unroll 14 +#endif +#ifdef DEVICE_GEFORCE_GTX_1080_TI +#pragma unroll 10 +#endif +#ifdef DEVICE_QUADRO_M6000 +#pragma unroll 2 +#endif +#ifdef DEVICE_GFX1010 +#pragma unroll 4 +#endif + union converter tex; + for(int proj = 0; proj < n_projections; proj++) { + float h = -by * sin_lut[angle_offset + proj] + bx * cos_lut[angle_offset + proj] + axis_pos; + tex.storage = read_imageui (sinogram, volumeSampler_int8, (float4)(h, proj + 0.5f,idz,0.0f)).xy; + + sum.s0 += (uint)tex.a.s0; + sum.s1 += (uint)tex.a.s1; + sum.s2 += (uint)tex.a.s2; + sum.s3 += (uint)tex.a.s3; + sum.s4 += (uint)tex.a.s4; + sum.s5 += (uint)tex.a.s5; + sum.s6 += (uint)tex.a.s6; + sum.s7 += (uint)tex.a.s7; + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; +}*/ + kernel void texture_uint ( read_only image2d_array_t sinogram, - global uint4 *reconstructed_buffer, + global uint8 *reconstructed_buffer, constant float *sin_lut, constant float *cos_lut, const unsigned int x_offset, @@ -308,12 +494,12 @@ texture_uint ( int global_sizex = get_global_size(0); int global_sizey = get_global_size(1); - /* Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant */ + // Computing sequential numbers of 4x4 square, quadrant, and pixel within quadrant int square = local_idy%4; int quadrant = local_idx/4; int pixel = local_idx%4; - /* Computing projection and pixel offsets */ + // Computing projection and pixel offsets int projection_index = local_idy/4; int2 remapped_index_local = {(4*square + 2*(quadrant%2) + (pixel%2)),(2* (quadrant/2) + (pixel/2))}; int2 remapped_index_global = {(get_group_id(0)*get_local_size(0)+remapped_index_local.x), @@ -321,23 +507,56 @@ texture_uint ( float2 pixel_coord = {(remapped_index_global.x-axis_pos+x_offset+0.5f), (remapped_index_global.y-axis_pos+y_offset+0.5f)}; //bx and by - uint4 sum[4] = {0,0,0,0}; - __local uint4 shared_mem[64][4]; - __local uint4 reconstructed_cache[16][16]; + uint8 sum[4] = {0,0,0,0}; + __local uint8 shared_mem[64][4]; + __local uint8 reconstructed_cache[16][16]; + + union converter tex; + +#ifdef DEVICE_TESLA_K20XM +#pragma unroll 4 +#endif +#ifdef DEVICE_TESLA_P100_PCIE_16GB +#pragma unroll 2 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN_BLACK +#pragma unroll 8 +#endif +#ifdef DEVICE_GEFORCE_GTX_TITAN +#pragma unroll 14 +#endif +#ifdef DEVICE_GEFORCE_GTX_1080_TI +#pragma unroll 10 +#endif +#ifdef DEVICE_QUADRO_M6000 +#pragma unroll 2 +#endif +#ifdef DEVICE_GFX1010 +#pragma unroll 4 +#endif for(int proj = projection_index; proj < n_projections; proj+=4) { float sine_value = sin_lut[angle_offset + proj]; float h = pixel_coord.x * cos_lut[angle_offset + proj] - pixel_coord.y * sin_lut[angle_offset + proj] + axis_pos; for(int q=0; q<4; q+=1){ - sum[q] += read_imageui(sinogram, volumeSampler_int8, (float4)(h-4*q*sine_value, proj + 0.5f,idz, 0.0)); + tex.storage = read_imageui(sinogram, volumeSampler_int8, (float4)(h-4*q*sine_value, proj + 0.5f,idz, 0.0)).xy; + + sum[q].s0 += (uint)tex.a.s0; + sum[q].s1 += (uint)tex.a.s1; + sum[q].s2 += (uint)tex.a.s2; + sum[q].s3 += (uint)tex.a.s3; + sum[q].s4 += (uint)tex.a.s4; + sum[q].s5 += (uint)tex.a.s5; + sum[q].s6 += (uint)tex.a.s6; + sum[q].s7 += (uint)tex.a.s7; + } } - int2 remapped_index = {(local_idx%4), (4*local_idy + (local_idx/4))}; for(int q=0; q<4;q+=1){ - /* Moving partial sums to shared memory */ + // Moving partial sums to shared memory shared_mem[(local_sizex*remapped_index_local.y + remapped_index_local.x)][projection_index] = sum[q]; barrier(CLK_LOCAL_MEM_FENCE); // syncthreads @@ -359,7 +578,7 @@ texture_uint ( } kernel void -uninterleave_uint (global uint4 *reconstructed_buffer, +uninterleave_uint (global uint8 *reconstructed_buffer, global float *output, const float min, const float max, @@ -371,11 +590,15 @@ uninterleave_uint (global uint4 *reconstructed_buffer, const int sizex = get_global_size(0); const int sizey = get_global_size(1); - int output_offset = idz*4; + int output_offset = idz*8; float scale = (max-min)/255.0f; - output[idx + idy*sizex + (output_offset)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].x)*scale+min)* M_PI_F / n_projections; - output[idx + idy*sizex + (output_offset+1)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].y)*scale+min)* M_PI_F / n_projections; - output[idx + idy*sizex + (output_offset+2)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].z)*scale+min)* M_PI_F / n_projections; - output[idx + idy*sizex + (output_offset+3)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].w)*scale+min)* M_PI_F / n_projections; -}
\ No newline at end of file + output[idx + idy*sizex + (output_offset)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s0)*scale+min) ; + output[idx + idy*sizex + (output_offset+1)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s1)*scale+min) ; + output[idx + idy*sizex + (output_offset+2)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s2)*scale+min) ; + output[idx + idy*sizex + (output_offset+3)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s3)*scale+min) ; + output[idx + idy*sizex + (output_offset+4)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s4)*scale+min) ; + output[idx + idy*sizex + (output_offset+5)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s5)*scale+min) ; + output[idx + idy*sizex + (output_offset+6)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s6)*scale+min) ; + output[idx + idy*sizex + (output_offset+7)*sizex*sizey] = ((reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s7)*scale+min) ; +} diff --git a/src/kernels/stacked-forwardproject.cl b/src/kernels/stacked-forwardproject.cl new file mode 100644 index 0000000..c37eb53 --- /dev/null +++ b/src/kernels/stacked-forwardproject.cl @@ -0,0 +1,274 @@ +/* + * 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/>. + */ + + constant sampler_t volumeSampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP | + CLK_FILTER_NEAREST; + + kernel void + interleave_single ( global float *slices, + write_only image2d_array_t interleaved_slices) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + + int slice_offset = idz*2; + + float x = slices[idx + idy * sizex + (slice_offset) * sizex * sizey]; + float y = slices[idx + idy * sizex + (slice_offset+1) * sizex * sizey]; + + write_imagef(interleaved_slices, (int4)(idx, idy, idz, 0),(float4)(x,y,0.0f,0.0f)); + } + + kernel void + texture_single ( + read_only image2d_array_t slices, + global float2 *reconstructed_buffer, + float axis_pos, + float angle_step, + unsigned long size){ + + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + float angle = idy * angle_step; + float r = fmin (axis_pos, size - axis_pos); + float d = idx - axis_pos + 0.5f; + float l = sqrt(8.0f*r*r - 4.0f*d*d); + + float2 D = (float2) (cos(angle), -sin(angle)); + + D = normalize(D); + + float2 N = (float2) (D.y, -D.x); + + float2 sample = d * D - l/2.0f * N + ((float2) (axis_pos, axis_pos)); + + float2 sum = {0.0f,0.0f}; + + for (int i = 0; i < l; i++) { + sum += read_imagef(slices, volumeSampler, (float4)((float2)sample,idz,0.0f)).xy; + sample += N; + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; + } + + + kernel void + uninterleave_single (global float2 *reconstructed_buffer, + global float *output) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + int output_offset = idz*2; + + output[idx + idy*sizex + (output_offset)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].x; + output[idx + idy*sizex + (output_offset+1)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].y; + } + + kernel void + interleave_half ( global float *slices, + write_only image2d_array_t interleaved_slices) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + + int slice_offset = idz*4; + + float4 b = {slices[idx + idy * sizex + (slice_offset) * sizex * sizey], + slices[idx + idy * sizex + (slice_offset+1) * sizex * sizey], + slices[idx + idy * sizex + (slice_offset+2) * sizex * sizey], + slices[idx + idy * sizex + (slice_offset+3) * sizex * sizey]}; + + write_imagef(interleaved_slices, (int4)(idx, idy, idz, 0),(float4)(b)); + } + +kernel void +texture_half ( + read_only image2d_array_t slices, + global float4 *reconstructed_buffer, + float axis_pos, + float angle_step, + unsigned long size){ + + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + float angle = idy * angle_step; + float r = fmin (axis_pos, size - axis_pos); + float d = idx - axis_pos + 0.5f; + float l = sqrt(8.0f*r*r - 4.0f*d*d); + + float2 D = (float2) (cos(angle), -sin(angle)); + + D = normalize(D); + + float2 N = (float2) (D.y, -D.x); + + float2 sample = d * D - l/2.0f * N + ((float2) (axis_pos, axis_pos)); + + float4 sum = {0.0f,0.0f,0.0f,0.0f}; + + for (int i = 0; i < l; i++) { + sum += read_imagef(slices, volumeSampler, (float4)((float2)sample,idz,0.0f)); + sample += N; + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; +} + + kernel void + uninterleave_half (global float4 *reconstructed_buffer, + global float *output) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + + int output_offset = idz*4; + + output[idx + idy*sizex + (output_offset)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].x; + output[idx + idy*sizex + (output_offset+1)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].y; + output[idx + idy*sizex + (output_offset+2)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].z; + output[idx + idy*sizex + (output_offset+3)*sizex*sizey] = reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].w; + } + + union converter { + uint2 storage; + uchar8 a; + }; + + kernel void + interleave_uint ( global float *slices, + write_only image2d_array_t interleaved_slices, + const float min, + const float max) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + + int slice_offset = idz*8; + + const float scale = 255.0f / (max - min); + + union converter il; + il.a.s0 = (slices[idx + idy * sizex + (slice_offset) * sizex * sizey] - min)*scale; + il.a.s1 = (slices[idx + idy * sizex + (slice_offset+1) * sizex * sizey] - min)*scale; + il.a.s2 = (slices[idx + idy * sizex + (slice_offset+2) * sizex * sizey] - min)*scale; + il.a.s3 = (slices[idx + idy * sizex + (slice_offset+3) * sizex * sizey] - min)*scale; + il.a.s4 = (slices[idx + idy * sizex + (slice_offset+4) * sizex * sizey] - min)*scale; + il.a.s5 = (slices[idx + idy * sizex + (slice_offset+5) * sizex * sizey] - min)*scale; + il.a.s6 = (slices[idx + idy * sizex + (slice_offset+6) * sizex * sizey] - min)*scale; + il.a.s7 = (slices[idx + idy * sizex + (slice_offset+7) * sizex * sizey] - min)*scale; + + write_imageui(interleaved_slices, (int4)(idx, idy, idz, 0),(uint4)((uint)il.storage.x,(uint)il.storage.y,0,0)); + } + +kernel void +texture_uint ( + read_only image2d_array_t slices, + global uint8 *reconstructed_buffer, + float axis_pos, + float angle_step, + unsigned long size){ + + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + float angle = idy * angle_step; + float r = fmin (axis_pos, size - axis_pos); + float d = idx - axis_pos + 0.5f; + float l = sqrt(8.0f*r*r - 4.0f*d*d); + + float2 D = (float2) (cos(angle), -sin(angle)); + + D = normalize(D); + + float2 N = (float2) (D.y, -D.x); + + float2 sample = d * D - l/2.0f * N + ((float2) (axis_pos, axis_pos)); + + uint8 sum = {0,0,0,0,0,0,0,0}; + + union converter tex; + for (int i = 0; i < l; i++) { + tex.storage = read_imageui(slices, volumeSampler, (float4)((float2)sample,idz,0.0f)).xy; + + sum.s0 += (uint)tex.a.s0; + sum.s1 += (uint)tex.a.s1; + sum.s2 += (uint)tex.a.s2; + sum.s3 += (uint)tex.a.s3; + sum.s4 += (uint)tex.a.s4; + sum.s5 += (uint)tex.a.s5; + sum.s6 += (uint)tex.a.s6; + sum.s7 += (uint)tex.a.s7; + + sample += N; + } + + reconstructed_buffer[idx + idy*size + idz*size*size] = sum; +} + + kernel void + uninterleave_uint (global uint8 *reconstructed_buffer, + global float *output, + const float min, + const float max) + { + const int idx = get_global_id(0); + const int idy = get_global_id(1); + const int idz = get_global_id(2); + + const int sizex = get_global_size(0); + const int sizey = get_global_size(1); + + int output_offset = idz*8; + float scale = (max-min)/255.0f; + + output[idx + idy*sizex + (output_offset)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s0)*scale+min; + output[idx + idy*sizex + (output_offset+1)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s1)*scale+min; + output[idx + idy*sizex + (output_offset+2)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s2)*scale+min; + output[idx + idy*sizex + (output_offset+3)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s3)*scale+min; + output[idx + idy*sizex + (output_offset+4)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s4)*scale+min; + output[idx + idy*sizex + (output_offset+5)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s5)*scale+min; + output[idx + idy*sizex + (output_offset+6)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s6)*scale+min; + output[idx + idy*sizex + (output_offset+7)*sizex*sizey] = (reconstructed_buffer[idx + idy*sizex + idz*sizex*sizey].s7)*scale+min; + } + diff --git a/src/kernels/z_kernel.cl b/src/kernels/z_kernel.cl new file mode 100644 index 0000000..12da57a --- /dev/null +++ b/src/kernels/z_kernel.cl @@ -0,0 +1,478 @@ +#define rotate() pixel.x -= x_center.x; \ + pixel.y -= y_center; \ + pixel.x = pixel.x * cos_roll + pixel.y * sin_roll; \ + pixel.y = -pixel.x * sin_roll + pixel.y * cos_roll; \ + pixel.x += x_center.x; \ + pixel.y += y_center; + +/* + * Copyright (C) 2015-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 backproject_burst_1 ( +read_only image2d_t projection_0, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float sines, + const float cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + voxel.z = mad((float) idz, z_region.y, z_region.x); + tmp = mad(voxel.z, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines, mad(voxel.y, sines, x_center.x)); + pixel.y = mad(tmp_x, sines, mad(tmp_y, cosines, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_2 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float2 sines, + const float2 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + voxel.z = mad((float) idz, z_region.y, z_region.x); + tmp = mad(voxel.z, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_4 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float4 sines, + const float4 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + voxel.z = mad((float) idz, z_region.y, z_region.x); + tmp = mad(voxel.z, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_8 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float8 sines, + const float8 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + voxel.z = mad((float) idz, z_region.y, z_region.x); + tmp = mad(voxel.z, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} +/* + * Copyright (C) 2015-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 backproject_burst_16 ( +read_only image2d_t projection_0, +read_only image2d_t projection_1, +read_only image2d_t projection_2, +read_only image2d_t projection_3, +read_only image2d_t projection_4, +read_only image2d_t projection_5, +read_only image2d_t projection_6, +read_only image2d_t projection_7, +read_only image2d_t projection_8, +read_only image2d_t projection_9, +read_only image2d_t projection_10, +read_only image2d_t projection_11, +read_only image2d_t projection_12, +read_only image2d_t projection_13, +read_only image2d_t projection_14, +read_only image2d_t projection_15, + global float *volume, + const sampler_t sampler, + const int3 real_size, + const float2 x_center, + const float y_center, + const float2 x_region, + const float2 y_region, + const float2 z_region, + const float2 lamino_region, + const float2 roll_region, + float sin_lamino, + float cos_lamino, + const float16 sines, + const float16 cosines, + const float norm_factor, + float sin_roll, + float cos_roll, + const int cumulate) +{ + int idx = get_global_id (0); + int idy = get_global_id (1); + int idz = get_global_id (2); + float result, tmp, tmp_x, tmp_y; + float2 pixel; + float3 voxel; + + if (idx < real_size.x && idy < real_size.y && idz < real_size.z) { + voxel.x = mad((float) idx, x_region.y, x_region.x); + voxel.y = mad((float) idy, y_region.y, y_region.x); + voxel.z = mad((float) idz, z_region.y, z_region.x); + tmp = mad(voxel.z, sin_lamino, y_center); + tmp_x = voxel.x * cos_lamino; + tmp_y = -voxel.y * cos_lamino; + + pixel.x = mad(voxel.x, cosines.s0, mad(voxel.y, sines.s0, x_center.x)); + pixel.y = mad(tmp_x, sines.s0, mad(tmp_y, cosines.s0, tmp)); + rotate (); + result = read_imagef (projection_0, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s1, mad(voxel.y, sines.s1, x_center.x)); + pixel.y = mad(tmp_x, sines.s1, mad(tmp_y, cosines.s1, tmp)); + rotate (); + result += read_imagef (projection_1, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s2, mad(voxel.y, sines.s2, x_center.x)); + pixel.y = mad(tmp_x, sines.s2, mad(tmp_y, cosines.s2, tmp)); + rotate (); + result += read_imagef (projection_2, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s3, mad(voxel.y, sines.s3, x_center.x)); + pixel.y = mad(tmp_x, sines.s3, mad(tmp_y, cosines.s3, tmp)); + rotate (); + result += read_imagef (projection_3, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s4, mad(voxel.y, sines.s4, x_center.x)); + pixel.y = mad(tmp_x, sines.s4, mad(tmp_y, cosines.s4, tmp)); + rotate (); + result += read_imagef (projection_4, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s5, mad(voxel.y, sines.s5, x_center.x)); + pixel.y = mad(tmp_x, sines.s5, mad(tmp_y, cosines.s5, tmp)); + rotate (); + result += read_imagef (projection_5, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s6, mad(voxel.y, sines.s6, x_center.x)); + pixel.y = mad(tmp_x, sines.s6, mad(tmp_y, cosines.s6, tmp)); + rotate (); + result += read_imagef (projection_6, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s7, mad(voxel.y, sines.s7, x_center.x)); + pixel.y = mad(tmp_x, sines.s7, mad(tmp_y, cosines.s7, tmp)); + rotate (); + result += read_imagef (projection_7, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s8, mad(voxel.y, sines.s8, x_center.x)); + pixel.y = mad(tmp_x, sines.s8, mad(tmp_y, cosines.s8, tmp)); + rotate (); + result += read_imagef (projection_8, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.s9, mad(voxel.y, sines.s9, x_center.x)); + pixel.y = mad(tmp_x, sines.s9, mad(tmp_y, cosines.s9, tmp)); + rotate (); + result += read_imagef (projection_9, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sa, mad(voxel.y, sines.sa, x_center.x)); + pixel.y = mad(tmp_x, sines.sa, mad(tmp_y, cosines.sa, tmp)); + rotate (); + result += read_imagef (projection_10, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sb, mad(voxel.y, sines.sb, x_center.x)); + pixel.y = mad(tmp_x, sines.sb, mad(tmp_y, cosines.sb, tmp)); + rotate (); + result += read_imagef (projection_11, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sc, mad(voxel.y, sines.sc, x_center.x)); + pixel.y = mad(tmp_x, sines.sc, mad(tmp_y, cosines.sc, tmp)); + rotate (); + result += read_imagef (projection_12, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sd, mad(voxel.y, sines.sd, x_center.x)); + pixel.y = mad(tmp_x, sines.sd, mad(tmp_y, cosines.sd, tmp)); + rotate (); + result += read_imagef (projection_13, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.se, mad(voxel.y, sines.se, x_center.x)); + pixel.y = mad(tmp_x, sines.se, mad(tmp_y, cosines.se, tmp)); + rotate (); + result += read_imagef (projection_14, sampler, pixel).x; + pixel.x = mad(voxel.x, cosines.sf, mad(voxel.y, sines.sf, x_center.x)); + pixel.y = mad(tmp_x, sines.sf, mad(tmp_y, cosines.sf, tmp)); + rotate (); + result += read_imagef (projection_15, sampler, pixel).x; + + + if (cumulate) { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] += result * norm_factor; + } else { + volume[idz * real_size.x * real_size.y + idy * real_size.x + idx] = result * norm_factor; + } + } +} + |