summaryrefslogtreecommitdiffstats
path: root/cuda
diff options
context:
space:
mode:
Diffstat (limited to 'cuda')
-rw-r--r--cuda/2d/astra.cu91
-rw-r--r--cuda/2d/astra.h13
-rw-r--r--cuda/3d/fdk.cu23
-rw-r--r--cuda/3d/fdk.h8
4 files changed, 131 insertions, 4 deletions
diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu
index 1c2e623..e317f6c 100644
--- a/cuda/2d/astra.cu
+++ b/cuda/2d/astra.cu
@@ -33,6 +33,7 @@ $Id$
#include "par_fp.h"
#include "fan_fp.h"
#include "par_bp.h"
+#include "fan_bp.h"
#include "arith.h"
#include "astra.h"
@@ -43,6 +44,9 @@ $Id$
#include "../../include/astra/Logger.h"
+// For fan beam FBP weighting
+#include "../3d/fdk.h"
+
using namespace astraCUDA;
using namespace std;
@@ -61,8 +65,14 @@ public:
float* angles;
float* TOffsets;
+ float fOriginSourceDistance;
+ float fOriginDetectorDistance;
+
float fPixelSize;
+ bool bFanBeam;
+ bool bShortScan;
+
bool initialized;
bool setStartReconstruction;
@@ -152,9 +162,33 @@ bool AstraFBP::setProjectionGeometry(unsigned int iProjAngles,
pData->angles = new float[iProjAngles];
memcpy(pData->angles, pfAngles, iProjAngles * sizeof(pfAngles[0]));
+ pData->bFanBeam = false;
+
+ return true;
+}
+
+bool AstraFBP::setFanGeometry(unsigned int iProjAngles,
+ unsigned int iProjDets,
+ const float* pfAngles,
+ float fOriginSourceDistance,
+ float fOriginDetectorDistance,
+ float fDetSize,
+ bool bShortScan)
+{
+ // Slightly abusing setProjectionGeometry for this...
+ if (!setProjectionGeometry(iProjAngles, iProjDets, pfAngles, fDetSize))
+ return false;
+
+ pData->fOriginSourceDistance = fOriginSourceDistance;
+ pData->fOriginDetectorDistance = fOriginDetectorDistance;
+
+ pData->bFanBeam = true;
+ pData->bShortScan = bShortScan;
+
return true;
}
+
bool AstraFBP::setPixelSuperSampling(unsigned int iPixelSuperSampling)
{
if (pData->initialized)
@@ -274,6 +308,34 @@ bool AstraFBP::run()
bool ok = false;
+ if (pData->bFanBeam) {
+ // Call FDK_PreWeight to handle fan beam geometry. We treat
+ // this as a cone beam setup of a single slice:
+
+ // TODO: TOffsets...
+
+ // We create a fake cudaPitchedPtr
+ cudaPitchedPtr tmp;
+ tmp.ptr = pData->D_sinoData;
+ tmp.pitch = pData->sinoPitch * sizeof(float);
+ tmp.xsize = pData->dims.iProjDets;
+ tmp.ysize = pData->dims.iProjAngles;
+ // and a fake Dimensions3D
+ astraCUDA3d::SDimensions3D dims3d;
+ dims3d.iVolX = pData->dims.iVolWidth;
+ dims3d.iVolY = pData->dims.iVolHeight;
+ dims3d.iVolZ = 1;
+ dims3d.iProjAngles = pData->dims.iProjAngles;
+ dims3d.iProjU = pData->dims.iProjDets;
+ dims3d.iProjV = 1;
+ dims3d.iRaysPerDetDim = dims3d.iRaysPerVoxelDim = 1;
+
+ astraCUDA3d::FDK_PreWeight(tmp, pData->fOriginSourceDistance,
+ pData->fOriginDetectorDistance, 0.0f, 0.0f,
+ pData->dims.fDetScale, 1.0f, // TODO: Are these correct?
+ pData->bShortScan, dims3d, pData->angles);
+ }
+
if (pData->m_pDevFilter) {
int iFFTRealDetCount = calcNextPowerOfTwo(2 * pData->dims.iProjDets);
@@ -293,7 +355,34 @@ bool AstraFBP::run()
}
- ok = BP(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, pData->angles, pData->TOffsets);
+ if (pData->bFanBeam) {
+ // TODO: TOffsets?
+ // TODO: Remove this code duplication with CudaReconstructionAlgorithm
+ SFanProjection* projs;
+ projs = new SFanProjection[pData->dims.iProjAngles];
+
+ float fSrcX0 = 0.0f;
+ float fSrcY0 = -pData->fOriginSourceDistance / pData->fPixelSize;
+ float fDetUX0 = pData->dims.fDetScale;
+ float fDetUY0 = 0.0f;
+ float fDetSX0 = pData->dims.iProjDets * fDetUX0 / -2.0f;
+ float fDetSY0 = pData->fOriginDetectorDistance / pData->fPixelSize;
+
+#define ROTATE0(name,i,alpha) do { projs[i].f##name##X = f##name##X0 * cos(alpha) - f##name##Y0 * sin(alpha); projs[i].f##name##Y = f##name##X0 * sin(alpha) + f##name##Y0 * cos(alpha); } while(0)
+ for (unsigned int i = 0; i < pData->dims.iProjAngles; ++i) {
+ ROTATE0(Src, i, pData->angles[i]);
+ ROTATE0(DetS, i, pData->angles[i]);
+ ROTATE0(DetU, i, pData->angles[i]);
+ }
+
+#undef ROTATE0
+ ok = FanBP_FBPWeighted(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, projs);
+
+ delete[] projs;
+
+ } else {
+ ok = BP(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, pData->angles, pData->TOffsets);
+ }
if(!ok)
{
return false;
diff --git a/cuda/2d/astra.h b/cuda/2d/astra.h
index 9e58301..42b15e8 100644
--- a/cuda/2d/astra.h
+++ b/cuda/2d/astra.h
@@ -68,6 +68,19 @@ public:
unsigned int iProjDets,
const float *pfAngles,
float fDetSize = 1.0f);
+ // Set the projection angles and number of detector pixels per angle.
+ // pfAngles must be a float array of length iProjAngles.
+ // fDetSize indicates the size of a detector pixel compared to a
+ // volume pixel edge.
+ //
+ // pfAngles will only be read from during this call.
+ bool setFanGeometry(unsigned int iProjAngles,
+ unsigned int iProjDets,
+ const float *pfAngles,
+ float fOriginSourceDistance,
+ float fOriginDetectorDistance,
+ float fDetSize = 1.0f,
+ bool bShortScan = false);
// Set linear supersampling factor for the BP.
// (The number of rays is the square of this)
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,
generated by cgit v1.2.3 (git 2.25.1) at 2024-11-05 16:43:48 +0000