summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2014-12-03 11:44:20 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2014-12-03 16:13:38 +0100
commit899faabba87628698fbc02a06f4a91ba6469fd8d (patch)
tree25c784cc6f2720d9e00530c239a632e8483a9e9f
parentb3e8338a7fa4c7ed9a5954ca02fa3126aefff530 (diff)
downloadastra-899faabba87628698fbc02a06f4a91ba6469fd8d.tar.gz
astra-899faabba87628698fbc02a06f4a91ba6469fd8d.tar.bz2
astra-899faabba87628698fbc02a06f4a91ba6469fd8d.tar.xz
astra-899faabba87628698fbc02a06f4a91ba6469fd8d.zip
Move BP coordinate transformation to utility function
-rw-r--r--build/linux/Makefile.in1
-rw-r--r--cuda/3d/cone_bp.cu3
-rw-r--r--cuda/3d/dims3d.cu74
-rw-r--r--cuda/3d/dims3d.h9
-rw-r--r--cuda/3d/par3d_bp.cu3
-rw-r--r--src/ConeVecProjectionGeometry3D.cpp17
-rw-r--r--src/ParallelVecProjectionGeometry3D.cpp16
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;
+
}
//----------------------------------------------------------------------------------------