summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be>2014-04-16 11:13:22 +0000
committerwpalenst <WillemJan.Palenstijn@uantwerpen.be>2014-04-16 11:13:22 +0000
commit0c77eee16e2f4161c1ebc110b15ce6563d4a96c2 (patch)
treed824a763da5578391839c57057d5ecee26554873 /cuda/3d
parent8b046691e7cf5ba603b690e5a820869f7aba0bb6 (diff)
downloadastra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.gz
astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.bz2
astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.xz
astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.zip
Add fan beam support to FBP_CUDA
Diffstat (limited to 'cuda/3d')
-rw-r--r--cuda/3d/fdk.cu23
-rw-r--r--cuda/3d/fdk.h8
2 files changed, 28 insertions, 3 deletions
diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu
index 883187d..fd5a8a2 100644
--- a/cuda/3d/fdk.cu
+++ b/cuda/3d/fdk.cu
@@ -298,8 +298,7 @@ __global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, un
// Perform the FDK pre-weighting and filtering
-bool FDK_Filter(cudaPitchedPtr D_projData,
- cufftComplex * D_filter,
+bool FDK_PreWeight(cudaPitchedPtr D_projData,
float fSrcOrigin, float fDetOrigin,
float fSrcZ, float fDetZ,
float fDetUSize, float fDetVSize, bool bShortScan,
@@ -335,13 +334,22 @@ bool FDK_Filter(cudaPitchedPtr D_projData,
}
cudaTextForceKernelsCompletion();
+ return true;
+}
+bool FDK_Filter(cudaPitchedPtr D_projData,
+ cufftComplex * D_filter,
+ float fSrcOrigin, float fDetOrigin,
+ float fSrcZ, float fDetZ,
+ float fDetUSize, float fDetVSize, bool bShortScan,
+ const SDimensions3D& dims, const float* angles)
+{
// The filtering is a regular ramp filter per detector line.
int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU);
int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount);
-
+ int projPitch = D_projData.pitch/sizeof(float);
// We process one sinogram at a time.
@@ -390,11 +398,18 @@ bool FDK(cudaPitchedPtr D_volumeData,
int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU);
int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount);
+ ok = FDK_PreWeight(D_projData, fSrcOrigin, fDetOrigin,
+ fSrcZ, fDetZ, fDetUSize, fDetVSize,
+ bShortScan, dims, angles);
+ if (!ok)
+ return false;
+
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);
+
allocateComplexOnDevice(dims.iProjAngles, iHalfFFTSize, &D_filter);
uploadComplexArrayToDevice(dims.iProjAngles, iHalfFFTSize, pHostFilter, D_filter);
@@ -403,6 +418,8 @@ bool FDK(cudaPitchedPtr D_volumeData,
// Perform filtering
+
+
ok = FDK_Filter(D_projData, D_filter, fSrcOrigin, fDetOrigin,
fSrcZ, fDetZ, fDetUSize, fDetVSize,
bShortScan, dims, angles);
diff --git a/cuda/3d/fdk.h b/cuda/3d/fdk.h
index 5443b19..c9123b9 100644
--- a/cuda/3d/fdk.h
+++ b/cuda/3d/fdk.h
@@ -29,8 +29,16 @@ $Id$
#ifndef _CUDA_FDK_H
#define _CUDA_FDK_H
+#include "dims3d.h"
+
namespace astraCUDA3d {
+bool FDK_PreWeight(cudaPitchedPtr D_projData,
+ float fSrcOrigin, float fDetOrigin,
+ float fSrcZ, float fDetZ,
+ float fDetUSize, float fDetVSize, bool bShortScan,
+ const SDimensions3D& dims, const float* angles);
+
bool FDK(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
float fSrcOrigin, float fDetOrigin,