From 399422985fd27a1e6a1f8cea3642402128b050fa Mon Sep 17 00:00:00 2001 From: "Daniel M. Pelt" Date: Fri, 20 May 2016 15:10:03 +0200 Subject: Add option to specify custom filter for FDK --- cuda/3d/astra3d.cu | 4 ++-- cuda/3d/astra3d.h | 2 +- cuda/3d/fdk.cu | 12 ++++++++++-- cuda/3d/fdk.h | 3 ++- 4 files changed, 15 insertions(+), 6 deletions(-) (limited to 'cuda/3d') diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu index 5670873..35e3cd4 100644 --- a/cuda/3d/astra3d.cu +++ b/cuda/3d/astra3d.cu @@ -1311,7 +1311,7 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections, const CVolumeGeometry3D* pVolGeom, const CConeProjectionGeometry3D* pProjGeom, bool bShortScan, - int iGPUIndex, int iVoxelSuperSampling) + int iGPUIndex, int iVoxelSuperSampling, const float* filter) { SDimensions3D dims; @@ -1369,7 +1369,7 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections, // TODO: Offer interface for SrcZ, DetZ ok &= FDK(D_volumeData, D_projData, fOriginSourceDistance, fOriginDetectorDistance, 0, 0, fDetUSize, fDetVSize, - dims, pfAngles, bShortScan); + dims, pfAngles, bShortScan, filter); ok &= copyVolumeFromDevice(pfVolume, D_volumeData, dims, dims.iVolX); diff --git a/cuda/3d/astra3d.h b/cuda/3d/astra3d.h index 2137587..dde1347 100644 --- a/cuda/3d/astra3d.h +++ b/cuda/3d/astra3d.h @@ -314,7 +314,7 @@ _AstraExport bool astraCudaFDK(float* pfVolume, const float* pfProjections, const CVolumeGeometry3D* pVolGeom, const CConeProjectionGeometry3D* pProjGeom, bool bShortScan, - int iGPUIndex, int iVoxelSuperSampling); + int iGPUIndex, int iVoxelSuperSampling, const float* filter); } diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index 0e13be1..4899ad1 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -394,7 +394,8 @@ bool FDK(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, float fSrcOrigin, float fDetOrigin, float fSrcZ, float fDetZ, float fDetUSize, float fDetVSize, - const SDimensions3D& dims, const float* angles, bool bShortScan) + const SDimensions3D& dims, const float* angles, bool bShortScan, + const float* filter) { bool ok; // Generate filter @@ -412,7 +413,14 @@ bool FDK(cudaPitchedPtr D_volumeData, cufftComplex *pHostFilter = new cufftComplex[dims.iProjAngles * iHalfFFTSize]; memset(pHostFilter, 0, sizeof(cufftComplex) * dims.iProjAngles * iHalfFFTSize); - genFilter(FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); + if (filter==NULL){ + genFilter(FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); + }else{ + for(int i=0;i Date: Fri, 7 Oct 2016 16:42:41 +0200 Subject: Increase max angle count for Parker Weighting --- cuda/3d/fdk.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'cuda/3d') diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index d847eee..7c3ee18 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -62,7 +62,7 @@ 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]; -- cgit v1.2.3 From c599eac7c9576a74707a3fa9b3c02cff05b09760 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 7 Oct 2016 16:42:12 +0200 Subject: Clean up unused variables --- cuda/3d/fdk.cu | 29 ----------------------------- 1 file changed, 29 deletions(-) (limited to 'cuda/3d') diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index 7c3ee18..19ef338 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -46,48 +46,19 @@ $Id$ #include "../../include/astra/Logging.h" -typedef texture 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 = 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(); - - 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) { -- cgit v1.2.3 From 71db6331f8dd0d5abbeee92977af01293be4f427 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Mon, 28 Nov 2016 15:54:07 +0100 Subject: Update headers (website+2016) --- cuda/3d/algo3d.cu | 7 +++---- cuda/3d/algo3d.h | 7 +++---- cuda/3d/arith3d.cu | 7 +++---- cuda/3d/arith3d.h | 7 +++---- cuda/3d/astra3d.cu | 7 +++---- cuda/3d/astra3d.h | 7 +++---- cuda/3d/cgls3d.cu | 7 +++---- cuda/3d/cgls3d.h | 7 +++---- cuda/3d/cone_bp.cu | 7 +++---- cuda/3d/cone_bp.h | 7 +++---- cuda/3d/cone_fp.cu | 7 +++---- cuda/3d/cone_fp.h | 7 +++---- cuda/3d/darthelper3d.cu | 7 +++---- cuda/3d/darthelper3d.h | 7 +++---- cuda/3d/dims3d.h | 7 +++---- cuda/3d/fdk.cu | 7 +++---- cuda/3d/fdk.h | 7 +++---- cuda/3d/mem3d.cu | 7 +++---- cuda/3d/mem3d.h | 6 +++--- cuda/3d/par3d_bp.cu | 7 +++---- cuda/3d/par3d_bp.h | 7 +++---- cuda/3d/par3d_fp.cu | 7 +++---- cuda/3d/par3d_fp.h | 7 +++---- cuda/3d/sirt3d.cu | 7 +++---- cuda/3d/sirt3d.h | 7 +++---- cuda/3d/util3d.cu | 7 +++---- cuda/3d/util3d.h | 7 +++---- 27 files changed, 81 insertions(+), 107 deletions(-) (limited to 'cuda/3d') 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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_ARITH3D_H diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu index 3bd56ea..8e0946e 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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CONE_BP_H diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index fefcdc1..e6f01e9 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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_CONE_DIMS_H diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index bbac0be..bf41b62 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 . ----------------------------------------------------------------------- -$Id$ */ #include diff --git a/cuda/3d/fdk.h b/cuda/3d/fdk.h index fc0df20..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 . ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_FDK_H diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index cb15ab9..eae9676 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 . ----------------------------------------------------------------------- -$Id$ */ #include diff --git a/cuda/3d/mem3d.h b/cuda/3d/mem3d.h index bb8b091..0e29774 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. 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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$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 . ----------------------------------------------------------------------- -$Id$ */ #include 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 . ----------------------------------------------------------------------- -$Id$ */ #ifndef _CUDA_UTIL3D_H -- cgit v1.2.3 From 08b14da8cf13d4fe9dcc4787306991e453976494 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 30 Nov 2016 16:14:09 +0100 Subject: Tune cone_fp block size --- cuda/3d/cone_fp.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'cuda/3d') diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index e6f01e9..181f2fe 100644 --- a/cuda/3d/cone_fp.cu +++ b/cuda/3d/cone_fp.cu @@ -48,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; -- cgit v1.2.3 From 55dabbf035b55f71c4261c9de7ef572da46300ff Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Thu, 1 Dec 2016 14:03:56 +0100 Subject: Expose the density weighting option of cone_bp It is now exposed via the new DensityWeighting option of CudaProjector3D. --- cuda/3d/mem3d.cu | 4 +++- cuda/3d/mem3d.h | 2 +- 2 files changed, 4 insertions(+), 2 deletions(-) (limited to 'cuda/3d') diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index eae9676..2b26fe1 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -249,7 +249,7 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con 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; @@ -269,6 +269,8 @@ 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 diff --git a/cuda/3d/mem3d.h b/cuda/3d/mem3d.h index 0e29774..a0829e2 100644 --- a/cuda/3d/mem3d.h +++ b/cuda/3d/mem3d.h @@ -93,7 +93,7 @@ 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 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); -- cgit v1.2.3 From d85a660f064e8130b27e11c7fd762221c754c315 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Thu, 26 Jan 2017 14:57:57 +0100 Subject: Start work on CFloat32Data3DGPU to allow persistent/external GPU memory --- cuda/3d/mem3d.cu | 24 ++++++++++++++++++++++++ cuda/3d/mem3d.h | 5 ++++- 2 files changed, 28 insertions(+), 1 deletion(-) (limited to 'cuda/3d') diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index 2b26fe1..97be8a4 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -118,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(); @@ -307,6 +314,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(new SMemHandle3D_internal); + *hnd.d = h; + + return hnd; +} diff --git a/cuda/3d/mem3d.h b/cuda/3d/mem3d.h index a0829e2..7a87ae6 100644 --- a/cuda/3d/mem3d.h +++ b/cuda/3d/mem3d.h @@ -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,6 +90,8 @@ 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); @@ -97,7 +101,6 @@ bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan, const float *pfFilter = 0); - } #endif -- cgit v1.2.3 From f591df5ff62deb8de2956cc8200db4034fb6ea05 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Tue, 14 Feb 2017 14:10:08 +0100 Subject: Fix FDK shortscan weighting It was computing weights for angles in [0,2pi) but using them on angles in (-2pi,0]. --- cuda/3d/fdk.cu | 21 +++++++-------------- 1 file changed, 7 insertions(+), 14 deletions(-) (limited to 'cuda/3d') diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index bf41b62..27357ad 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -121,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 @@ -195,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, -- cgit v1.2.3 From aa325035704f932e29217274aeab26ae46847a7b Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 17 Feb 2017 15:56:36 +0100 Subject: Fix cuda3d geometry memory leaks --- cuda/3d/astra3d.cu | 38 ++++++++++++++++++++++++++++++++++---- cuda/3d/mem3d.cu | 14 +++++++++++++- 2 files changed, 47 insertions(+), 5 deletions(-) (limited to 'cuda/3d') diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu index 8e0946e..30cbe1b 100644 --- a/cuda/3d/astra3d.cu +++ b/cuda/3d/astra3d.cu @@ -319,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; @@ -702,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; @@ -1133,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; } @@ -1156,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; @@ -1168,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); @@ -1208,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; } @@ -1228,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; @@ -1244,6 +1272,8 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, processVol3D(D_pixelWeight, dims); if (!ok) { + delete[] pParProjs; + delete[] pConeProjs; cudaFree(D_pixelWeight.ptr); cudaFree(D_volumeData.ptr); cudaFree(D_projData.ptr); diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index 97be8a4..ed779fa 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -253,6 +253,9 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con } } + delete[] pParProjs; + delete[] pConeProjs; + return ok; } @@ -283,6 +286,9 @@ bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con else ok &= ConeBP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, params); + delete[] pParProjs; + delete[] pConeProjs; + return ok; } @@ -303,11 +309,17 @@ 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, pfFilter); + delete[] pParProjs; + delete[] pConeProjs; + return ok; -- cgit v1.2.3