diff options
-rw-r--r-- | build/linux/Makefile.in | 1 | ||||
-rw-r--r-- | cuda/3d/cone_bp.cu | 3 | ||||
-rw-r--r-- | cuda/3d/dims3d.cu | 74 | ||||
-rw-r--r-- | cuda/3d/dims3d.h | 9 | ||||
-rw-r--r-- | cuda/3d/par3d_bp.cu | 3 | ||||
-rw-r--r-- | src/ConeVecProjectionGeometry3D.cpp | 17 | ||||
-rw-r--r-- | src/ParallelVecProjectionGeometry3D.cpp | 16 |
7 files changed, 117 insertions, 6 deletions
diff --git a/build/linux/Makefile.in b/build/linux/Makefile.in index 49bc981..d39c044 100644 --- a/build/linux/Makefile.in +++ b/build/linux/Makefile.in @@ -168,6 +168,7 @@ CUDA_OBJECTS=\ cuda/3d/cgls3d.lo \ cuda/3d/cone_fp.lo \ cuda/3d/cone_bp.lo \ + cuda/3d/dims3d.lo \ cuda/3d/fdk.lo \ cuda/3d/par3d_fp.lo \ cuda/3d/par3d_bp.lo \ diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu index 3b0fd70..149a655 100644 --- a/cuda/3d/cone_bp.cu +++ b/cuda/3d/cone_bp.cu @@ -267,6 +267,9 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData, // NB: We increment angles at the end of the loop body. + + // TODO: Use functions from dims3d.cu for this: + #define TRANSFER_TO_CONSTANT(expr,name) do { for (unsigned int i = 0; i < angleCount; ++i) tmp[i] = (expr) ; cudaMemcpyToSymbol(gC_##name, tmp, angleCount*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0) TRANSFER_TO_CONSTANT( (angles[i].fDetSZ - angles[i].fSrcZ)*angles[i].fDetVY - (angles[i].fDetSY - angles[i].fSrcY)*angles[i].fDetVZ , Cux ); diff --git a/cuda/3d/dims3d.cu b/cuda/3d/dims3d.cu new file mode 100644 index 0000000..83a8eba --- /dev/null +++ b/cuda/3d/dims3d.cu @@ -0,0 +1,74 @@ +/* +----------------------------------------------------------------------- +Copyright 2012 iMinds-Vision Lab, University of Antwerp + +Contact: astra@ua.ac.be +Website: http://astra.ua.ac.be + + +This file is part of the +All Scale Tomographic Reconstruction Antwerp Toolbox ("ASTRA Toolbox"). + +The ASTRA Toolbox is free software: you can redistribute it and/or modify +it under the terms of the GNU General Public License as published by +the Free Software Foundation, either version 3 of the License, or +(at your option) any later version. + +The ASTRA Toolbox is distributed in the hope that it will be useful, +but WITHOUT ANY WARRANTY; without even the implied warranty of +MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +GNU General Public License for more details. + +You should have received a copy of the GNU General Public License +along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. + +----------------------------------------------------------------------- +$Id$ +*/ + +#include "dims3d.h" + +namespace astra { + +// (See declaration in header for (mathematical) description of these functions) + + +void computeBP_UV_Coeffs(const SPar3DProjection& proj, double &fUX, double &fUY, double &fUZ, double &fUC, + double &fVX, double &fVY, double &fVZ, double &fVC) +{ + double denom = (proj.fRayX*proj.fDetUY*proj.fDetVZ - proj.fRayX*proj.fDetUZ*proj.fDetVY - proj.fRayY*proj.fDetUX*proj.fDetVZ + proj.fRayY*proj.fDetUZ*proj.fDetVX + proj.fRayZ*proj.fDetUX*proj.fDetVY - proj.fRayZ*proj.fDetUY*proj.fDetVX); + + fUX = ( - (proj.fRayY*proj.fDetVZ - proj.fRayZ*proj.fDetVY)) / denom; + fUY = ( (proj.fRayX*proj.fDetVZ - proj.fRayZ*proj.fDetVX)) / denom; + fUZ = (- (proj.fRayX*proj.fDetVY - proj.fRayY*proj.fDetVX) ) / denom; + fUC = (-(proj.fDetSY*proj.fDetVZ - proj.fDetSZ*proj.fDetVY)*proj.fRayX + (proj.fRayY*proj.fDetVZ - proj.fRayZ*proj.fDetVY)*proj.fDetSX - (proj.fRayY*proj.fDetSZ - proj.fRayZ*proj.fDetSY)*proj.fDetVX) / denom; + + fVX = ((proj.fRayY*proj.fDetUZ - proj.fRayZ*proj.fDetUY) ) / denom; + fVY = (- (proj.fRayX*proj.fDetUZ - proj.fRayZ*proj.fDetUX) ) / denom; + fVZ = ((proj.fRayX*proj.fDetUY - proj.fRayY*proj.fDetUX) ) / denom; + fVC = ((proj.fDetSY*proj.fDetUZ - proj.fDetSZ*proj.fDetUY)*proj.fRayX - (proj.fRayY*proj.fDetUZ - proj.fRayZ*proj.fDetUY)*proj.fDetSX + (proj.fRayY*proj.fDetSZ - proj.fRayZ*proj.fDetSY)*proj.fDetUX ) / denom; +} + + + +void computeBP_UV_Coeffs(const SConeProjection& proj, double &fUX, double &fUY, double &fUZ, double &fUC, + double &fVX, double &fVY, double &fVZ, double &fVC, + double &fDX, double &fDY, double &fDZ, double &fDC) +{ + fUX = (proj.fDetSZ - proj.fSrcZ)*proj.fDetVY - (proj.fDetSY - proj.fSrcY)*proj.fDetVZ; + fUY = (proj.fDetSX - proj.fSrcX)*proj.fDetVZ -(proj.fDetSZ - proj.fSrcZ)*proj.fDetVX; + fUZ = (proj.fDetSY - proj.fSrcY)*proj.fDetVX - (proj.fDetSX - proj.fSrcX)*proj.fDetVY; + fUC = (proj.fDetSY*proj.fDetVZ - proj.fDetSZ*proj.fDetVY)*proj.fSrcX - (proj.fDetSX*proj.fDetVZ - proj.fDetSZ*proj.fDetVX)*proj.fSrcY + (proj.fDetSX*proj.fDetVY - proj.fDetSY*proj.fDetVX)*proj.fSrcZ; + + fVX = (proj.fDetSY - proj.fSrcY)*proj.fDetUZ-(proj.fDetSZ - proj.fSrcZ)*proj.fDetUY; + fVY = (proj.fDetSZ - proj.fSrcZ)*proj.fDetUX - (proj.fDetSX - proj.fSrcX)*proj.fDetUZ; + fVZ = (proj.fDetSX - proj.fSrcX)*proj.fDetUY-(proj.fDetSY - proj.fSrcY)*proj.fDetUX; + fVC = -(proj.fDetSY*proj.fDetUZ - proj.fDetSZ*proj.fDetUY)*proj.fSrcX + (proj.fDetSX*proj.fDetUZ - proj.fDetSZ*proj.fDetUX)*proj.fSrcY - (proj.fDetSX*proj.fDetUY - proj.fDetSY*proj.fDetUX)*proj.fSrcZ; + + fDX = proj.fDetUY*proj.fDetVZ - proj.fDetUZ*proj.fDetVY; + fDY = proj.fDetUZ*proj.fDetVX - proj.fDetUX*proj.fDetVZ; + fDZ = proj.fDetUX*proj.fDetVY - proj.fDetUY*proj.fDetVX; + fDC = -proj.fSrcX * (proj.fDetUY*proj.fDetVZ - proj.fDetUZ*proj.fDetVY) - proj.fSrcY * (proj.fDetUZ*proj.fDetVX - proj.fDetUX*proj.fDetVZ) - proj.fSrcZ * (proj.fDetUX*proj.fDetVY - proj.fDetUY*proj.fDetVX); +} + +} diff --git a/cuda/3d/dims3d.h b/cuda/3d/dims3d.h index ec3c4a3..b5e02a4 100644 --- a/cuda/3d/dims3d.h +++ b/cuda/3d/dims3d.h @@ -59,9 +59,18 @@ struct SPar3DProjection { double fDetVX, fDetVY, fDetVZ; }; +void computeBP_UV_Coeffs(const SPar3DProjection& proj, double &fUX, double &fUY, double &fUZ, double &fUC, + double &fVX, double &fVY, double &fVZ, double &fVC); + +void computeBP_UV_Coeffs(const SConeProjection& proj, double &fUX, double &fUY, double &fUZ, double &fUC, + double &fVX, double &fVY, double &fVZ, double &fVC, + double &fDX, double &fDY, double &fDZ, double &fDC); + } + + namespace astraCUDA3d { using astra::SConeProjection; diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 58b19fe..5f8fb77 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -244,6 +244,9 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData, // NB: We increment angles at the end of the loop body. + + // TODO: Use functions from dims3d.cu for this: + #define TRANSFER_TO_CONSTANT(expr,name) do { for (unsigned int i = 0; i < angleCount; ++i) tmp[i] = (expr) ; cudaMemcpyToSymbol(gC_##name, tmp, angleCount*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0) #define DENOM (angles[i].fRayX*angles[i].fDetUY*angles[i].fDetVZ - angles[i].fRayX*angles[i].fDetUZ*angles[i].fDetVY - angles[i].fRayY*angles[i].fDetUX*angles[i].fDetVZ + angles[i].fRayY*angles[i].fDetUZ*angles[i].fDetVX + angles[i].fRayZ*angles[i].fDetUX*angles[i].fDetVY - angles[i].fRayZ*angles[i].fDetUY*angles[i].fDetVX) diff --git a/src/ConeVecProjectionGeometry3D.cpp b/src/ConeVecProjectionGeometry3D.cpp index 29d84b7..79f56c4 100644 --- a/src/ConeVecProjectionGeometry3D.cpp +++ b/src/ConeVecProjectionGeometry3D.cpp @@ -224,9 +224,20 @@ void CConeVecProjectionGeometry3D::projectPoint(float32 fX, float32 fY, float32 int iAngleIndex, float32 &fU, float32 &fV) const { -#warning implementme - fU = 0.0f/0.0f; - fV = 0.0f/0.0f; + ASTRA_ASSERT(iAngleIndex >= 0); + ASTRA_ASSERT(iAngleIndex < m_iProjectionAngleCount); + + double fUX, fUY, fUZ, fUC; + double fVX, fVY, fVZ, fVC; + double fDX, fDY, fDZ, fDC; + + computeBP_UV_Coeffs(m_pProjectionAngles[iAngleIndex], + fUX, fUY, fUZ, fUC, fVX, fVY, fVZ, fVC, fDX, fDY, fDZ, fDC); + + // The -0.5f shifts from corner to center of detector pixels + double fD = fDX*fX + fDY*fY + fDZ*fZ + fDC; + fU = (fUX*fX + fUY*fY + fUZ*fZ + fUC) / fD - 0.5f; + fV = (fVX*fX + fVY*fY + fVZ*fZ + fVC) / fD - 0.5f; } diff --git a/src/ParallelVecProjectionGeometry3D.cpp b/src/ParallelVecProjectionGeometry3D.cpp index 7c2d2cd..56597fc 100644 --- a/src/ParallelVecProjectionGeometry3D.cpp +++ b/src/ParallelVecProjectionGeometry3D.cpp @@ -222,9 +222,19 @@ void CParallelVecProjectionGeometry3D::projectPoint(float32 fX, float32 fY, floa int iAngleIndex, float32 &fU, float32 &fV) const { -#warning implementme - fU = 0.0f/0.0f; - fV = 0.0f/0.0f; + ASTRA_ASSERT(iAngleIndex >= 0); + ASTRA_ASSERT(iAngleIndex < m_iProjectionAngleCount); + + double fUX, fUY, fUZ, fUC; + double fVX, fVY, fVZ, fVC; + + computeBP_UV_Coeffs(m_pProjectionAngles[iAngleIndex], + fUX, fUY, fUZ, fUC, fVX, fVY, fVZ, fVC); + + // The -0.5f shifts from corner to center of detector pixels + fU = (fUX*fX + fUY*fY + fUZ*fZ + fUC) - 0.5f; + fV = (fVX*fX + fVY*fY + fVZ*fZ + fVC) - 0.5f; + } //---------------------------------------------------------------------------------------- |