diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2017-09-15 16:38:08 +0200 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2017-09-15 16:38:08 +0200 |
| commit | aa31a06235496c0d808e57c8ce914cb4b640bc6e (patch) | |
| tree | 33385e828ddca0b2857bac9e3dac4dd3723a3eee /cuda | |
| parent | f6aa2db83dfea89f9d2cfc6fcbd3da141ee77e02 (diff) | |
| parent | 00a1c6118b2d64b867c8e640c295462bcccfc7c9 (diff) | |
| download | astra-aa31a06235496c0d808e57c8ce914cb4b640bc6e.tar.gz astra-aa31a06235496c0d808e57c8ce914cb4b640bc6e.tar.bz2 astra-aa31a06235496c0d808e57c8ce914cb4b640bc6e.tar.xz astra-aa31a06235496c0d808e57c8ce914cb4b640bc6e.zip | |
Merge branch 'master' into parallel_vec
Diffstat (limited to 'cuda')
57 files changed, 274 insertions, 288 deletions
diff --git a/cuda/2d/algo.cu b/cuda/2d/algo.cu index bde4f42..3275f6d 100644 --- a/cuda/2d/algo.cu +++ b/cuda/2d/algo.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cassert> diff --git a/cuda/2d/algo.h b/cuda/2d/algo.h index 0b6a066..c1d932c 100644 --- a/cuda/2d/algo.h +++ b/cuda/2d/algo.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ALGO_H diff --git a/cuda/2d/arith.cu b/cuda/2d/arith.cu index 110ed51..9805093 100644 --- a/cuda/2d/arith.cu +++ b/cuda/2d/arith.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include "util.h" diff --git a/cuda/2d/arith.h b/cuda/2d/arith.h index 2c648a1..5b1c270 100644 --- a/cuda/2d/arith.h +++ b/cuda/2d/arith.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ARITH_H diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu index b39e0a9..279b57d 100644 --- a/cuda/2d/astra.cu +++ b/cuda/2d/astra.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/astra.h b/cuda/2d/astra.h index 8e91feb..a20b830 100644 --- a/cuda/2d/astra.h +++ b/cuda/2d/astra.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ASTRA_H diff --git a/cuda/2d/cgls.cu b/cuda/2d/cgls.cu index 3f06581..1a31438 100644 --- a/cuda/2d/cgls.cu +++ b/cuda/2d/cgls.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/cgls.h b/cuda/2d/cgls.h index c55dd42..804f943 100644 --- a/cuda/2d/cgls.h +++ b/cuda/2d/cgls.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CGLS_H diff --git a/cuda/2d/darthelper.cu b/cuda/2d/darthelper.cu index 1d10d49..744184e 100644 --- a/cuda/2d/darthelper.cu +++ b/cuda/2d/darthelper.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include "util.h" diff --git a/cuda/2d/darthelper.h b/cuda/2d/darthelper.h index acd14e3..a2f1f45 100644 --- a/cuda/2d/darthelper.h +++ b/cuda/2d/darthelper.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ARITH2_H diff --git a/cuda/2d/dims.h b/cuda/2d/dims.h index f74d0e4..10eaabf 100644 --- a/cuda/2d/dims.h +++ b/cuda/2d/dims.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_DIMS_H diff --git a/cuda/2d/em.cu b/cuda/2d/em.cu index 47ec548..03cc652 100644 --- a/cuda/2d/em.cu +++ b/cuda/2d/em.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/em.h b/cuda/2d/em.h index 9b62fa5..f99e798 100644 --- a/cuda/2d/em.h +++ b/cuda/2d/em.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_EM_H diff --git a/cuda/2d/fan_bp.cu b/cuda/2d/fan_bp.cu index b4321ba..2025528 100644 --- a/cuda/2d/fan_bp.cu +++ b/cuda/2d/fan_bp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/fan_bp.h b/cuda/2d/fan_bp.h index 3ebe1e8..4e466fc 100644 --- a/cuda/2d/fan_bp.h +++ b/cuda/2d/fan_bp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_FAN_BP_H diff --git a/cuda/2d/fan_fp.cu b/cuda/2d/fan_fp.cu index 7c9d8cc..494e89b 100644 --- a/cuda/2d/fan_fp.cu +++ b/cuda/2d/fan_fp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/fan_fp.h b/cuda/2d/fan_fp.h index 3962ee3..4bc07eb 100644 --- a/cuda/2d/fan_fp.h +++ b/cuda/2d/fan_fp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_FAN_FP_H diff --git a/cuda/2d/fbp_filters.h b/cuda/2d/fbp_filters.h index 71c7572..bf1342d 100644 --- a/cuda/2d/fbp_filters.h +++ b/cuda/2d/fbp_filters.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef FBP_FILTERS_H diff --git a/cuda/2d/fft.cu b/cuda/2d/fft.cu index f6d3cc7..5500e14 100644 --- a/cuda/2d/fft.cu +++ b/cuda/2d/fft.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include "fft.h" diff --git a/cuda/2d/fft.h b/cuda/2d/fft.h index b29f17a..3e91043 100644 --- a/cuda/2d/fft.h +++ b/cuda/2d/fft.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef FFT_H diff --git a/cuda/2d/par_bp.cu b/cuda/2d/par_bp.cu index cf0a684..307c561 100644 --- a/cuda/2d/par_bp.cu +++ b/cuda/2d/par_bp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/par_bp.h b/cuda/2d/par_bp.h index 102cb2b..3fa8a28 100644 --- a/cuda/2d/par_bp.h +++ b/cuda/2d/par_bp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_PAR_BP_H diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu index 3c010b3..10ce683 100644 --- a/cuda/2d/par_fp.cu +++ b/cuda/2d/par_fp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/par_fp.h b/cuda/2d/par_fp.h index 5fd593c..1f0ff69 100644 --- a/cuda/2d/par_fp.h +++ b/cuda/2d/par_fp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_PAR_FP_H diff --git a/cuda/2d/sart.cu b/cuda/2d/sart.cu index a1085cd..3a3c7e6 100644 --- a/cuda/2d/sart.cu +++ b/cuda/2d/sart.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/sart.h b/cuda/2d/sart.h index eff9ecf..ab99e66 100644 --- a/cuda/2d/sart.h +++ b/cuda/2d/sart.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_SART_H diff --git a/cuda/2d/sirt.cu b/cuda/2d/sirt.cu index 9dc4f64..2516c6c 100644 --- a/cuda/2d/sirt.cu +++ b/cuda/2d/sirt.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/sirt.h b/cuda/2d/sirt.h index bc913f4..488ab0a 100644 --- a/cuda/2d/sirt.h +++ b/cuda/2d/sirt.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_SIRT_H diff --git a/cuda/2d/util.cu b/cuda/2d/util.cu index a4f8f3e..09d1a2b 100644 --- a/cuda/2d/util.cu +++ b/cuda/2d/util.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/2d/util.h b/cuda/2d/util.h index ff438c3..875aae3 100644 --- a/cuda/2d/util.h +++ b/cuda/2d/util.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_UTIL_H diff --git a/cuda/3d/algo3d.cu b/cuda/3d/algo3d.cu index e2c1e43..437c493 100644 --- a/cuda/3d/algo3d.cu +++ b/cuda/3d/algo3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cassert> diff --git a/cuda/3d/algo3d.h b/cuda/3d/algo3d.h index ce331ad..6ba8b6f 100644 --- a/cuda/3d/algo3d.h +++ b/cuda/3d/algo3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ALGO_H diff --git a/cuda/3d/arith3d.cu b/cuda/3d/arith3d.cu index 471b5cd..fd41812 100644 --- a/cuda/3d/arith3d.cu +++ b/cuda/3d/arith3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include "util3d.h" diff --git a/cuda/3d/arith3d.h b/cuda/3d/arith3d.h index 85d3d3b..c42603e 100644 --- a/cuda/3d/arith3d.h +++ b/cuda/3d/arith3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ARITH3D_H diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu index 91f4530..30cbe1b 100644 --- a/cuda/3d/astra3d.cu +++ b/cuda/3d/astra3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> @@ -320,6 +319,9 @@ AstraSIRT3d::~AstraSIRT3d() delete[] pData->projs; pData->projs = 0; + delete[] pData->parprojs; + pData->parprojs = 0; + cudaFree(pData->D_projData.ptr); pData->D_projData.ptr = 0; @@ -703,6 +705,9 @@ AstraCGLS3d::~AstraCGLS3d() delete[] pData->projs; pData->projs = 0; + delete[] pData->parprojs; + pData->parprojs = 0; + cudaFree(pData->D_projData.ptr); pData->D_projData.ptr = 0; @@ -1134,19 +1139,27 @@ bool astraCudaBP(float* pfVolume, const float* pfProjections, cudaError_t err = cudaGetLastError(); // Ignore errors caused by calling cudaSetDevice multiple times - if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess) + if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess) { + delete[] pParProjs; + delete[] pConeProjs; return false; + } } cudaPitchedPtr D_volumeData = allocateVolumeData(dims); ok = D_volumeData.ptr; - if (!ok) + if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; return false; + } cudaPitchedPtr D_projData = allocateProjectionData(dims); ok = D_projData.ptr; if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_volumeData.ptr); return false; } @@ -1157,6 +1170,8 @@ bool astraCudaBP(float* pfVolume, const float* pfProjections, ok &= zeroVolumeData(D_volumeData, dims); if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_volumeData.ptr); cudaFree(D_projData.ptr); return false; @@ -1169,6 +1184,8 @@ bool astraCudaBP(float* pfVolume, const float* pfProjections, ok &= copyVolumeFromDevice(pfVolume, D_volumeData, dims, dims.iVolX); + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_volumeData.ptr); cudaFree(D_projData.ptr); @@ -1209,19 +1226,27 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, cudaError_t err = cudaGetLastError(); // Ignore errors caused by calling cudaSetDevice multiple times - if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess) + if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess) { + delete[] pParProjs; + delete[] pConeProjs; return false; + } } cudaPitchedPtr D_pixelWeight = allocateVolumeData(dims); ok = D_pixelWeight.ptr; - if (!ok) + if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; return false; + } cudaPitchedPtr D_volumeData = allocateVolumeData(dims); ok = D_volumeData.ptr; if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_pixelWeight.ptr); return false; } @@ -1229,6 +1254,8 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, cudaPitchedPtr D_projData = allocateProjectionData(dims); ok = D_projData.ptr; if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_pixelWeight.ptr); cudaFree(D_volumeData.ptr); return false; @@ -1245,6 +1272,8 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, processVol3D<opInvert>(D_pixelWeight, dims); if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_pixelWeight.ptr); cudaFree(D_volumeData.ptr); cudaFree(D_projData.ptr); @@ -1289,6 +1318,4 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, } - - } diff --git a/cuda/3d/astra3d.h b/cuda/3d/astra3d.h index 93bf576..5267899 100644 --- a/cuda/3d/astra3d.h +++ b/cuda/3d/astra3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ASTRA3D_H diff --git a/cuda/3d/cgls3d.cu b/cuda/3d/cgls3d.cu index 0b0331e..7200144 100644 --- a/cuda/3d/cgls3d.cu +++ b/cuda/3d/cgls3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/cgls3d.h b/cuda/3d/cgls3d.h index bfb4af8..e09fcfb 100644 --- a/cuda/3d/cgls3d.h +++ b/cuda/3d/cgls3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CGLS3D_H diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu index 6bd9d16..c3405b8 100644 --- a/cuda/3d/cone_bp.cu +++ b/cuda/3d/cone_bp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/cone_bp.h b/cuda/3d/cone_bp.h index a6e2c23..3952ab9 100644 --- a/cuda/3d/cone_bp.h +++ b/cuda/3d/cone_bp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CONE_BP_H diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index fefcdc1..181f2fe 100644 --- a/cuda/3d/cone_fp.cu +++ b/cuda/3d/cone_fp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> @@ -49,7 +48,7 @@ namespace astraCUDA3d { static const unsigned int g_anglesPerBlock = 4; // thickness of the slices we're splitting the volume up into -static const unsigned int g_blockSlices = 32; +static const unsigned int g_blockSlices = 4; static const unsigned int g_detBlockU = 32; static const unsigned int g_detBlockV = 32; diff --git a/cuda/3d/cone_fp.h b/cuda/3d/cone_fp.h index 62c9caa..d46990b 100644 --- a/cuda/3d/cone_fp.h +++ b/cuda/3d/cone_fp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CONE_FP_H diff --git a/cuda/3d/darthelper3d.cu b/cuda/3d/darthelper3d.cu index d9bd636..1457017 100644 --- a/cuda/3d/darthelper3d.cu +++ b/cuda/3d/darthelper3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include "util3d.h" diff --git a/cuda/3d/darthelper3d.h b/cuda/3d/darthelper3d.h index fdf43b9..71ea5b0 100644 --- a/cuda/3d/darthelper3d.h +++ b/cuda/3d/darthelper3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_DARTHELPER3_H diff --git a/cuda/3d/dims3d.h b/cuda/3d/dims3d.h index eb7045f..24c334a 100644 --- a/cuda/3d/dims3d.h +++ b/cuda/3d/dims3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CONE_DIMS_H diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index bb2393c..91b0219 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> @@ -46,48 +45,19 @@ $Id$ #include "../../include/astra/Logging.h" -typedef texture<float, 3, cudaReadModeElementType> texture3D; - -static texture3D gT_coneProjTexture; - namespace astraCUDA3d { -static const unsigned int g_volBlockZ = 16; - -static const unsigned int g_anglesPerBlock = 64; -static const unsigned int g_volBlockX = 32; -static const unsigned int g_volBlockY = 16; - static const unsigned int g_anglesPerWeightBlock = 16; static const unsigned int g_detBlockU = 32; static const unsigned int g_detBlockV = 32; -static const unsigned g_MaxAngles = 2048; +static const unsigned g_MaxAngles = 12000; -__constant__ float gC_angle_sin[g_MaxAngles]; -__constant__ float gC_angle_cos[g_MaxAngles]; __constant__ float gC_angle[g_MaxAngles]; // per-detector u/v shifts? -static bool bindProjDataTexture(const cudaArray* array) -{ - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); - - gT_coneProjTexture.addressMode[0] = cudaAddressModeBorder; - gT_coneProjTexture.addressMode[1] = cudaAddressModeBorder; - gT_coneProjTexture.addressMode[2] = cudaAddressModeBorder; - gT_coneProjTexture.filterMode = cudaFilterModeLinear; - gT_coneProjTexture.normalized = false; - - cudaBindTextureToArray(gT_coneProjTexture, array, channelDesc); - - // TODO: error value? - - return true; -} - __global__ void devFDK_preweight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fZShift, float fDetUSize, float fDetVSize, const SDimensions3D dims) { @@ -151,14 +121,8 @@ __global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, un // TODO: Detector pixel size const float fU = (detectorU - 0.5f*dims.iProjU + 0.5f) * fDetUSize; - //const float fGamma = atanf(fU / fCentralRayLength); - //const float fBeta = gC_angle[angle]; const float fGamma = atanf(fU / fCentralRayLength); - float fBeta = -gC_angle[angle]; - if (fBeta < 0.0f) - fBeta += 2*M_PI; - if (fBeta >= 2*M_PI) - fBeta -= 2*M_PI; + float fBeta = gC_angle[angle]; // compute the weight depending on the location in the central fan's radon // space @@ -225,24 +189,23 @@ bool FDK_PreWeight(cudaPitchedPtr D_projData, float fAngleBase; if (fdA >= 0.0f) { // going up from angles[0] - fAngleBase = angles[dims.iProjAngles - 1]; + fAngleBase = angles[0]; } else { // going down from angles[0] - fAngleBase = angles[0]; + fAngleBase = angles[dims.iProjAngles - 1]; } - // We pick the highest end of the range, and then - // move all angles so they fall in (-2pi,0] + // We pick the lowest end of the range, and then + // move all angles so they fall in [0,2pi) float *fRelAngles = new float[dims.iProjAngles]; for (unsigned int i = 0; i < dims.iProjAngles; ++i) { float f = angles[i] - fAngleBase; - while (f > 0) + while (f >= 2*M_PI) f -= 2*M_PI; - while (f <= -2*M_PI) + while (f < 0) f += 2*M_PI; fRelAngles[i] = f; - } cudaError_t e1 = cudaMemcpyToSymbol(gC_angle, fRelAngles, @@ -310,7 +273,8 @@ bool FDK_Filter(cudaPitchedPtr D_projData, bool FDK(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SConeProjection* angles, - const SDimensions3D& dims, SProjectorParams3D params, bool bShortScan) + const SDimensions3D& dims, SProjectorParams3D params, bool bShortScan, + const float* pfFilter) { bool ok; // Generate filter @@ -361,7 +325,14 @@ bool FDK(cudaPitchedPtr D_volumeData, cufftComplex *pHostFilter = new cufftComplex[dims.iProjAngles * iHalfFFTSize]; memset(pHostFilter, 0, sizeof(cufftComplex) * dims.iProjAngles * iHalfFFTSize); - astraCUDA::genFilter(astra::FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); + if (pfFilter == 0){ + astraCUDA::genFilter(astra::FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); + } else { + for (int i = 0; i < dims.iProjAngles * iHalfFFTSize; i++) { + pHostFilter[i].x = pfFilter[i]; + pHostFilter[i].y = 0; + } + } astraCUDA::allocateComplexOnDevice(dims.iProjAngles, iHalfFFTSize, &D_filter); diff --git a/cuda/3d/fdk.h b/cuda/3d/fdk.h index 8d2a128..0165af9 100644 --- a/cuda/3d/fdk.h +++ b/cuda/3d/fdk.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_FDK_H @@ -42,7 +41,8 @@ bool FDK_PreWeight(cudaPitchedPtr D_projData, bool FDK(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SConeProjection* angles, - const SDimensions3D& dims, SProjectorParams3D params, bool bShortScan); + const SDimensions3D& dims, SProjectorParams3D params, bool bShortScan, + const float* filter); } diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index 4fb30a2..ed779fa 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> @@ -119,6 +118,13 @@ MemHandle3D allocateGPUMemory(unsigned int x, unsigned int y, unsigned int z, Me return ret; } +bool zeroGPUMemory(MemHandle3D handle, unsigned int x, unsigned int y, unsigned int z) +{ + SMemHandle3D_internal& hnd = *handle.d.get(); + cudaError_t err = cudaMemset3D(hnd.ptr, 0, make_cudaExtent(sizeof(float)*x, y, z)); + return err == cudaSuccess; +} + bool freeGPUMemory(MemHandle3D handle) { size_t free = availableGPUMemory(); @@ -247,10 +253,13 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con } } + delete[] pParProjs; + delete[] pConeProjs; + return ok; } -bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling) +bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling, bool bFDKWeighting) { SDimensions3D dims; SProjectorParams3D params; @@ -270,16 +279,21 @@ bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con pParProjs, pConeProjs, params); + params.bFDKWeighting = bFDKWeighting; + if (pParProjs) ok &= Par3DBP(volData.d->ptr, projData.d->ptr, dims, pParProjs, params); else ok &= ConeBP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, params); + delete[] pParProjs; + delete[] pConeProjs; + return ok; } -bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan) +bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan, const float *pfFilter) { SDimensions3D dims; SProjectorParams3D params; @@ -295,10 +309,16 @@ bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, co pParProjs, pConeProjs, params); - if (!ok || !pConeProjs) + if (!ok || !pConeProjs) { + delete[] pParProjs; + delete[] pConeProjs; return false; + } - ok &= FDK(volData.d->ptr, projData.d->ptr, pConeProjs, dims, params, bShortScan); + ok &= FDK(volData.d->ptr, projData.d->ptr, pConeProjs, dims, params, bShortScan, pfFilter); + + delete[] pParProjs; + delete[] pConeProjs; return ok; @@ -306,6 +326,23 @@ bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, co } +MemHandle3D wrapHandle(float *D_ptr, unsigned int x, unsigned int y, unsigned int z, unsigned int pitch) +{ + cudaPitchedPtr ptr; + ptr.ptr = D_ptr; + ptr.xsize = sizeof(float) * x; + ptr.pitch = sizeof(float) * pitch; + ptr.ysize = y; + + SMemHandle3D_internal h; + h.ptr = ptr; + + MemHandle3D hnd; + hnd.d = boost::shared_ptr<SMemHandle3D_internal>(new SMemHandle3D_internal); + *hnd.d = h; + + return hnd; +} diff --git a/cuda/3d/mem3d.h b/cuda/3d/mem3d.h index 2d0fb27..7a87ae6 100644 --- a/cuda/3d/mem3d.h +++ b/cuda/3d/mem3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -80,6 +80,8 @@ enum Mem3DZeroMode { size_t availableGPUMemory(); int maxBlockDimension(); +MemHandle3D wrapHandle(float *D_ptr, unsigned int x, unsigned int y, unsigned int z, unsigned int pitch); + MemHandle3D allocateGPUMemory(unsigned int x, unsigned int y, unsigned int z, Mem3DZeroMode zero); bool copyToGPUMemory(const float *src, MemHandle3D dst, const SSubDimensions3D &pos); @@ -88,15 +90,16 @@ bool copyFromGPUMemory(float *dst, MemHandle3D src, const SSubDimensions3D &pos) bool freeGPUMemory(MemHandle3D handle); +bool zeroGPUMemory(MemHandle3D handle, unsigned int x, unsigned int y, unsigned int z); + bool setGPUIndex(int index); bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iDetectorSuperSampling, astra::Cuda3DProjectionKernel projKernel); -bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling); - -bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan); +bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling, bool bFDKWeighting); +bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan, const float *pfFilter = 0); } diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 491ee2f..e43479a 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/par3d_bp.h b/cuda/3d/par3d_bp.h index 3efad19..4454395 100644 --- a/cuda/3d/par3d_bp.h +++ b/cuda/3d/par3d_bp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_PAR3D_BP_H diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu index 6f9ce40..a99308f 100644 --- a/cuda/3d/par3d_fp.cu +++ b/cuda/3d/par3d_fp.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/par3d_fp.h b/cuda/3d/par3d_fp.h index 5f65cd9..c69ab51 100644 --- a/cuda/3d/par3d_fp.h +++ b/cuda/3d/par3d_fp.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_PAR3D_FP_H diff --git a/cuda/3d/sirt3d.cu b/cuda/3d/sirt3d.cu index ba6a159..4d1ed81 100644 --- a/cuda/3d/sirt3d.cu +++ b/cuda/3d/sirt3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/sirt3d.h b/cuda/3d/sirt3d.h index 5e93deb..69031b8 100644 --- a/cuda/3d/sirt3d.h +++ b/cuda/3d/sirt3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_SIRT3D_H diff --git a/cuda/3d/util3d.cu b/cuda/3d/util3d.cu index 537ed69..e7b8fbb 100644 --- a/cuda/3d/util3d.cu +++ b/cuda/3d/util3d.cu @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #include <cstdio> diff --git a/cuda/3d/util3d.h b/cuda/3d/util3d.h index fa0a03f..7dca762 100644 --- a/cuda/3d/util3d.h +++ b/cuda/3d/util3d.h @@ -1,10 +1,10 @@ /* ----------------------------------------------------------------------- -Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp - 2014-2015, CWI, Amsterdam +Copyright: 2010-2016, iMinds-Vision Lab, University of Antwerp + 2014-2016, CWI, Amsterdam Contact: astra@uantwerpen.be -Website: http://sf.net/projects/astra-toolbox +Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. @@ -23,7 +23,6 @@ You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_UTIL3D_H |