summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/3d')
-rw-r--r--cuda/3d/algo3d.cu7
-rw-r--r--cuda/3d/algo3d.h7
-rw-r--r--cuda/3d/arith3d.cu7
-rw-r--r--cuda/3d/arith3d.h7
-rw-r--r--cuda/3d/astra3d.cu47
-rw-r--r--cuda/3d/astra3d.h7
-rw-r--r--cuda/3d/cgls3d.cu7
-rw-r--r--cuda/3d/cgls3d.h7
-rw-r--r--cuda/3d/cone_bp.cu7
-rw-r--r--cuda/3d/cone_bp.h7
-rw-r--r--cuda/3d/cone_fp.cu9
-rw-r--r--cuda/3d/cone_fp.h7
-rw-r--r--cuda/3d/darthelper3d.cu7
-rw-r--r--cuda/3d/darthelper3d.h7
-rw-r--r--cuda/3d/dims3d.h7
-rw-r--r--cuda/3d/fdk.cu71
-rw-r--r--cuda/3d/fdk.h10
-rw-r--r--cuda/3d/mem3d.cu53
-rw-r--r--cuda/3d/mem3d.h15
-rw-r--r--cuda/3d/par3d_bp.cu7
-rw-r--r--cuda/3d/par3d_bp.h7
-rw-r--r--cuda/3d/par3d_fp.cu7
-rw-r--r--cuda/3d/par3d_fp.h7
-rw-r--r--cuda/3d/sirt3d.cu7
-rw-r--r--cuda/3d/sirt3d.h7
-rw-r--r--cuda/3d/util3d.cu7
-rw-r--r--cuda/3d/util3d.h7
27 files changed, 184 insertions, 168 deletions
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
generated by cgit v1.2.3 (git 2.25.1) at 2024-11-11 05:03:46 +0000