diff options
author | Willem Jan Palenstijn <wjp@usecode.org> | 2016-10-07 16:25:52 +0200 |
---|---|---|
committer | GitHub <noreply@github.com> | 2016-10-07 16:25:52 +0200 |
commit | e4b8b6e94be7c5f7dbaad51543c5eace8882a115 (patch) | |
tree | 435152709bffadcac0ce6dba4966fa683e97164b /cuda | |
parent | 7bb42ddd9e26fc7c01734d26bc114b5a935d9110 (diff) | |
parent | e38a4dc774d3f7ca78cec7f16710afd583709b10 (diff) | |
download | astra-e4b8b6e94be7c5f7dbaad51543c5eace8882a115.tar.gz astra-e4b8b6e94be7c5f7dbaad51543c5eace8882a115.tar.bz2 astra-e4b8b6e94be7c5f7dbaad51543c5eace8882a115.tar.xz astra-e4b8b6e94be7c5f7dbaad51543c5eace8882a115.zip |
Merge pull request #41 from wjp/aniso
Add support for non-cube voxels
Diffstat (limited to 'cuda')
-rw-r--r-- | cuda/2d/astra.cu | 1 | ||||
-rw-r--r-- | cuda/3d/algo3d.cu | 21 | ||||
-rw-r--r-- | cuda/3d/algo3d.h | 5 | ||||
-rw-r--r-- | cuda/3d/astra3d.cu | 141 | ||||
-rw-r--r-- | cuda/3d/astra3d.h | 11 | ||||
-rw-r--r-- | cuda/3d/cgls3d.cu | 2 | ||||
-rw-r--r-- | cuda/3d/cone_bp.cu | 30 | ||||
-rw-r--r-- | cuda/3d/cone_bp.h | 4 | ||||
-rw-r--r-- | cuda/3d/cone_fp.cu | 95 | ||||
-rw-r--r-- | cuda/3d/cone_fp.h | 4 | ||||
-rw-r--r-- | cuda/3d/dims3d.h | 22 | ||||
-rw-r--r-- | cuda/3d/mem3d.cu | 27 | ||||
-rw-r--r-- | cuda/3d/par3d_bp.cu | 30 | ||||
-rw-r--r-- | cuda/3d/par3d_bp.h | 4 | ||||
-rw-r--r-- | cuda/3d/par3d_fp.cu | 156 | ||||
-rw-r--r-- | cuda/3d/par3d_fp.h | 6 | ||||
-rw-r--r-- | cuda/3d/sirt3d.cu | 2 |
17 files changed, 349 insertions, 212 deletions
diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu index 4c69628..7317d69 100644 --- a/cuda/2d/astra.cu +++ b/cuda/2d/astra.cu @@ -341,7 +341,6 @@ bool AstraFBP::run() 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, diff --git a/cuda/3d/algo3d.cu b/cuda/3d/algo3d.cu index cc86b70..e2c1e43 100644 --- a/cuda/3d/algo3d.cu +++ b/cuda/3d/algo3d.cu @@ -41,7 +41,6 @@ ReconAlgo3D::ReconAlgo3D() coneProjs = 0; par3DProjs = 0; shouldAbort = false; - fOutputScale = 1.0f; } ReconAlgo3D::~ReconAlgo3D() @@ -58,10 +57,10 @@ void ReconAlgo3D::reset() shouldAbort = false; } -bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProjection* _angles, float _outputScale) +bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProjection* _angles, const SProjectorParams3D& _params) { dims = _dims; - fOutputScale = _outputScale; + params = _params; coneProjs = new SConeProjection[dims.iProjAngles]; par3DProjs = 0; @@ -71,10 +70,10 @@ bool ReconAlgo3D::setConeGeometry(const SDimensions3D& _dims, const SConeProject return true; } -bool ReconAlgo3D::setPar3DGeometry(const SDimensions3D& _dims, const SPar3DProjection* _angles, float _outputScale) +bool ReconAlgo3D::setPar3DGeometry(const SDimensions3D& _dims, const SPar3DProjection* _angles, const SProjectorParams3D& _params) { dims = _dims; - fOutputScale = _outputScale; + params = _params; par3DProjs = new SPar3DProjection[dims.iProjAngles]; coneProjs = 0; @@ -89,10 +88,12 @@ bool ReconAlgo3D::callFP(cudaPitchedPtr& D_volumeData, cudaPitchedPtr& D_projData, float outputScale) { + SProjectorParams3D p = params; + p.fOutputScale *= outputScale; if (coneProjs) { - return ConeFP(D_volumeData, D_projData, dims, coneProjs, outputScale * this->fOutputScale); + return ConeFP(D_volumeData, D_projData, dims, coneProjs, p); } else { - return Par3DFP(D_volumeData, D_projData, dims, par3DProjs, outputScale * this->fOutputScale); + return Par3DFP(D_volumeData, D_projData, dims, par3DProjs, p); } } @@ -100,10 +101,12 @@ bool ReconAlgo3D::callBP(cudaPitchedPtr& D_volumeData, cudaPitchedPtr& D_projData, float outputScale) { + SProjectorParams3D p = params; + p.fOutputScale *= outputScale; if (coneProjs) { - return ConeBP(D_volumeData, D_projData, dims, coneProjs, outputScale * this->fOutputScale); + return ConeBP(D_volumeData, D_projData, dims, coneProjs, p); } else { - return Par3DBP(D_volumeData, D_projData, dims, par3DProjs, outputScale * this->fOutputScale); + return Par3DBP(D_volumeData, D_projData, dims, par3DProjs, p); } } diff --git a/cuda/3d/algo3d.h b/cuda/3d/algo3d.h index 886b092..ce331ad 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, float fOutputScale); - bool setPar3DGeometry(const SDimensions3D& dims, const SPar3DProjection* projs, float fOutputScale); + bool setConeGeometry(const SDimensions3D& dims, const SConeProjection* projs, const SProjectorParams3D& params); + bool setPar3DGeometry(const SDimensions3D& dims, const SPar3DProjection* projs, const SProjectorParams3D& params); void signalAbort() { shouldAbort = true; } @@ -55,6 +55,7 @@ protected: float outputScale); SDimensions3D dims; + SProjectorParams3D params; SConeProjection* coneProjs; SPar3DProjection* par3DProjs; diff --git a/cuda/3d/astra3d.cu b/cuda/3d/astra3d.cu index 5670873..dd6d551 100644 --- a/cuda/3d/astra3d.cu +++ b/cuda/3d/astra3d.cu @@ -67,34 +67,41 @@ template<typename ProjectionT> static bool convertAstraGeometry_internal(const CVolumeGeometry3D* pVolGeom, unsigned int iProjectionAngleCount, ProjectionT*& pProjs, - float& fOutputScale) + SProjectorParams3D& params) { assert(pVolGeom); assert(pProjs); +#if 0 // 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; - +#endif // 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(); + float fx = 1.0f / pVolGeom->getPixelLengthX(); + float fy = 1.0f / pVolGeom->getPixelLengthY(); + float fz = 1.0f / pVolGeom->getPixelLengthZ(); for (int i = 0; i < iProjectionAngleCount; ++i) { // CHECKME: Order of scaling and translation pProjs[i].translate(dx, dy, dz); - pProjs[i].scale(factor); + pProjs[i].scale(fx, fy, fz); } + params.fVolScaleX = pVolGeom->getPixelLengthX(); + params.fVolScaleY = pVolGeom->getPixelLengthY(); + params.fVolScaleZ = pVolGeom->getPixelLengthZ(); + // CHECKME: Check factor - fOutputScale *= pVolGeom->getPixelLengthX(); + //params.fOutputScale *= pVolGeom->getPixelLengthX(); return true; } @@ -108,10 +115,8 @@ bool convertAstraGeometry_dims(const CVolumeGeometry3D* pVolGeom, 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; + dims.iProjU = pProjGeom->getDetectorColCount(); + dims.iProjV = pProjGeom->getDetectorRowCount(); if (dims.iVolX <= 0 || dims.iVolX <= 0 || dims.iVolX <= 0) return false; @@ -124,7 +129,7 @@ bool convertAstraGeometry_dims(const CVolumeGeometry3D* pVolGeom, bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CParallelProjectionGeometry3D* pProjGeom, - SPar3DProjection*& pProjs, float& fOutputScale) + SPar3DProjection*& pProjs, SProjectorParams3D& params) { assert(pVolGeom); assert(pProjGeom); @@ -141,16 +146,14 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, bool ok; - fOutputScale = 1.0f; - - ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale); + ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, params); return ok; } bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CParallelVecProjectionGeometry3D* pProjGeom, - SPar3DProjection*& pProjs, float& fOutputScale) + SPar3DProjection*& pProjs, SProjectorParams3D& params) { assert(pVolGeom); assert(pProjGeom); @@ -164,16 +167,14 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, bool ok; - fOutputScale = 1.0f; - - ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale); + ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, params); return ok; } bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CConeProjectionGeometry3D* pProjGeom, - SConeProjection*& pProjs, float& fOutputScale) + SConeProjection*& pProjs, SProjectorParams3D& params) { assert(pVolGeom); assert(pProjGeom); @@ -192,16 +193,14 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, bool ok; - fOutputScale = 1.0f; - - ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale); + ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, params); return ok; } bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CConeVecProjectionGeometry3D* pProjGeom, - SConeProjection*& pProjs, float& fOutputScale) + SConeProjection*& pProjs, SProjectorParams3D& params) { assert(pVolGeom); assert(pProjGeom); @@ -215,9 +214,7 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, bool ok; - fOutputScale = 1.0f; - - ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, fOutputScale); + ok = convertAstraGeometry_internal(pVolGeom, nth, pProjs, params); return ok; } @@ -227,7 +224,7 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CProjectionGeometry3D* pProjGeom, SPar3DProjection*& pParProjs, SConeProjection*& pConeProjs, - float& fOutputScale) + SProjectorParams3D& params) { const CConeProjectionGeometry3D* conegeom = dynamic_cast<const CConeProjectionGeometry3D*>(pProjGeom); const CParallelProjectionGeometry3D* par3dgeom = dynamic_cast<const CParallelProjectionGeometry3D*>(pProjGeom); @@ -240,13 +237,13 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, bool ok; if (conegeom) { - ok = convertAstraGeometry(pVolGeom, conegeom, pConeProjs, fOutputScale); + ok = convertAstraGeometry(pVolGeom, conegeom, pConeProjs, params); } else if (conevec3dgeom) { - ok = convertAstraGeometry(pVolGeom, conevec3dgeom, pConeProjs, fOutputScale); + ok = convertAstraGeometry(pVolGeom, conevec3dgeom, pConeProjs, params); } else if (par3dgeom) { - ok = convertAstraGeometry(pVolGeom, par3dgeom, pParProjs, fOutputScale); + ok = convertAstraGeometry(pVolGeom, par3dgeom, pParProjs, params); } else if (parvec3dgeom) { - ok = convertAstraGeometry(pVolGeom, parvec3dgeom, pParProjs, fOutputScale); + ok = convertAstraGeometry(pVolGeom, parvec3dgeom, pParProjs, params); } else { ok = false; } @@ -260,20 +257,17 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, class AstraSIRT3d_internal { public: SDimensions3D dims; + SProjectorParams3D params; CUDAProjectionType3d projType; float* angles; float fOriginSourceDistance; float fOriginDetectorDistance; - float fSourceZ; - float fDetSize; float fRelaxation; SConeProjection* projs; SPar3DProjection* parprojs; - float fOutputScale; - bool initialized; bool setStartReconstruction; @@ -305,12 +299,9 @@ AstraSIRT3d::AstraSIRT3d() pData->dims.iProjAngles = 0; pData->dims.iProjU = 0; pData->dims.iProjV = 0; - pData->dims.iRaysPerDetDim = 1; - pData->dims.iRaysPerVoxelDim = 1; pData->projs = 0; pData->parprojs = 0; - pData->fOutputScale = 1.0f; pData->fRelaxation = 1.0f; @@ -361,7 +352,7 @@ bool AstraSIRT3d::setGeometry(const CVolumeGeometry3D* pVolGeom, ok = convertAstraGeometry(pVolGeom, pProjGeom, pData->parprojs, pData->projs, - pData->fOutputScale); + pData->params); if (!ok) return false; @@ -386,8 +377,8 @@ bool AstraSIRT3d::enableSuperSampling(unsigned int iVoxelSuperSampling, if (iVoxelSuperSampling == 0 || iDetectorSuperSampling == 0) return false; - pData->dims.iRaysPerVoxelDim = iVoxelSuperSampling; - pData->dims.iRaysPerDetDim = iDetectorSuperSampling; + pData->params.iRaysPerVoxelDim = iVoxelSuperSampling; + pData->params.iRaysPerDetDim = iDetectorSuperSampling; return true; } @@ -447,9 +438,9 @@ bool AstraSIRT3d::init() bool ok; if (pData->projType == PROJ_PARALLEL) { - ok = pData->sirt.setPar3DGeometry(pData->dims, pData->parprojs, pData->fOutputScale); + ok = pData->sirt.setPar3DGeometry(pData->dims, pData->parprojs, pData->params); } else { - ok = pData->sirt.setConeGeometry(pData->dims, pData->projs, pData->fOutputScale); + ok = pData->sirt.setConeGeometry(pData->dims, pData->projs, pData->params); } if (!ok) @@ -652,19 +643,16 @@ float AstraSIRT3d::computeDiffNorm() class AstraCGLS3d_internal { public: SDimensions3D dims; + SProjectorParams3D params; CUDAProjectionType3d projType; float* angles; float fOriginSourceDistance; float fOriginDetectorDistance; - float fSourceZ; - float fDetSize; SConeProjection* projs; SPar3DProjection* parprojs; - float fOutputScale; - bool initialized; bool setStartReconstruction; @@ -696,12 +684,9 @@ AstraCGLS3d::AstraCGLS3d() pData->dims.iProjAngles = 0; pData->dims.iProjU = 0; pData->dims.iProjV = 0; - pData->dims.iRaysPerDetDim = 1; - pData->dims.iRaysPerVoxelDim = 1; pData->projs = 0; pData->parprojs = 0; - pData->fOutputScale = 1.0f; pData->initialized = false; pData->setStartReconstruction = false; @@ -750,7 +735,7 @@ bool AstraCGLS3d::setGeometry(const CVolumeGeometry3D* pVolGeom, ok = convertAstraGeometry(pVolGeom, pProjGeom, pData->parprojs, pData->projs, - pData->fOutputScale); + pData->params); if (!ok) return false; @@ -774,8 +759,8 @@ bool AstraCGLS3d::enableSuperSampling(unsigned int iVoxelSuperSampling, if (iVoxelSuperSampling == 0 || iDetectorSuperSampling == 0) return false; - pData->dims.iRaysPerVoxelDim = iVoxelSuperSampling; - pData->dims.iRaysPerDetDim = iDetectorSuperSampling; + pData->params.iRaysPerVoxelDim = iVoxelSuperSampling; + pData->params.iRaysPerDetDim = iDetectorSuperSampling; return true; } @@ -829,9 +814,9 @@ bool AstraCGLS3d::init() bool ok; if (pData->projType == PROJ_PARALLEL) { - ok = pData->cgls.setPar3DGeometry(pData->dims, pData->parprojs, pData->fOutputScale); + ok = pData->cgls.setPar3DGeometry(pData->dims, pData->parprojs, pData->params); } else { - ok = pData->cgls.setConeGeometry(pData->dims, pData->projs, pData->fOutputScale); + ok = pData->cgls.setConeGeometry(pData->dims, pData->projs, pData->params); } if (!ok) @@ -1039,23 +1024,23 @@ bool astraCudaFP(const float* pfVolume, float* pfProjections, Cuda3DProjectionKernel projKernel) { SDimensions3D dims; + SProjectorParams3D params; + + params.iRaysPerDetDim = iDetectorSuperSampling; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); if (!ok) return false; - dims.iRaysPerDetDim = iDetectorSuperSampling; if (iDetectorSuperSampling == 0) return false; SPar3DProjection* pParProjs; SConeProjection* pConeProjs; - float outputScale; - ok = convertAstraGeometry(pVolGeom, pProjGeom, pParProjs, pConeProjs, - outputScale); + params); if (iGPUIndex != -1) { @@ -1093,10 +1078,10 @@ bool astraCudaFP(const float* pfVolume, float* pfProjections, if (pParProjs) { switch (projKernel) { case ker3d_default: - ok &= Par3DFP(D_volumeData, D_projData, dims, pParProjs, outputScale); + ok &= Par3DFP(D_volumeData, D_projData, dims, pParProjs, params); break; case ker3d_sum_square_weights: - ok &= Par3DFP_SumSqW(D_volumeData, D_projData, dims, pParProjs, outputScale*outputScale); + ok &= Par3DFP_SumSqW(D_volumeData, D_projData, dims, pParProjs, params); break; default: assert(false); @@ -1104,7 +1089,7 @@ bool astraCudaFP(const float* pfVolume, float* pfProjections, } else { switch (projKernel) { case ker3d_default: - ok &= ConeFP(D_volumeData, D_projData, dims, pConeProjs, outputScale); + ok &= ConeFP(D_volumeData, D_projData, dims, pConeProjs, params); break; default: assert(false); @@ -1129,21 +1114,20 @@ bool astraCudaBP(float* pfVolume, const float* pfProjections, int iGPUIndex, int iVoxelSuperSampling) { SDimensions3D dims; + SProjectorParams3D params; + + params.iRaysPerVoxelDim = iVoxelSuperSampling; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); if (!ok) return false; - dims.iRaysPerVoxelDim = iVoxelSuperSampling; - SPar3DProjection* pParProjs; SConeProjection* pConeProjs; - float outputScale; - ok = convertAstraGeometry(pVolGeom, pProjGeom, pParProjs, pConeProjs, - outputScale); + params); if (iGPUIndex != -1) { cudaSetDevice(iGPUIndex); @@ -1179,9 +1163,9 @@ bool astraCudaBP(float* pfVolume, const float* pfProjections, } if (pParProjs) - ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, outputScale); + ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, params); else - ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, outputScale); + ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, params); ok &= copyVolumeFromDevice(pfVolume, D_volumeData, dims, dims.iVolX); @@ -1204,21 +1188,21 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, int iGPUIndex, int iVoxelSuperSampling) { SDimensions3D dims; + SProjectorParams3D params; + + params.iRaysPerVoxelDim = iVoxelSuperSampling; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); if (!ok) return false; - dims.iRaysPerVoxelDim = iVoxelSuperSampling; SPar3DProjection* pParProjs; SConeProjection* pConeProjs; - float outputScale; - ok = convertAstraGeometry(pVolGeom, pProjGeom, pParProjs, pConeProjs, - outputScale); + params); if (iGPUIndex != -1) { cudaSetDevice(iGPUIndex); @@ -1255,9 +1239,9 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, processSino3D<opSet>(D_projData, 1.0f, dims); if (pParProjs) - ok &= Par3DBP(D_pixelWeight, D_projData, dims, pParProjs, outputScale); + ok &= Par3DBP(D_pixelWeight, D_projData, dims, pParProjs, params); else - ok &= ConeBP(D_pixelWeight, D_projData, dims, pConeProjs, outputScale); + ok &= ConeBP(D_pixelWeight, D_projData, dims, pConeProjs, params); processVol3D<opInvert>(D_pixelWeight, dims); if (!ok) { @@ -1272,9 +1256,9 @@ bool astraCudaBP_SIRTWeighted(float* pfVolume, ok &= zeroVolumeData(D_volumeData, dims); // Do BP into D_volumeData if (pParProjs) - ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, outputScale); + ok &= Par3DBP(D_volumeData, D_projData, dims, pParProjs, params); else - ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, outputScale); + ok &= ConeBP(D_volumeData, D_projData, dims, pConeProjs, params); // Multiply with weights processVol3D<opMul>(D_volumeData, D_pixelWeight, dims); @@ -1314,6 +1298,9 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections, int iGPUIndex, int iVoxelSuperSampling) { SDimensions3D dims; + SProjectorParams3D params; + + params.iRaysPerVoxelDim = iVoxelSuperSampling; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); @@ -1323,7 +1310,6 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections, if (!ok) return false; - dims.iRaysPerVoxelDim = iVoxelSuperSampling; if (iVoxelSuperSampling == 0) return false; @@ -1367,6 +1353,7 @@ bool astraCudaFDK(float* pfVolume, const float* pfProjections, // TODO: Offer interface for SrcZ, DetZ + // TODO: Voxel scaling ok &= FDK(D_volumeData, D_projData, fOriginSourceDistance, fOriginDetectorDistance, 0, 0, fDetUSize, fDetVSize, dims, pfAngles, bShortScan); diff --git a/cuda/3d/astra3d.h b/cuda/3d/astra3d.h index 2137587..3345ab8 100644 --- a/cuda/3d/astra3d.h +++ b/cuda/3d/astra3d.h @@ -37,11 +37,6 @@ namespace astra { // TODO: Switch to a class hierarchy as with the 2D algorithms -enum Cuda3DProjectionKernel { - ker3d_default = 0, - ker3d_sum_square_weights -}; - class CProjectionGeometry3D; class CParallelProjectionGeometry3D; class CParallelVecProjectionGeometry3D; @@ -50,6 +45,10 @@ class CConeVecProjectionGeometry3D; class CVolumeGeometry3D; class AstraSIRT3d_internal; +using astraCUDA3d::Cuda3DProjectionKernel; +using astraCUDA3d::ker3d_default; +using astraCUDA3d::ker3d_sum_square_weights; + class _AstraExport AstraSIRT3d { public: @@ -291,7 +290,7 @@ bool convertAstraGeometry(const CVolumeGeometry3D* pVolGeom, const CProjectionGeometry3D* pProjGeom, SPar3DProjection*& pParProjs, SConeProjection*& pConeProjs, - float& fOutputScale); + astraCUDA3d::SProjectorParams3D& params); _AstraExport bool astraCudaFP(const float* pfVolume, float* pfProjections, const CVolumeGeometry3D* pVolGeom, diff --git a/cuda/3d/cgls3d.cu b/cuda/3d/cgls3d.cu index dd0e8a0..0b0331e 100644 --- a/cuda/3d/cgls3d.cu +++ b/cuda/3d/cgls3d.cu @@ -242,7 +242,7 @@ bool doCGLS(cudaPitchedPtr& D_volumeData, CGLS cgls; bool ok = true; - ok &= cgls.setConeGeometry(dims, angles, 1.0f); + ok &= cgls.setConeGeometry(dims, angles, SProjectorParams3D()); if (D_maskData.ptr) ok &= cgls.enableVolumeMask(); diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu index 4a41f6a..f077f0d 100644 --- a/cuda/3d/cone_bp.cu +++ b/cuda/3d/cone_bp.cu @@ -154,7 +154,7 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng // supersampling version -__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) +__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale) { float* volData = (float*)D_volData; @@ -185,12 +185,12 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start if (endZ > dims.iVolZ) endZ = dims.iVolZ; - float fX = X - 0.5f*dims.iVolX + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; - float fY = Y - 0.5f*dims.iVolY + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; - float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; - const float fSubStep = 1.0f/dims.iRaysPerVoxelDim; + float fX = X - 0.5f*dims.iVolX + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; + float fY = Y - 0.5f*dims.iVolY + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; + float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; + const float fSubStep = 1.0f/iRaysPerVoxelDim; - fOutputScale /= (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim); + fOutputScale /= (iRaysPerVoxelDim*iRaysPerVoxelDim*iRaysPerVoxelDim); for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f) @@ -216,11 +216,11 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start const float fCdc = gC_C[12*angle+11]; float fXs = fX; - for (int iSubX = 0; iSubX < dims.iRaysPerVoxelDim; ++iSubX) { + for (int iSubX = 0; iSubX < iRaysPerVoxelDim; ++iSubX) { float fYs = fY; - for (int iSubY = 0; iSubY < dims.iRaysPerVoxelDim; ++iSubY) { + for (int iSubY = 0; iSubY < iRaysPerVoxelDim; ++iSubY) { float fZs = fZ; - for (int iSubZ = 0; iSubZ < dims.iRaysPerVoxelDim; ++iSubZ) { + for (int iSubZ = 0; iSubZ < iRaysPerVoxelDim; ++iSubZ) { const float fUNum = fCuc + fXs * fCux + fYs * fCuy + fZs * fCuz; const float fVNum = fCvc + fXs * fCvx + fYs * fCvy + fZs * fCvz; @@ -248,10 +248,12 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start bool ConeBP_Array(cudaPitchedPtr D_volumeData, cudaArray *D_projArray, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { bindProjDataTexture(D_projArray); + float fOutputScale = params.fOutputScale * params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ; + for (unsigned int th = 0; th < dims.iProjAngles; th += g_MaxAngles) { unsigned int angleCount = g_MaxAngles; if (th + angleCount > dims.iProjAngles) @@ -295,10 +297,10 @@ 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) + if (params.iRaysPerVoxelDim == 1) 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, fOutputScale); + dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale); } cudaTextForceKernelsCompletion(); @@ -315,14 +317,14 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData, bool ConeBP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer projections to array cudaArray* cuArray = allocateProjectionArray(dims); transferProjectionsToArray(D_projData, cuArray, dims); - bool ret = ConeBP_Array(D_volumeData, cuArray, dims, angles, fOutputScale); + bool ret = ConeBP_Array(D_volumeData, cuArray, dims, angles, params); cudaFreeArray(cuArray); diff --git a/cuda/3d/cone_bp.h b/cuda/3d/cone_bp.h index 4d3d2dd..a6e2c23 100644 --- a/cuda/3d/cone_bp.h +++ b/cuda/3d/cone_bp.h @@ -34,12 +34,12 @@ namespace astraCUDA3d { _AstraExport bool ConeBP_Array(cudaPitchedPtr D_volumeData, cudaArray *D_projArray, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); _AstraExport bool ConeBP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); } diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index 13b184f..fefcdc1 100644 --- a/cuda/3d/cone_fp.cu +++ b/cuda/3d/cone_fp.cu @@ -128,6 +128,18 @@ struct DIR_Z { __device__ float z(float f0, float f1, float f2) const { return f0; } }; +struct SCALE_CUBE { + float fOutputScale; + __device__ float scale(float a1, float a2) const { return sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; } +}; + +struct SCALE_NONCUBE { + float fScale1; + float fScale2; + float fOutputScale; + __device__ float scale(float a1, float a2) const { return sqrt(a1*a1*fScale1+a2*a2*fScale2+1.0f) * fOutputScale; } +}; + // threadIdx: x = ??? detector (u?) // y = relative angle @@ -135,11 +147,12 @@ struct DIR_Z { // blockIdx: x = ??? detector (u+v?) // y = angle block -template<class COORD> +template<class COORD, class SCALE> __global__ void cone_FP_t(float* D_projData, unsigned int projPitch, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, - const SDimensions3D dims, float fOutputScale) + const SDimensions3D dims, + SCALE sc) { COORD c; @@ -188,7 +201,7 @@ __global__ void cone_FP_t(float* D_projData, unsigned int projPitch, const float b1 = c.c1(fSrcX,fSrcY,fSrcZ) - a1 * c.c0(fSrcX,fSrcY,fSrcZ); const float b2 = c.c2(fSrcX,fSrcY,fSrcZ) - a2 * c.c0(fSrcX,fSrcY,fSrcZ); - const float fDistCorr = sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; + const float fDistCorr = sc.scale(a1, a2); float fVal = 0.0f; @@ -214,7 +227,8 @@ template<class COORD> __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, - const SDimensions3D dims, float fOutputScale) + const SDimensions3D dims, int iRaysPerDetDim, + SCALE_NONCUBE sc) { COORD c; @@ -245,7 +259,7 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, if (endSlice > c.nSlices(dims)) endSlice = c.nSlices(dims); - const float fSubStep = 1.0f/dims.iRaysPerDetDim; + const float fSubStep = 1.0f/iRaysPerDetDim; for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV) { @@ -255,9 +269,9 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, float fV = 0.0f; float fdU = detectorU - 0.5f + 0.5f*fSubStep; - for (int iSubU = 0; iSubU < dims.iRaysPerDetDim; ++iSubU, fdU+=fSubStep) { + for (int iSubU = 0; iSubU < iRaysPerDetDim; ++iSubU, fdU+=fSubStep) { float fdV = detectorV - 0.5f + 0.5f*fSubStep; - for (int iSubV = 0; iSubV < dims.iRaysPerDetDim; ++iSubV, fdV+=fSubStep) { + for (int iSubV = 0; iSubV < iRaysPerDetDim; ++iSubV, fdV+=fSubStep) { const float fDetX = fDetSX + fdU*fDetUX + fdV*fDetVX; const float fDetY = fDetSY + fdU*fDetUY + fdV*fDetVY; @@ -272,7 +286,7 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, const float b1 = c.c1(fSrcX,fSrcY,fSrcZ) - a1 * c.c0(fSrcX,fSrcY,fSrcZ); const float b2 = c.c2(fSrcX,fSrcY,fSrcZ) - a2 * c.c0(fSrcX,fSrcY,fSrcZ); - const float fDistCorr = sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; + const float fDistCorr = sc.scale(a1, a2); float fVal = 0.0f; @@ -294,14 +308,14 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, } } - D_projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] += fV / (dims.iRaysPerDetDim * dims.iRaysPerDetDim); + D_projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] += fV / (iRaysPerDetDim * iRaysPerDetDim); } } bool ConeFP_Array_internal(cudaPitchedPtr D_projData, const SDimensions3D& dims, unsigned int angleCount, const SConeProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer angles to constant memory float* tmp = new float[angleCount]; @@ -336,6 +350,36 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData, unsigned int blockEnd = 0; int blockDirection = 0; + bool cube = true; + if (abs(params.fVolScaleX / params.fVolScaleY - 1.0) > 0.00001) + cube = false; + if (abs(params.fVolScaleX / params.fVolScaleZ - 1.0) > 0.00001) + cube = false; + + SCALE_CUBE scube; + scube.fOutputScale = params.fOutputScale * params.fVolScaleX; + + SCALE_NONCUBE snoncubeX; + float fS1 = params.fVolScaleY / params.fVolScaleX; + snoncubeX.fScale1 = fS1 * fS1; + float fS2 = params.fVolScaleZ / params.fVolScaleX; + snoncubeX.fScale2 = fS2 * fS2; + snoncubeX.fOutputScale = params.fOutputScale * params.fVolScaleX; + + SCALE_NONCUBE snoncubeY; + fS1 = params.fVolScaleX / params.fVolScaleY; + snoncubeY.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleY; + snoncubeY.fScale2 = fS2 * fS2; + snoncubeY.fOutputScale = params.fOutputScale * params.fVolScaleY; + + SCALE_NONCUBE snoncubeZ; + fS1 = params.fVolScaleX / params.fVolScaleZ; + snoncubeZ.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleZ; + snoncubeZ.fScale2 = fS2 * fS2; + snoncubeZ.fOutputScale = params.fOutputScale * params.fVolScaleZ; + // timeval t; // tic(t); @@ -373,22 +417,31 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData, if (blockDirection == 0) { for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX); else - cone_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + cone_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX); } else if (blockDirection == 1) { for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY); else - cone_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + cone_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY); } else if (blockDirection == 2) { for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ); else - cone_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + cone_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ); } } @@ -414,7 +467,7 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData, bool ConeFP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer volume to array @@ -434,7 +487,7 @@ bool ConeFP(cudaPitchedPtr D_volumeData, ret = ConeFP_Array_internal(D_subprojData, dims, iEndAngle - iAngle, angles + iAngle, - fOutputScale); + params); if (!ret) break; } diff --git a/cuda/3d/cone_fp.h b/cuda/3d/cone_fp.h index 969a6d8..62c9caa 100644 --- a/cuda/3d/cone_fp.h +++ b/cuda/3d/cone_fp.h @@ -34,12 +34,12 @@ namespace astraCUDA3d { _AstraExport bool ConeFP_Array(cudaArray *D_volArray, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); _AstraExport bool ConeFP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SConeProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); } diff --git a/cuda/3d/dims3d.h b/cuda/3d/dims3d.h index 5437a85..a15c67a 100644 --- a/cuda/3d/dims3d.h +++ b/cuda/3d/dims3d.h @@ -37,6 +37,13 @@ namespace astraCUDA3d { using astra::SConeProjection; using astra::SPar3DProjection; + +enum Cuda3DProjectionKernel { + ker3d_default = 0, + ker3d_sum_square_weights +}; + + struct SDimensions3D { unsigned int iVolX; unsigned int iVolY; @@ -44,8 +51,23 @@ struct SDimensions3D { unsigned int iProjAngles; unsigned int iProjU; // number of detectors in the U direction unsigned int iProjV; // number of detectors in the V direction +}; + +struct SProjectorParams3D { + SProjectorParams3D() : + iRaysPerDetDim(1), iRaysPerVoxelDim(1), + fOutputScale(1.0f), + fVolScaleX(1.0f), fVolScaleY(1.0f), fVolScaleZ(1.0f), + ker(ker3d_default) + { } + unsigned int iRaysPerDetDim; unsigned int iRaysPerVoxelDim; + float fOutputScale; + float fVolScaleX; + float fVolScaleY; + float fVolScaleZ; + Cuda3DProjectionKernel ker; }; } diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index 0320117..9a239da 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -193,17 +193,17 @@ bool copyFromGPUMemory(float *dst, MemHandle3D src, const SSubDimensions3D &pos) bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iDetectorSuperSampling, astra::Cuda3DProjectionKernel projKernel) { SDimensions3D dims; + SProjectorParams3D params; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); if (!ok) return false; #if 1 - dims.iRaysPerDetDim = iDetectorSuperSampling; + params.iRaysPerDetDim = iDetectorSuperSampling; if (iDetectorSuperSampling == 0) return false; #else - dims.iRaysPerDetDim = 1; astra::Cuda3DProjectionKernel projKernel = astra::ker3d_default; #endif @@ -211,11 +211,9 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con SPar3DProjection* pParProjs; SConeProjection* pConeProjs; - float outputScale = 1.0f; - ok = convertAstraGeometry(pVolGeom, pProjGeom, pParProjs, pConeProjs, - outputScale); + params); if (pParProjs) { #if 0 @@ -230,10 +228,10 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con switch (projKernel) { case astra::ker3d_default: - ok &= Par3DFP(volData.d->ptr, projData.d->ptr, dims, pParProjs, outputScale); + ok &= Par3DFP(volData.d->ptr, projData.d->ptr, dims, pParProjs, params); break; case astra::ker3d_sum_square_weights: - ok &= Par3DFP_SumSqW(volData.d->ptr, projData.d->ptr, dims, pParProjs, outputScale*outputScale); + ok &= Par3DFP_SumSqW(volData.d->ptr, projData.d->ptr, dims, pParProjs, params); break; default: ok = false; @@ -241,7 +239,7 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con } else { switch (projKernel) { case astra::ker3d_default: - ok &= ConeFP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, outputScale); + ok &= ConeFP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, params); break; default: ok = false; @@ -254,30 +252,27 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling) { SDimensions3D dims; + SProjectorParams3D params; bool ok = convertAstraGeometry_dims(pVolGeom, pProjGeom, dims); if (!ok) return false; #if 1 - dims.iRaysPerVoxelDim = iVoxelSuperSampling; -#else - dims.iRaysPerVoxelDim = 1; + params.iRaysPerVoxelDim = iVoxelSuperSampling; #endif SPar3DProjection* pParProjs; SConeProjection* pConeProjs; - float outputScale = 1.0f; - ok = convertAstraGeometry(pVolGeom, pProjGeom, pParProjs, pConeProjs, - outputScale); + params); if (pParProjs) - ok &= Par3DBP(volData.d->ptr, projData.d->ptr, dims, pParProjs, outputScale); + ok &= Par3DBP(volData.d->ptr, projData.d->ptr, dims, pParProjs, params); else - ok &= ConeBP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, outputScale); + ok &= ConeBP(volData.d->ptr, projData.d->ptr, dims, pConeProjs, params); return ok; diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index cafab46..491ee2f 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -143,7 +143,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn } // supersampling version -__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) +__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale) { float* volData = (float*)D_volData; @@ -174,13 +174,13 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star if (endZ > dims.iVolZ) endZ = dims.iVolZ; - float fX = X - 0.5f*dims.iVolX + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; - float fY = Y - 0.5f*dims.iVolY + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; - float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim; + float fX = X - 0.5f*dims.iVolX + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; + float fY = Y - 0.5f*dims.iVolY + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; + float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/iRaysPerVoxelDim; - const float fSubStep = 1.0f/dims.iRaysPerVoxelDim; + const float fSubStep = 1.0f/iRaysPerVoxelDim; - fOutputScale /= (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim); + fOutputScale /= (iRaysPerVoxelDim*iRaysPerVoxelDim*iRaysPerVoxelDim); for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f) @@ -201,11 +201,11 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star const float fCvc = gC_C[8*angle+7]; float fXs = fX; - for (int iSubX = 0; iSubX < dims.iRaysPerVoxelDim; ++iSubX) { + for (int iSubX = 0; iSubX < iRaysPerVoxelDim; ++iSubX) { float fYs = fY; - for (int iSubY = 0; iSubY < dims.iRaysPerVoxelDim; ++iSubY) { + for (int iSubY = 0; iSubY < iRaysPerVoxelDim; ++iSubY) { float fZs = fZ; - for (int iSubZ = 0; iSubZ < dims.iRaysPerVoxelDim; ++iSubZ) { + for (int iSubZ = 0; iSubZ < iRaysPerVoxelDim; ++iSubZ) { const float fU = fCuc + fXs * fCux + fYs * fCuy + fZs * fCuz; const float fV = fCvc + fXs * fCvx + fYs * fCvy + fZs * fCvz; @@ -228,10 +228,12 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star bool Par3DBP_Array(cudaPitchedPtr D_volumeData, cudaArray *D_projArray, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { bindProjDataTexture(D_projArray); + float fOutputScale = params.fOutputScale * params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ; + for (unsigned int th = 0; th < dims.iProjAngles; th += g_MaxAngles) { unsigned int angleCount = g_MaxAngles; if (th + angleCount > dims.iProjAngles) @@ -274,10 +276,10 @@ 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) + if (params.iRaysPerVoxelDim == 1) 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, fOutputScale); + dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale); } cudaTextForceKernelsCompletion(); @@ -293,14 +295,14 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData, bool Par3DBP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer projections to array cudaArray* cuArray = allocateProjectionArray(dims); transferProjectionsToArray(D_projData, cuArray, dims); - bool ret = Par3DBP_Array(D_volumeData, cuArray, dims, angles, fOutputScale); + bool ret = Par3DBP_Array(D_volumeData, cuArray, dims, angles, params); cudaFreeArray(cuArray); diff --git a/cuda/3d/par3d_bp.h b/cuda/3d/par3d_bp.h index f1fc62d..3efad19 100644 --- a/cuda/3d/par3d_bp.h +++ b/cuda/3d/par3d_bp.h @@ -34,12 +34,12 @@ namespace astraCUDA3d { _AstraExport bool Par3DBP_Array(cudaPitchedPtr D_volumeData, cudaArray *D_projArray, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); _AstraExport bool Par3DBP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); } diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu index 3ce3d42..6f9ce40 100644 --- a/cuda/3d/par3d_fp.cu +++ b/cuda/3d/par3d_fp.cu @@ -128,6 +128,18 @@ struct DIR_Z { __device__ float z(float f0, float f1, float f2) const { return f0; } }; +struct SCALE_CUBE { + float fOutputScale; + __device__ float scale(float a1, float a2) const { return sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; } +}; + +struct SCALE_NONCUBE { + float fScale1; + float fScale2; + float fOutputScale; + __device__ float scale(float a1, float a2) const { return sqrt(a1*a1*fScale1+a2*a2*fScale2+1.0f) * fOutputScale; } +}; + // threadIdx: x = u detector @@ -136,11 +148,12 @@ struct DIR_Z { // y = angle block -template<class COORD> +template<class COORD, class SCALE> __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, - const SDimensions3D dims, float fOutputScale) + const SDimensions3D dims, + SCALE sc) { COORD c; @@ -161,6 +174,9 @@ __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch, const float fDetSY = gC_DetSY[angle] + 0.5f * fDetUY + 0.5f * fDetVY; const float fDetSZ = gC_DetSZ[angle] + 0.5f * fDetUZ + 0.5f * fDetVZ; + const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float fDistCorr = sc.scale(a1, a2); const int detectorU = (blockIdx.x%((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockU + threadIdx.x; @@ -186,13 +202,9 @@ __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch, /* ray: (y) = (ay) * x + (by) */ /* (z) (az) (bz) */ - const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); - const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); const float b1 = c.c1(fDetX,fDetY,fDetZ) - a1 * c.c0(fDetX,fDetY,fDetZ); const float b2 = c.c2(fDetX,fDetY,fDetZ) - a2 * c.c0(fDetX,fDetY,fDetZ); - const float fDistCorr = sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; - float fVal = 0.0f; float f0 = startSlice + 0.5f; @@ -218,7 +230,8 @@ template<class COORD> __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, - const SDimensions3D dims, float fOutputScale) + const SDimensions3D dims, int iRaysPerDetDim, + SCALE_NONCUBE sc) { COORD c; @@ -239,7 +252,9 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, const float fDetSY = gC_DetSY[angle] + 0.5f * fDetUY + 0.5f * fDetVY; const float fDetSZ = gC_DetSZ[angle] + 0.5f * fDetUZ + 0.5f * fDetVZ; - + const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float fDistCorr = sc.scale(a1, a2); const int detectorU = (blockIdx.x%((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockU + threadIdx.x; const int startDetectorV = (blockIdx.x/((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockV; @@ -251,7 +266,7 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, if (endSlice > c.nSlices(dims)) endSlice = c.nSlices(dims); - const float fSubStep = 1.0f/dims.iRaysPerDetDim; + const float fSubStep = 1.0f/iRaysPerDetDim; for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV) { @@ -259,9 +274,9 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, float fV = 0.0f; float fdU = detectorU - 0.5f + 0.5f*fSubStep; - for (int iSubU = 0; iSubU < dims.iRaysPerDetDim; ++iSubU, fdU+=fSubStep) { + for (int iSubU = 0; iSubU < iRaysPerDetDim; ++iSubU, fdU+=fSubStep) { float fdV = detectorV - 0.5f + 0.5f*fSubStep; - for (int iSubV = 0; iSubV < dims.iRaysPerDetDim; ++iSubV, fdV+=fSubStep) { + for (int iSubV = 0; iSubV < iRaysPerDetDim; ++iSubV, fdV+=fSubStep) { /* Trace ray in direction Ray to (detectorU,detectorV) from */ /* X = startSlice to X = endSlice */ @@ -274,12 +289,9 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, /* ray: (y) = (ay) * x + (by) */ /* (z) (az) (bz) */ - const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); - const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); const float b1 = c.c1(fDetX,fDetY,fDetZ) - a1 * c.c0(fDetX,fDetY,fDetZ); const float b2 = c.c2(fDetX,fDetY,fDetZ) - a2 * c.c0(fDetX,fDetY,fDetZ); - const float fDistCorr = sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; float fVal = 0.0f; @@ -295,13 +307,13 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch, f2 += a2; } - fVal *= fDistCorr; fV += fVal; } } - D_projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] += fV / (dims.iRaysPerDetDim * dims.iRaysPerDetDim); + fV *= fDistCorr; + D_projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] += fV / (iRaysPerDetDim * iRaysPerDetDim); } } @@ -324,7 +336,8 @@ template<class COORD> __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, - const SDimensions3D dims, float fOutputScale) + const SDimensions3D dims, + SCALE_NONCUBE sc) { COORD c; @@ -345,6 +358,9 @@ __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch, const float fDetSY = gC_DetSY[angle] + 0.5f * fDetUY + 0.5f * fDetVY; const float fDetSZ = gC_DetSZ[angle] + 0.5f * fDetUZ + 0.5f * fDetVZ; + const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); + const float fDistCorr = sc.scale(a1, a2); const int detectorU = (blockIdx.x%((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockU + threadIdx.x; @@ -370,13 +386,9 @@ __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch, /* ray: (y) = (ay) * x + (by) */ /* (z) (az) (bz) */ - const float a1 = c.c1(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); - const float a2 = c.c2(fRayX,fRayY,fRayZ) / c.c0(fRayX,fRayY,fRayZ); const float b1 = c.c1(fDetX,fDetY,fDetZ) - a1 * c.c0(fDetX,fDetY,fDetZ); const float b2 = c.c2(fDetX,fDetY,fDetZ) - a2 * c.c0(fDetX,fDetY,fDetZ); - const float fDistCorr = sqrt(a1*a1+a2*a2+1.0f) * fOutputScale; - float fVal = 0.0f; float f0 = startSlice + 0.5f; @@ -385,12 +397,13 @@ __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch, for (int s = startSlice; s < endSlice; ++s) { - fVal += dirWeights(f1, c.nDim1(dims)) * dirWeights(f2, c.nDim2(dims)) * fDistCorr * fDistCorr; + fVal += dirWeights(f1, c.nDim1(dims)) * dirWeights(f2, c.nDim2(dims)); f0 += 1.0f; f1 += a1; f2 += a2; } + fVal *= fDistCorr * fDistCorr; D_projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] += fVal; } } @@ -401,7 +414,7 @@ __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch, bool Par3DFP_Array_internal(cudaPitchedPtr D_projData, const SDimensions3D& dims, unsigned int angleCount, const SPar3DProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer angles to constant memory float* tmp = new float[dims.iProjAngles]; @@ -436,6 +449,36 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData, unsigned int blockEnd = 0; int blockDirection = 0; + bool cube = true; + if (abs(params.fVolScaleX / params.fVolScaleY - 1.0) > 0.00001) + cube = false; + if (abs(params.fVolScaleX / params.fVolScaleZ - 1.0) > 0.00001) + cube = false; + + SCALE_CUBE scube; + scube.fOutputScale = params.fOutputScale * params.fVolScaleX; + + SCALE_NONCUBE snoncubeX; + float fS1 = params.fVolScaleY / params.fVolScaleX; + snoncubeX.fScale1 = fS1 * fS1; + float fS2 = params.fVolScaleZ / params.fVolScaleX; + snoncubeX.fScale2 = fS2 * fS2; + snoncubeX.fOutputScale = params.fOutputScale * params.fVolScaleX; + + SCALE_NONCUBE snoncubeY; + fS1 = params.fVolScaleX / params.fVolScaleY; + snoncubeY.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleY; + snoncubeY.fScale2 = fS2 * fS2; + snoncubeY.fOutputScale = params.fOutputScale * params.fVolScaleY; + + SCALE_NONCUBE snoncubeZ; + fS1 = params.fVolScaleX / params.fVolScaleZ; + snoncubeZ.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleZ; + snoncubeZ.fScale2 = fS2 * fS2; + snoncubeZ.fOutputScale = params.fOutputScale * params.fVolScaleZ; + // timeval t; // tic(t); @@ -473,22 +516,31 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData, if (blockDirection == 0) { for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX); else - par3D_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + par3D_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX); } else if (blockDirection == 1) { for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY); else - par3D_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + par3D_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY); } else if (blockDirection == 2) { for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + if (cube) + par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + else + par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ); else - par3D_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + par3D_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ); } } @@ -514,7 +566,7 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData, bool Par3DFP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer volume to array cudaArray* cuArray = allocateVolumeArray(dims); @@ -533,7 +585,7 @@ bool Par3DFP(cudaPitchedPtr D_volumeData, ret = Par3DFP_Array_internal(D_subprojData, dims, iEndAngle - iAngle, angles + iAngle, - fOutputScale); + params); if (!ret) break; } @@ -548,7 +600,7 @@ bool Par3DFP(cudaPitchedPtr D_volumeData, bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale) + const SProjectorParams3D& params) { // transfer angles to constant memory float* tmp = new float[dims.iProjAngles]; @@ -583,6 +635,28 @@ bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, unsigned int blockEnd = 0; int blockDirection = 0; + SCALE_NONCUBE snoncubeX; + float fS1 = params.fVolScaleY / params.fVolScaleX; + snoncubeX.fScale1 = fS1 * fS1; + float fS2 = params.fVolScaleZ / params.fVolScaleX; + snoncubeX.fScale2 = fS2 * fS2; + snoncubeX.fOutputScale = params.fOutputScale * params.fVolScaleX; + + SCALE_NONCUBE snoncubeY; + fS1 = params.fVolScaleX / params.fVolScaleY; + snoncubeY.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleY; + snoncubeY.fScale2 = fS2 * fS2; + snoncubeY.fOutputScale = params.fOutputScale * params.fVolScaleY; + + SCALE_NONCUBE snoncubeZ; + fS1 = params.fVolScaleX / params.fVolScaleZ; + snoncubeZ.fScale1 = fS1 * fS1; + fS2 = params.fVolScaleY / params.fVolScaleZ; + snoncubeZ.fScale2 = fS2 * fS2; + snoncubeZ.fOutputScale = params.fOutputScale * params.fVolScaleZ; + + // timeval t; // tic(t); @@ -620,8 +694,8 @@ bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, if (blockDirection == 0) { for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_SumSqW_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + par3D_FP_SumSqW_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX); else #if 0 par3D_FP_SS_SumSqW_dirX<<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); @@ -630,8 +704,8 @@ bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, #endif } else if (blockDirection == 1) { for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_SumSqW_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + par3D_FP_SumSqW_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY); else #if 0 par3D_FP_SS_SumSqW_dirY<<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); @@ -640,8 +714,8 @@ bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, #endif } else if (blockDirection == 2) { for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices) - if (dims.iRaysPerDetDim == 1) - par3D_FP_SumSqW_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); + if (params.iRaysPerDetDim == 1) + par3D_FP_SumSqW_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ); else #if 0 par3D_FP_SS_SumSqW_dirZ<<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, fOutputScale); diff --git a/cuda/3d/par3d_fp.h b/cuda/3d/par3d_fp.h index 41f204f..5f65cd9 100644 --- a/cuda/3d/par3d_fp.h +++ b/cuda/3d/par3d_fp.h @@ -34,17 +34,17 @@ namespace astraCUDA3d { _AstraExport bool Par3DFP_Array(cudaArray *D_volArray, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); _AstraExport bool Par3DFP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); _AstraExport bool Par3DFP_SumSqW(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SDimensions3D& dims, const SPar3DProjection* angles, - float fOutputScale); + const SProjectorParams3D& params); } diff --git a/cuda/3d/sirt3d.cu b/cuda/3d/sirt3d.cu index 713944b..ba6a159 100644 --- a/cuda/3d/sirt3d.cu +++ b/cuda/3d/sirt3d.cu @@ -353,7 +353,7 @@ bool doSIRT(cudaPitchedPtr& D_volumeData, SIRT sirt; bool ok = true; - ok &= sirt.setConeGeometry(dims, angles, 1.0f); + ok &= sirt.setConeGeometry(dims, angles, SProjectorParams3D()); if (D_maskData.ptr) ok &= sirt.enableVolumeMask(); |