summaryrefslogtreecommitdiffstats
path: root/cuda
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-04-04 14:37:36 +0200
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-09-25 14:10:10 +0200
commit980f443c9fea3dc30b89bb728268fb23de2b9814 (patch)
treea30f955da1a13cc64c23c51ab3e44261b4512a8d /cuda
parent9950fc9bf91073c0168c8847a8f6a8814690f97c (diff)
downloadastra-980f443c9fea3dc30b89bb728268fb23de2b9814.tar.gz
astra-980f443c9fea3dc30b89bb728268fb23de2b9814.tar.bz2
astra-980f443c9fea3dc30b89bb728268fb23de2b9814.tar.xz
astra-980f443c9fea3dc30b89bb728268fb23de2b9814.zip
Fix part of non-cube cone scaling
Diffstat (limited to 'cuda')
-rw-r--r--cuda/3d/cone_fp.cu2
-rw-r--r--cuda/3d/fdk.cu5
2 files changed, 5 insertions, 2 deletions
diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu
index 7e0fae8..418f8fd 100644
--- a/cuda/3d/cone_fp.cu
+++ b/cuda/3d/cone_fp.cu
@@ -368,7 +368,7 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData,
SCALE_NONCUBE snoncubeY;
fS1 = params.fVolScaleX / params.fVolScaleY;
snoncubeY.fScale1 = fS1 * fS1;
- fS2 = params.fVolScaleY / params.fVolScaleY;
+ fS2 = params.fVolScaleZ / params.fVolScaleY;
snoncubeY.fScale2 = fS2 * fS2;
snoncubeY.fOutputScale = params.fOutputScale * params.fVolScaleY;
diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu
index d765201..af95b6b 100644
--- a/cuda/3d/fdk.cu
+++ b/cuda/3d/fdk.cu
@@ -57,7 +57,10 @@ static const unsigned g_MaxAngles = 12000;
__constant__ float gC_angle[g_MaxAngles];
-// per-detector u/v shifts?
+
+// TODO: To support non-cube voxels, preweighting needs per-view
+// parameters. NB: Need to properly take into account the
+// anisotropic volume normalization done for that too.
__global__ void devFDK_preweight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fZShift, float fDetUSize, float fDetVSize, const SDimensions3D dims)