summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/3d')
-rw-r--r--cuda/3d/algo3d.cu18
-rw-r--r--cuda/3d/algo3d.h9
-rw-r--r--cuda/3d/astra3d.cu967
-rw-r--r--cuda/3d/astra3d.h208
-rw-r--r--cuda/3d/cgls3d.cu6
-rw-r--r--cuda/3d/cone_bp.cu26
-rw-r--r--cuda/3d/cone_bp.h7
-rw-r--r--cuda/3d/par3d_bp.cu25
-rw-r--r--cuda/3d/par3d_bp.h6
-rw-r--r--cuda/3d/sirt3d.cu8
10 files changed, 425 insertions, 855 deletions
diff --git a/cuda/3d/algo3d.cu b/cuda/3d/algo3d.cu
index 7f61280..cc86b70 100644
--- a/cuda/3d/algo3d.cu
+++ b/cuda/3d/algo3d.cu
@@ -41,6 +41,7 @@ ReconAlgo3D::ReconAlgo3D()
coneProjs = 0;
par3DProjs = 0;
shouldAbort = false;
+ fOutputScale = 1.0f;
}
ReconAlgo3D::~ReconAlgo3D()
@@ -57,9 +58,10 @@ void ReconAlgo3D::reset()
shouldAbort = false;
}
-bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProjection* _angles)
+bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProjection* _angles, float _outputScale)
{
dims = _dims;
+ fOutputScale = _outputScale;
coneProjs = new SConeProjection[dims.iProjAngles];
par3DProjs = 0;
@@ -69,9 +71,10 @@ bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProject
return true;
}
-bool ReconAlgo3D::setPar3DGeometry(const SDimensions3D& _dims, const SPar3DProjection* _angles)
+bool ReconAlgo3D::setPar3DGeometry(const SDimensions3D& _dims, const SPar3DProjection* _angles, float _outputScale)
{
dims = _dims;
+ fOutputScale = _outputScale;
par3DProjs = new SPar3DProjection[dims.iProjAngles];
coneProjs = 0;
@@ -87,19 +90,20 @@ bool ReconAlgo3D::callFP(cudaPitchedPtr& D_volumeData,
float outputScale)
{
if (coneProjs) {
- return ConeFP(D_volumeData, D_projData, dims, coneProjs, outputScale);
+ return ConeFP(D_volumeData, D_projData, dims, coneProjs, outputScale * this->fOutputScale);
} else {
- return Par3DFP(D_volumeData, D_projData, dims, par3DProjs, outputScale);
+ return Par3DFP(D_volumeData, D_projData, dims, par3DProjs, outputScale * this->fOutputScale);
}
}
bool ReconAlgo3D::callBP(cudaPitchedPtr& D_volumeData,
- cudaPitchedPtr& D_projData)
+ cudaPitchedPtr& D_projData,
+ float outputScale)
{
if (coneProjs) {
- return ConeBP(D_volumeData, D_projData, dims, coneProjs);
+ return ConeBP(D_volumeData, D_projData, dims, coneProjs, outputScale * this->fOutputScale);
} else {
- return Par3DBP(D_volumeData, D_projData, dims, par3DProjs);
+ return Par3DBP(D_volumeData, D_projData, dims, par3DProjs, outputScale * this->fOutputScale);
}
}
diff --git a/cuda/3d/algo3d.h b/cuda/3d/algo3d.h
index f4c6a87..886b092 100644
--- a/cuda/3d/algo3d.h
+++ b/cuda/3d/algo3d.h
@@ -39,8 +39,8 @@ public:
ReconAlgo3D();
~ReconAlgo3D();
- bool setConeGeometry(const SDimensions3D& dims, const SConeProjection* projs);
- bool setPar3DGeometry(const SDimensions3D& dims, const SPar3DProjection* projs);
+ bool setConeGeometry(const SDimensions3D& dims, const SConeProjection* projs, float fOutputScale);
+ bool setPar3DGeometry(const SDimensions3D& dims, const SPar3DProjection* projs, float fOutputScale);
void signalAbort() { shouldAbort = true; }
@@ -51,12 +51,15 @@ protected:
cudaPitchedPtr& D_projData,
float outputScale);
bool callBP(cudaPitchedPtr& D_volumeData,
- cudaPitchedPtr& D_projData);
+ cudaPitchedPtr& D_projData,
+ float outputScale);
SDimensions3D dims;
SConeProjection* coneProjs;
SPar3DProjection* par3DProjs;
+ float fOutputScale;
+
volatile bool shouldAbort;
};
diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu
index 0b9c70b..3815a1a 100644
--- a/cuda/3d/astra3d.cu
+++ b/cuda/3d/astra3d.cu
@@ -40,6 +40,12 @@ $Id$
#include "arith3d.h"
#include "astra3d.h"
+#include "astra/ParallelProjectionGeometry3D.h"
+#include "astra/ParallelVecProjectionGeometry3D.h"
+#include "astra/ConeProjectionGeometry3D.h"
+#include "astra/ConeVecProjectionGeometry3D.h"
+#include "astra/VolumeGeometry3D.h"
+
#include <iostream>
using namespace astraCUDA3d;
@@ -137,6 +143,202 @@ static SPar3DProjection* genPar3DProjections(unsigned int iProjAngles,
+
+// adjust pProjs to normalize volume geometry
+template<typename ProjectionT>
+static bool convertAstraGeometry_internal(const CVolumeGeometry3D* pVolGeom,
+ unsigned int iProjectionAngleCount,
+ ProjectionT*& pProjs,
+ float& fOutputScale)
+{
+ assert(pVolGeom);
+ assert(pProjs);
+
+ // TODO: Relative instead of absolute
+ const float EPS = 0.00001f;
+ if (abs(pVolGeom->getPixelLengthX() - pVolGeom->getPixelLengthY()) > EPS)
+ return false;
+ if (abs(pVolGeom->getPixelLengthX() - pVolGeom->getPixelLengthZ()) > EPS)
+ return false;
+
+
+ // Translate
+ float dx = -(pVolGeom->getWindowMinX() + pVolGeom->getWindowMaxX()) / 2;
+ float dy = -(pVolGeom->getWindowMinY() + pVolGeom->getWindowMaxY()) / 2;
+ float dz = -(pVolGeom->getWindowMinZ() + pVolGeom->getWindowMaxZ()) / 2;
+
+ float factor = 1.0f / pVolGeom->getPixelLengthX();
+
+ for (int i = 0; i < iProjectionAngleCount; ++i) {
+ // CHECKME: Order of scaling and translation
+ pProjs[i].translate(dx, dy, dz);
+ pProjs[i].scale(factor);
+ }
+
+ // CHECKME: Check factor
+ fOutputScale *= pVolGeom->getPixelLengthX();
+
+ return true;
+}
+
+
+bool convertAstraGeometry_dims(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
+ SDimensions3D& dims)
+{
+ dims.iVolX = pVolGeom->getGridColCount();
+ dims.iVolY = pVolGeom->getGridRowCount();
+ dims.iVolZ = pVolGeom->getGridSliceCount();
+ dims.iProjAngles = pProjGeom->getProjectionCount();
+ dims.iProjU = pProjGeom->getDetectorColCount(),
+ dims.iProjV = pProjGeom->getDetectorRowCount(),
+ dims.iRaysPerDetDim = 1;
+ dims.iRaysPerVoxelDim = 1;
+
+ if (dims.iVolX <= 0 || dims.iVolX <= 0 || dims.iVolX <= 0)
+ return false;
+ if (dims.iProjAngles <= 0 || dims.iProjU <= 0 || dims.iProjV <= 0)
+ return false;
+
+ return true;
+}
+
+
+bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CParallelProjectionGeometry3D* pProjGeom,
+ SPar3DProjection*& pProjs, float& fOutputScale)
+{
+ assert(pVolGeom);
+ assert(pProjGeom);
+ assert(pProjGeom->getProjectionAngles());
+
+ int nth = pProjGeom->getProjectionCount();
+
+ pProjs = genPar3DProjections(nth,
+ pProjGeom->getDetectorColCount(),
+ pProjGeom->getDetectorRowCount(),
+ pProjGeom->getDetectorSpacingX(),
+ pProjGeom->getDetectorSpacingY(),
+ pProjGeom->getProjectionAngles());
+
+ bool ok;
+
+ fOutputScale = 1.0f;
+
+ ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale);
+
+ return ok;
+}
+
+bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CParallelVecProjectionGeometry3D* pProjGeom,
+ SPar3DProjection*& pProjs, float& fOutputScale)
+{
+ assert(pVolGeom);
+ assert(pProjGeom);
+ assert(pProjGeom->getProjectionVectors());
+
+ int nth = pProjGeom->getProjectionCount();
+
+ pProjs = new SPar3DProjection[nth];
+ for (int i = 0; i < nth; ++i)
+ pProjs[i] = pProjGeom->getProjectionVectors()[i];
+
+ bool ok;
+
+ fOutputScale = 1.0f;
+
+ ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale);
+
+ return ok;
+}
+
+bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CConeProjectionGeometry3D* pProjGeom,
+ SConeProjection*& pProjs, float& fOutputScale)
+{
+ assert(pVolGeom);
+ assert(pProjGeom);
+ assert(pProjGeom->getProjectionAngles());
+
+ int nth = pProjGeom->getProjectionCount();
+
+ pProjs = genConeProjections(nth,
+ pProjGeom->getDetectorColCount(),
+ pProjGeom->getDetectorRowCount(),
+ pProjGeom->getOriginSourceDistance(),
+ pProjGeom->getOriginDetectorDistance(),
+ pProjGeom->getDetectorSpacingX(),
+ pProjGeom->getDetectorSpacingY(),
+ pProjGeom->getProjectionAngles());
+
+ bool ok;
+
+ fOutputScale = 1.0f;
+
+ ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale);
+
+ return ok;
+}
+
+bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CConeVecProjectionGeometry3D* pProjGeom,
+ SConeProjection*& pProjs, float& fOutputScale)
+{
+ assert(pVolGeom);
+ assert(pProjGeom);
+ assert(pProjGeom->getProjectionVectors());
+
+ int nth = pProjGeom->getProjectionCount();
+
+ pProjs = new SConeProjection[nth];
+ for (int i = 0; i < nth; ++i)
+ pProjs[i] = pProjGeom->getProjectionVectors()[i];
+
+ bool ok;
+
+ fOutputScale = 1.0f;
+
+ ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale);
+
+ return ok;
+}
+
+
+bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
+ SPar3DProjection*& pParProjs,
+ SConeProjection*& pConeProjs,
+ float& fOutputScale)
+{
+ const CConeProjectionGeometry3D* conegeom = dynamic_cast<const CConeProjectionGeometry3D*>(pProjGeom);
+ const CParallelProjectionGeometry3D* par3dgeom = dynamic_cast<const CParallelProjectionGeometry3D*>(pProjGeom);
+ const CParallelVecProjectionGeometry3D* parvec3dgeom = dynamic_cast<const CParallelVecProjectionGeometry3D*>(pProjGeom);
+ const CConeVecProjectionGeometry3D* conevec3dgeom = dynamic_cast<const CConeVecProjectionGeometry3D*>(pProjGeom);
+
+ pConeProjs = 0;
+ pParProjs = 0;
+
+ bool ok;
+
+ if (conegeom) {
+ ok = convertAstraGeometry(pVolGeom, conegeom, pConeProjs, fOutputScale);
+ } else if (conevec3dgeom) {
+ ok = convertAstraGeometry(pVolGeom, conevec3dgeom, pConeProjs, fOutputScale);
+ } else if (par3dgeom) {
+ ok = convertAstraGeometry(pVolGeom, par3dgeom, pParProjs, fOutputScale);
+ } else if (parvec3dgeom) {
+ ok = convertAstraGeometry(pVolGeom, parvec3dgeom, pParProjs, fOutputScale);
+ } else {
+ ok = false;
+ }
+
+ return ok;
+}
+
+
+
+
class AstraSIRT3d_internal {
public:
SDimensions3D dims;
@@ -151,7 +353,7 @@ public:
SConeProjection* projs;
SPar3DProjection* parprojs;
- float fPixelSize;
+ float fOutputScale;
bool initialized;
bool setStartReconstruction;
@@ -188,6 +390,8 @@ AstraSIRT3d::AstraSIRT3d()
pData->dims.iRaysPerVoxelDim = 1;
pData->projs = 0;
+ pData->parprojs = 0;
+ pData->fOutputScale = 1.0f;
pData->initialized = false;
pData->setStartReconstruction = false;
@@ -220,127 +424,37 @@ AstraSIRT3d::~AstraSIRT3d()
pData = 0;
}
-bool AstraSIRT3d::setReconstructionGeometry(unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ/*,
- float fPixelSize = 1.0f*/)
-{
- if (pData->initialized)
- return false;
-
- pData->dims.iVolX = iVolX;
- pData->dims.iVolY = iVolY;
- pData->dims.iVolZ = iVolZ;
-
- return (iVolX > 0 && iVolY > 0 && iVolZ > 0);
-}
-
-
-bool AstraSIRT3d::setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection* projs)
-{
- if (pData->initialized)
- return false;
-
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || projs == 0)
- return false;
-
- pData->parprojs = new SPar3DProjection[iProjAngles];
- memcpy(pData->parprojs, projs, iProjAngles * sizeof(projs[0]));
-
- pData->projType = PROJ_PARALLEL;
-
- return true;
-}
-
-bool AstraSIRT3d::setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles)
+bool AstraSIRT3d::setGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom)
{
if (pData->initialized)
return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SPar3DProjection* p = genPar3DProjections(iProjAngles,
- iProjU, iProjV,
- fDetUSize, fDetVSize,
- pfAngles);
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- pData->parprojs = p;
- pData->projType = PROJ_PARALLEL;
-
- return true;
-}
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, pData->dims);
-
-
-bool AstraSIRT3d::setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection* projs)
-{
- if (pData->initialized)
+ if (!ok)
return false;
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
+ pData->projs = 0;
+ pData->parprojs = 0;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || projs == 0)
+ ok = convertAstraGeometry(pVolGeom, pProjGeom,
+ pData->parprojs, pData->projs,
+ pData->fOutputScale);
+ if (!ok)
return false;
- pData->projs = new SConeProjection[iProjAngles];
- memcpy(pData->projs, projs, iProjAngles * sizeof(projs[0]));
-
- pData->projType = PROJ_CONE;
+ if (pData->projs) {
+ assert(pData->parprojs == 0);
+ pData->projType = PROJ_CONE;
+ } else {
+ assert(pData->parprojs != 0);
+ pData->projType = PROJ_PARALLEL;
+ }
return true;
}
-bool AstraSIRT3d::setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles)
-{
- if (pData->initialized)
- return false;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SConeProjection* p = genConeProjections(iProjAngles,
- iProjU, iProjV,
- fOriginSourceDistance,
- fOriginDetectorDistance,
- fDetUSize, fDetVSize,
- pfAngles);
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- pData->projs = p;
- pData->projType = PROJ_CONE;
-
- return true;
-}
bool AstraSIRT3d::enableSuperSampling(unsigned int iVoxelSuperSampling,
unsigned int iDetectorSuperSampling)
@@ -404,9 +518,9 @@ bool AstraSIRT3d::init()
bool ok;
if (pData->projType == PROJ_PARALLEL) {
- ok = pData->sirt.setPar3DGeometry(pData->dims, pData->parprojs);
+ ok = pData->sirt.setPar3DGeometry(pData->dims, pData->parprojs, pData->fOutputScale);
} else {
- ok = pData->sirt.setConeGeometry(pData->dims, pData->projs);
+ ok = pData->sirt.setConeGeometry(pData->dims, pData->projs, pData->fOutputScale);
}
if (!ok)
@@ -618,7 +732,7 @@ public:
SConeProjection* projs;
SPar3DProjection* parprojs;
- float fPixelSize;
+ float fOutputScale;
bool initialized;
bool setStartReconstruction;
@@ -655,6 +769,8 @@ AstraCGLS3d::AstraCGLS3d()
pData->dims.iRaysPerVoxelDim = 1;
pData->projs = 0;
+ pData->parprojs = 0;
+ pData->fOutputScale = 1.0f;
pData->initialized = false;
pData->setStartReconstruction = false;
@@ -687,125 +803,33 @@ AstraCGLS3d::~AstraCGLS3d()
pData = 0;
}
-bool AstraCGLS3d::setReconstructionGeometry(unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ/*,
- float fPixelSize = 1.0f*/)
+bool AstraCGLS3d::setGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom)
{
if (pData->initialized)
return false;
- pData->dims.iVolX = iVolX;
- pData->dims.iVolY = iVolY;
- pData->dims.iVolZ = iVolZ;
-
- return (iVolX > 0 && iVolY > 0 && iVolZ > 0);
-}
-
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, pData->dims);
-bool AstraCGLS3d::setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection* projs)
-{
- if (pData->initialized)
- return false;
-
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || projs == 0)
- return false;
-
- pData->parprojs = new SPar3DProjection[iProjAngles];
- memcpy(pData->parprojs, projs, iProjAngles * sizeof(projs[0]));
-
- pData->projType = PROJ_PARALLEL;
-
- return true;
-}
-
-bool AstraCGLS3d::setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles)
-{
- if (pData->initialized)
- return false;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SPar3DProjection* p = genPar3DProjections(iProjAngles,
- iProjU, iProjV,
- fDetUSize, fDetVSize,
- pfAngles);
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- pData->parprojs = p;
- pData->projType = PROJ_PARALLEL;
-
- return true;
-}
-
-
-
-bool AstraCGLS3d::setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection* projs)
-{
- if (pData->initialized)
- return false;
-
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || projs == 0)
+ if (!ok)
return false;
- pData->projs = new SConeProjection[iProjAngles];
- memcpy(pData->projs, projs, iProjAngles * sizeof(projs[0]));
-
- pData->projType = PROJ_CONE;
-
- return true;
-}
-
-bool AstraCGLS3d::setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles)
-{
- if (pData->initialized)
- return false;
+ pData->projs = 0;
+ pData->parprojs = 0;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
+ ok = convertAstraGeometry(pVolGeom, pProjGeom,
+ pData->parprojs, pData->projs,
+ pData->fOutputScale);
+ if (!ok)
return false;
- SConeProjection* p = genConeProjections(iProjAngles,
- iProjU, iProjV,
- fOriginSourceDistance,
- fOriginDetectorDistance,
- fDetUSize, fDetVSize,
- pfAngles);
-
- pData->dims.iProjAngles = iProjAngles;
- pData->dims.iProjU = iProjU;
- pData->dims.iProjV = iProjV;
-
- pData->projs = p;
- pData->projType = PROJ_CONE;
+ if (pData->projs) {
+ assert(pData->parprojs == 0);
+ pData->projType = PROJ_CONE;
+ } else {
+ assert(pData->parprojs != 0);
+ pData->projType = PROJ_PARALLEL;
+ }
return true;
}
@@ -874,9 +898,9 @@ bool AstraCGLS3d::init()
bool ok;
if (pData->projType == PROJ_PARALLEL) {
- ok = pData->cgls.setPar3DGeometry(pData->dims, pData->parprojs);
+ ok = pData->cgls.setPar3DGeometry(pData->dims, pData->parprojs, pData->fOutputScale);
} else {
- ok = pData->cgls.setConeGeometry(pData->dims, pData->projs);
+ ok = pData->cgls.setConeGeometry(pData->dims, pData->projs, pData->fOutputScale);
}
if (!ok)
@@ -1077,179 +1101,31 @@ float AstraCGLS3d::computeDiffNorm()
-bool astraCudaConeFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling)
-{
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SConeProjection* p = genConeProjections(iProjAngles,
- iProjU, iProjV,
- fOriginSourceDistance,
- fOriginDetectorDistance,
- fDetUSize, fDetVSize,
- pfAngles);
-
- bool ok;
- ok = astraCudaConeFP(pfVolume, pfProjections, iVolX, iVolY, iVolZ,
- iProjAngles, iProjU, iProjV, p, iGPUIndex, iDetectorSuperSampling);
-
- delete[] p;
-
- return ok;
-}
-
-bool astraCudaConeFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling)
+bool astraCudaFP(const float* pfVolume, float* pfProjections,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
+ int iGPUIndex, int iDetectorSuperSampling,
+ Cuda3DProjectionKernel projKernel)
{
SDimensions3D dims;
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
-
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims);
+ if (!ok)
return false;
dims.iRaysPerDetDim = iDetectorSuperSampling;
-
if (iDetectorSuperSampling == 0)
return false;
- if (iGPUIndex != -1) {
- cudaSetDevice(iGPUIndex);
- cudaError_t err = cudaGetLastError();
+ SPar3DProjection* pParProjs;
+ SConeProjection* pConeProjs;
- // Ignore errors caused by calling cudaSetDevice multiple times
- if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess)
- return false;
- }
+ float outputScale;
- cudaPitchedPtr D_volumeData = allocateVolumeData(dims);
- bool ok = D_volumeData.ptr;
- if (!ok)
- return false;
+ ok = convertAstraGeometry(pVolGeom, pProjGeom,
+ pParProjs, pConeProjs,
+ outputScale);
- cudaPitchedPtr D_projData = allocateProjectionData(dims);
- ok = D_projData.ptr;
- if (!ok) {
- cudaFree(D_volumeData.ptr);
- return false;
- }
-
- ok &= copyVolumeToDevice(pfVolume, D_volumeData, dims, dims.iVolX);
-
- ok &= zeroProjectionData(D_projData, dims);
-
- if (!ok) {
- cudaFree(D_volumeData.ptr);
- cudaFree(D_projData.ptr);
- return false;
- }
-
- ok &= ConeFP(D_volumeData, D_projData, dims, pfAngles, 1.0f);
-
- ok &= copyProjectionsFromDevice(pfProjections, D_projData,
- dims, dims.iProjU);
-
-
- cudaFree(D_volumeData.ptr);
- cudaFree(D_projData.ptr);
-
- return ok;
-
-}
-
-bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling,
- Cuda3DProjectionKernel projKernel)
-{
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SPar3DProjection* p = genPar3DProjections(iProjAngles,
- iProjU, iProjV,
- fDetUSize, fDetVSize,
- pfAngles);
-
- bool ok;
- ok = astraCudaPar3DFP(pfVolume, pfProjections, iVolX, iVolY, iVolZ,
- iProjAngles, iProjU, iProjV, p, iGPUIndex, iDetectorSuperSampling,
- projKernel);
-
- delete[] p;
-
- return ok;
-}
-
-
-bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling,
- Cuda3DProjectionKernel projKernel)
-{
- SDimensions3D dims;
-
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
-
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- dims.iRaysPerDetDim = iDetectorSuperSampling;
-
- if (iDetectorSuperSampling == 0)
- return false;
if (iGPUIndex != -1) {
cudaSetDevice(iGPUIndex);
@@ -1262,7 +1138,7 @@ bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
cudaPitchedPtr D_volumeData = allocateVolumeData(dims);
- bool ok = D_volumeData.ptr;
+ ok = D_volumeData.ptr;
if (!ok)
return false;
@@ -1283,15 +1159,25 @@ bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
return false;
}
- switch (projKernel) {
- case ker3d_default:
- ok &= Par3DFP(D_volumeData, D_projData, dims, pfAngles, 1.0f);
- break;
- case ker3d_sum_square_weights:
- ok &= Par3DFP_SumSqW(D_volumeData, D_projData, dims, pfAngles, 1.0f);
- break;
- default:
- assert(false);
+ if (pParProjs) {
+ switch (projKernel) {
+ case ker3d_default:
+ ok &= Par3DFP(D_volumeData, D_projData, dims, pParProjs, outputScale);
+ break;
+ case ker3d_sum_square_weights:
+ ok &= Par3DFP_SumSqW(D_volumeData, D_projData, dims, pParProjs, outputScale*outputScale);
+ break;
+ default:
+ assert(false);
+ }
+ } else {
+ switch (projKernel) {
+ case ker3d_default:
+ ok &= ConeFP(D_volumeData, D_projData, dims, pConeProjs, outputScale);
+ break;
+ default:
+ assert(false);
+ }
}
ok &= copyProjectionsFromDevice(pfProjections, D_projData,
@@ -1305,207 +1191,28 @@ bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
}
-bool astraCudaConeBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling)
-{
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SConeProjection* p = genConeProjections(iProjAngles,
- iProjU, iProjV,
- fOriginSourceDistance,
- fOriginDetectorDistance,
- fDetUSize, fDetVSize,
- pfAngles);
-
- bool ok;
- ok = astraCudaConeBP(pfVolume, pfProjections, iVolX, iVolY, iVolZ,
- iProjAngles, iProjU, iProjV, p, iGPUIndex, iVoxelSuperSampling);
- delete[] p;
-
- return ok;
-}
-
-bool astraCudaConeBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling)
+bool astraCudaBP(float* pfVolume, const float* pfProjections,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
+ int iGPUIndex, int iVoxelSuperSampling)
{
SDimensions3D dims;
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
-
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- dims.iRaysPerVoxelDim = iVoxelSuperSampling;
-
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- if (iGPUIndex != -1) {
- cudaSetDevice(iGPUIndex);
- cudaError_t err = cudaGetLastError();
-
- // Ignore errors caused by calling cudaSetDevice multiple times
- if (err != cudaSuccess && err != cudaErrorSetOnActiveProcess)
- return false;
- }
-
- cudaPitchedPtr D_volumeData = allocateVolumeData(dims);
- bool ok = D_volumeData.ptr;
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims);
if (!ok)
return false;
- cudaPitchedPtr D_projData = allocateProjectionData(dims);
- ok = D_projData.ptr;
- if (!ok) {
- cudaFree(D_volumeData.ptr);
- return false;
- }
-
- ok &= copyProjectionsToDevice(pfProjections, D_projData,
- dims, dims.iProjU);
-
- ok &= zeroVolumeData(D_volumeData, dims);
-
- if (!ok) {
- cudaFree(D_volumeData.ptr);
- cudaFree(D_projData.ptr);
- return false;
- }
-
- ok &= ConeBP(D_volumeData, D_projData, dims, pfAngles);
-
- ok &= copyVolumeFromDevice(pfVolume, D_volumeData, dims, dims.iVolX);
-
-
- cudaFree(D_volumeData.ptr);
- cudaFree(D_projData.ptr);
-
- return ok;
-
-}
-
-bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling)
-{
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SPar3DProjection* p = genPar3DProjections(iProjAngles,
- iProjU, iProjV,
- fDetUSize, fDetVSize,
- pfAngles);
-
- bool ok;
- ok = astraCudaPar3DBP(pfVolume, pfProjections, iVolX, iVolY, iVolZ,
- iProjAngles, iProjU, iProjV, p, iGPUIndex, iVoxelSuperSampling);
-
- delete[] p;
-
- return ok;
-}
-
-// This computes the column weights, divides by them, and adds the
-// result to the current volume. This is both more expensive and more
-// GPU memory intensive than the regular BP, but allows saving system RAM.
-bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling)
-{
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
-
- SPar3DProjection* p = genPar3DProjections(iProjAngles,
- iProjU, iProjV,
- fDetUSize, fDetVSize,
- pfAngles);
-
- bool ok;
- ok = astraCudaPar3DBP_SIRTWeighted(pfVolume, pfProjections, iVolX, iVolY, iVolZ,
- iProjAngles, iProjU, iProjV, p, iGPUIndex, iVoxelSuperSampling);
-
- delete[] p;
-
- return ok;
-}
-
-
-bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling)
-{
- SDimensions3D dims;
-
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
+ dims.iRaysPerVoxelDim = iVoxelSuperSampling;
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
+ SPar3DProjection* pParProjs;
+ SConeProjection* pConeProjs;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
+ float outputScale;
- dims.iRaysPerVoxelDim = iVoxelSuperSampling;
+ ok = convertAstraGeometry(pVolGeom, pProjGeom,
+ pParProjs, pConeProjs,
+ outputScale);
if (iGPUIndex != -1) {
cudaSetDevice(iGPUIndex);
@@ -1518,7 +1225,7 @@ bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
cudaPitchedPtr D_volumeData = allocateVolumeData(dims);
- bool ok = D_volumeData.ptr;
+ ok = D_volumeData.ptr;
if (!ok)
return false;
@@ -1540,7 +1247,10 @@ bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
return false;
}
- ok &= Par3DBP(D_volumeData, D_projData, dims, pfAngles);
+ if (pParProjs)
+ ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, outputScale);
+ else
+ ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, outputScale);
ok &= copyVolumeFromDevice(pfVolume, D_volumeData, dims, dims.iVolX);
@@ -1556,33 +1266,28 @@ bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
// This computes the column weights, divides by them, and adds the
// result to the current volume. This is both more expensive and more
// GPU memory intensive than the regular BP, but allows saving system RAM.
-bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
+bool astraCudaBP_SIRTWeighted(float* pfVolume,
const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
int iGPUIndex, int iVoxelSuperSampling)
{
SDimensions3D dims;
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims);
+ if (!ok)
return false;
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
+ dims.iRaysPerVoxelDim = iVoxelSuperSampling;
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
- return false;
+ SPar3DProjection* pParProjs;
+ SConeProjection* pConeProjs;
- dims.iRaysPerVoxelDim = iVoxelSuperSampling;
+ float outputScale;
+
+ ok = convertAstraGeometry(pVolGeom, pProjGeom,
+ pParProjs, pConeProjs,
+ outputScale);
if (iGPUIndex != -1) {
cudaSetDevice(iGPUIndex);
@@ -1595,7 +1300,7 @@ bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
cudaPitchedPtr D_pixelWeight = allocateVolumeData(dims);
- bool ok = D_pixelWeight.ptr;
+ ok = D_pixelWeight.ptr;
if (!ok)
return false;
@@ -1617,7 +1322,12 @@ bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
// Compute weights
ok &= zeroVolumeData(D_pixelWeight, dims);
processSino3D<opSet>(D_projData, 1.0f, dims);
- ok &= Par3DBP(D_pixelWeight, D_projData, dims, pfAngles);
+
+ if (pParProjs)
+ ok &= Par3DBP(D_pixelWeight, D_projData, dims, pParProjs, outputScale);
+ else
+ ok &= ConeBP(D_pixelWeight, D_projData, dims, pConeProjs, outputScale);
+
processVol3D<opInvert>(D_pixelWeight, dims);
if (!ok) {
cudaFree(D_pixelWeight.ptr);
@@ -1630,7 +1340,11 @@ bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
dims, dims.iProjU);
ok &= zeroVolumeData(D_volumeData, dims);
// Do BP into D_volumeData
- ok &= Par3DBP(D_volumeData, D_projData, dims, pfAngles);
+ if (pParProjs)
+ ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, outputScale);
+ else
+ ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, outputScale);
+
// Multiply with weights
processVol3D<opMul>(D_volumeData, D_pixelWeight, dims);
@@ -1653,6 +1367,9 @@ bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
cudaFree(D_volumeData.ptr);
cudaFree(D_projData.ptr);
+ delete[] pParProjs;
+ delete[] pConeProjs;
+
return ok;
}
@@ -1660,33 +1377,19 @@ bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume,
bool astraCudaFDK(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
+ const CVolumeGeometry3D* pVolGeom,
+ const CConeProjectionGeometry3D* pProjGeom,
bool bShortScan,
int iGPUIndex, int iVoxelSuperSampling)
{
SDimensions3D dims;
- dims.iVolX = iVolX;
- dims.iVolY = iVolY;
- dims.iVolZ = iVolZ;
- if (iVolX == 0 || iVolY == 0 || iVolZ == 0)
- return false;
+ bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims);
- dims.iProjAngles = iProjAngles;
- dims.iProjU = iProjU;
- dims.iProjV = iProjV;
+ // TODO: Check that pVolGeom is normalized, since we don't support
+ // other volume geometries yet
- if (iProjAngles == 0 || iProjU == 0 || iProjV == 0 || pfAngles == 0)
+ if (!ok)
return false;
dims.iRaysPerVoxelDim = iVoxelSuperSampling;
@@ -1703,9 +1406,8 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections,
return false;
}
-
cudaPitchedPtr D_volumeData = allocateVolumeData(dims);
- bool ok = D_volumeData.ptr;
+ ok = D_volumeData.ptr;
if (!ok)
return false;
@@ -1726,6 +1428,13 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections,
return false;
}
+ float fOriginSourceDistance = pProjGeom->getOriginSourceDistance();
+ float fOriginDetectorDistance = pProjGeom->getOriginDetectorDistance();
+ float fDetUSize = pProjGeom->getDetectorSpacingX();
+ float fDetVSize = pProjGeom->getDetectorSpacingY();
+ const float *pfAngles = pProjGeom->getProjectionAngles();
+
+
// TODO: Offer interface for SrcZ, DetZ
ok &= FDK(D_volumeData, D_projData, fOriginSourceDistance,
fOriginDetectorDistance, 0, 0, fDetUSize, fDetVSize,
diff --git a/cuda/3d/astra3d.h b/cuda/3d/astra3d.h
index f91fe26..6c3fcfb 100644
--- a/cuda/3d/astra3d.h
+++ b/cuda/3d/astra3d.h
@@ -42,7 +42,12 @@ enum Cuda3DProjectionKernel {
ker3d_sum_square_weights
};
-
+class CProjectionGeometry3D;
+class CParallelProjectionGeometry3D;
+class CParallelVecProjectionGeometry3D;
+class CConeProjectionGeometry3D;
+class CConeVecProjectionGeometry3D;
+class CVolumeGeometry3D;
class AstraSIRT3d_internal;
@@ -52,37 +57,9 @@ public:
AstraSIRT3d();
~AstraSIRT3d();
- // Set the number of pixels in the reconstruction rectangle,
- // and the length of the edge of a pixel.
- // Volume pixels are assumed to be square.
- // This must be called before setting the projection geometry.
- bool setReconstructionGeometry(unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ/*,
- float fPixelSize = 1.0f*/);
-
- bool setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection* projs);
- bool setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fSourceZ,
- float fDetSize,
- const float *pfAngles);
- bool setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection* projs);
- bool setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fSourceZ,
- float fDetSize,
- const float *pfAngles);
+ // Set the volume and projection geometry
+ bool setGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom);
// Enable supersampling.
//
@@ -197,37 +174,9 @@ public:
AstraCGLS3d();
~AstraCGLS3d();
- // Set the number of pixels in the reconstruction rectangle,
- // and the length of the edge of a pixel.
- // Volume pixels are assumed to be square.
- // This must be called before setting the projection geometry.
- bool setReconstructionGeometry(unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ/*,
- float fPixelSize = 1.0f*/);
-
- bool setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection* projs);
- bool setConeGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fSourceZ,
- float fDetSize,
- const float *pfAngles);
- bool setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection* projs);
- bool setPar3DGeometry(unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fSourceZ,
- float fDetSize,
- const float *pfAngles);
+ // Set the volume and projection geometry
+ bool setGeometry(const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom);
// Enable supersampling.
//
@@ -333,139 +282,30 @@ protected:
};
-
-_AstraExport bool astraCudaConeFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling);
-
-_AstraExport bool astraCudaConeFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling);
-
-_AstraExport bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iDetectorSuperSampling,
- Cuda3DProjectionKernel projKernel);
-
-_AstraExport bool astraCudaPar3DFP(const float* pfVolume, float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
+_AstraExport bool astraCudaFP(const float* pfVolume, float* pfProjections,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
int iGPUIndex, int iDetectorSuperSampling,
Cuda3DProjectionKernel projKernel);
-_AstraExport bool astraCudaConeBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling);
-
-_AstraExport bool astraCudaConeBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SConeProjection *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling);
-
-_AstraExport bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
+_AstraExport bool astraCudaBP(float* pfVolume, const float* pfProjections,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
int iGPUIndex, int iVoxelSuperSampling);
-_AstraExport bool astraCudaPar3DBP(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling);
-
-_AstraExport bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
- int iGPUIndex, int iVoxelSuperSampling);
-
-_AstraExport bool astraCudaPar3DBP_SIRTWeighted(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- const SPar3DProjection *pfAngles,
+_AstraExport bool astraCudaBP_SIRTWeighted(float* pfVolume, const float* pfProjections,
+ const CVolumeGeometry3D* pVolGeom,
+ const CProjectionGeometry3D* pProjGeom,
int iGPUIndex, int iVoxelSuperSampling);
_AstraExport bool astraCudaFDK(float* pfVolume, const float* pfProjections,
- unsigned int iVolX,
- unsigned int iVolY,
- unsigned int iVolZ,
- unsigned int iProjAngles,
- unsigned int iProjU,
- unsigned int iProjV,
- float fOriginSourceDistance,
- float fOriginDetectorDistance,
- float fDetUSize,
- float fDetVSize,
- const float *pfAngles,
+ const CVolumeGeometry3D* pVolGeom,
+ const CConeProjectionGeometry3D* pProjGeom,
bool bShortScan,
int iGPUIndex, int iVoxelSuperSampling);
+
}
diff --git a/cuda/3d/cgls3d.cu b/cuda/3d/cgls3d.cu
index 5071a9b..dd0e8a0 100644
--- a/cuda/3d/cgls3d.cu
+++ b/cuda/3d/cgls3d.cu
@@ -165,7 +165,7 @@ bool CGLS::iterate(unsigned int iterations)
// p = A'*r
zeroVolumeData(D_p, dims);
- callBP(D_p, D_r);
+ callBP(D_p, D_r, 1.0f);
if (useVolumeMask)
processVol3D<opMul>(D_p, D_maskData, dims);
@@ -195,7 +195,7 @@ bool CGLS::iterate(unsigned int iterations)
// z = A'*r
zeroVolumeData(D_z, dims);
- callBP(D_z, D_r);
+ callBP(D_z, D_r, 1.0f);
if (useVolumeMask)
processVol3D<opMul>(D_z, D_maskData, dims);
@@ -242,7 +242,7 @@ bool doCGLS(cudaPitchedPtr& D_volumeData,
CGLS cgls;
bool ok = true;
- ok &= cgls.setConeGeometry(dims, angles);
+ ok &= cgls.setConeGeometry(dims, angles, 1.0f);
if (D_maskData.ptr)
ok &= cgls.enableVolumeMask();
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index 5648d6f..4a41f6a 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -78,7 +78,8 @@ bool bindProjDataTexture(const cudaArray* array)
//__launch_bounds__(32*16, 4)
__global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAngle,
- int angleOffset, const astraCUDA3d::SDimensions3D dims)
+ int angleOffset, const astraCUDA3d::SDimensions3D dims,
+ float fOutputScale)
{
float* volData = (float*)D_volData;
@@ -147,13 +148,13 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng
endZ = dims.iVolZ - startZ;
for(int i=0; i < endZ; i++)
- volData[((startZ+i)*dims.iVolY+Y)*volPitch+X] += Z[i];
+ volData[((startZ+i)*dims.iVolY+Y)*volPitch+X] += Z[i] * fOutputScale;
} //End kernel
// supersampling version
-__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims)
+__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
{
float* volData = (float*)D_volData;
@@ -189,6 +190,9 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start
float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim;
const float fSubStep = 1.0f/dims.iRaysPerVoxelDim;
+ fOutputScale /= (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim);
+
+
for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f)
{
@@ -236,14 +240,15 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start
}
- volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal / (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim);
+ volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal * fOutputScale;
}
}
bool ConeBP_Array(cudaPitchedPtr D_volumeData,
cudaArray *D_projArray,
- const SDimensions3D& dims, const SConeProjection* angles)
+ const SDimensions3D& dims, const SConeProjection* angles,
+ float fOutputScale)
{
bindProjDataTexture(D_projArray);
@@ -291,9 +296,9 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
for (unsigned int i = 0; i < angleCount; i += g_anglesPerBlock) {
// printf("Calling BP: %d, %dx%d, %dx%d to %p\n", i, dimBlock.x, dimBlock.y, dimGrid.x, dimGrid.y, (void*)D_volumeData.ptr);
if (dims.iRaysPerVoxelDim == 1)
- dev_cone_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims);
+ dev_cone_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
else
- dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims);
+ dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
}
cudaTextForceKernelsCompletion();
@@ -309,14 +314,15 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
bool ConeBP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
- const SDimensions3D& dims, const SConeProjection* angles)
+ const SDimensions3D& dims, const SConeProjection* angles,
+ float fOutputScale)
{
// transfer projections to array
cudaArray* cuArray = allocateProjectionArray(dims);
transferProjectionsToArray(D_projData, cuArray, dims);
- bool ret = ConeBP_Array(D_volumeData, cuArray, dims, angles);
+ bool ret = ConeBP_Array(D_volumeData, cuArray, dims, angles, fOutputScale);
cudaFreeArray(cuArray);
@@ -473,7 +479,7 @@ int main()
}
#endif
- astraCUDA3d::ConeBP(volData, projData, dims, angle);
+ astraCUDA3d::ConeBP(volData, projData, dims, angle, 1.0f);
#if 0
float* buf = new float[256*256];
diff --git a/cuda/3d/cone_bp.h b/cuda/3d/cone_bp.h
index cba6d9f..4d3d2dd 100644
--- a/cuda/3d/cone_bp.h
+++ b/cuda/3d/cone_bp.h
@@ -33,13 +33,14 @@ namespace astraCUDA3d {
_AstraExport bool ConeBP_Array(cudaPitchedPtr D_volumeData,
cudaArray *D_projArray,
- const SDimensions3D& dims, const SConeProjection* angles);
+ const SDimensions3D& dims, const SConeProjection* angles,
+ float fOutputScale);
_AstraExport bool ConeBP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
- const SDimensions3D& dims, const SConeProjection* angles);
+ const SDimensions3D& dims, const SConeProjection* angles,
+ float fOutputScale);
-
}
#endif
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu
index 0c33280..cafab46 100644
--- a/cuda/3d/par3d_bp.cu
+++ b/cuda/3d/par3d_bp.cu
@@ -77,7 +77,7 @@ static bool bindProjDataTexture(const cudaArray* array)
}
-__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims)
+__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
{
float* volData = (float*)D_volData;
@@ -139,11 +139,11 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn
endZ = dims.iVolZ - startZ;
for(int i=0; i < endZ; i++)
- volData[((startZ+i)*dims.iVolY+Y)*volPitch+X] += Z[i];
+ volData[((startZ+i)*dims.iVolY+Y)*volPitch+X] += Z[i] * fOutputScale;
}
// supersampling version
-__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims)
+__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
{
float* volData = (float*)D_volData;
@@ -180,6 +180,9 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star
const float fSubStep = 1.0f/dims.iRaysPerVoxelDim;
+ fOutputScale /= (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim);
+
+
for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f)
{
@@ -217,14 +220,15 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star
}
- volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal / (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim);
+ volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal * fOutputScale;
}
}
bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
cudaArray *D_projArray,
- const SDimensions3D& dims, const SPar3DProjection* angles)
+ const SDimensions3D& dims, const SPar3DProjection* angles,
+ float fOutputScale)
{
bindProjDataTexture(D_projArray);
@@ -271,9 +275,9 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
for (unsigned int i = 0; i < angleCount; i += g_anglesPerBlock) {
// printf("Calling BP: %d, %dx%d, %dx%d to %p\n", i, dimBlock.x, dimBlock.y, dimGrid.x, dimGrid.y, (void*)D_volumeData.ptr);
if (dims.iRaysPerVoxelDim == 1)
- dev_par3D_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims);
+ dev_par3D_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
else
- dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims);
+ dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
}
cudaTextForceKernelsCompletion();
@@ -288,14 +292,15 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
bool Par3DBP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
- const SDimensions3D& dims, const SPar3DProjection* angles)
+ const SDimensions3D& dims, const SPar3DProjection* angles,
+ float fOutputScale)
{
// transfer projections to array
cudaArray* cuArray = allocateProjectionArray(dims);
transferProjectionsToArray(D_projData, cuArray, dims);
- bool ret = Par3DBP_Array(D_volumeData, cuArray, dims, angles);
+ bool ret = Par3DBP_Array(D_volumeData, cuArray, dims, angles, fOutputScale);
cudaFreeArray(cuArray);
@@ -445,7 +450,7 @@ int main()
cudaMemcpy3D(&p);
}
- astraCUDA3d::Par3DBP(volData, projData, dims, angle);
+ astraCUDA3d::Par3DBP(volData, projData, dims, angle, 1.0f);
#if 1
float* buf = new float[256*256];
diff --git a/cuda/3d/par3d_bp.h b/cuda/3d/par3d_bp.h
index ece37d1..f1fc62d 100644
--- a/cuda/3d/par3d_bp.h
+++ b/cuda/3d/par3d_bp.h
@@ -33,11 +33,13 @@ namespace astraCUDA3d {
_AstraExport bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
cudaArray *D_projArray,
- const SDimensions3D& dims, const SPar3DProjection* angles);
+ const SDimensions3D& dims, const SPar3DProjection* angles,
+ float fOutputScale);
_AstraExport bool Par3DBP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
- const SDimensions3D& dims, const SPar3DProjection* angles);
+ const SDimensions3D& dims, const SPar3DProjection* angles,
+ float fOutputScale);
}
diff --git a/cuda/3d/sirt3d.cu b/cuda/3d/sirt3d.cu
index 389ee6b..484521e 100644
--- a/cuda/3d/sirt3d.cu
+++ b/cuda/3d/sirt3d.cu
@@ -160,10 +160,10 @@ bool SIRT::precomputeWeights()
zeroVolumeData(D_pixelWeight, dims);
if (useSinogramMask) {
- callBP(D_pixelWeight, D_smaskData);
+ callBP(D_pixelWeight, D_smaskData, 1.0f);
} else {
processSino3D<opSet>(D_projData, 1.0f, dims);
- callBP(D_pixelWeight, D_projData);
+ callBP(D_pixelWeight, D_projData, 1.0f);
}
#if 0
float* bufp = new float[512*512];
@@ -293,7 +293,7 @@ bool SIRT::iterate(unsigned int iterations)
#endif
- callBP(D_tmpData, D_projData);
+ callBP(D_tmpData, D_projData, 1.0f);
#if 0
printf("Dumping tmpData: %p\n", (void*)D_tmpData.ptr);
float* buf = new float[256*256];
@@ -347,7 +347,7 @@ bool doSIRT(cudaPitchedPtr& D_volumeData,
SIRT sirt;
bool ok = true;
- ok &= sirt.setConeGeometry(dims, angles);
+ ok &= sirt.setConeGeometry(dims, angles, 1.0f);
if (D_maskData.ptr)
ok &= sirt.enableVolumeMask();
generated by cgit v1.2.3 (git 2.25.1) at 2025-11-04 15:07:53 +0000