summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-04-04 12:09:07 +0200
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-09-25 14:10:10 +0200
commit9950fc9bf91073c0168c8847a8f6a8814690f97c (patch)
tree9dd1f95460ec7cd622256c9a82edd4296c33787f /cuda/3d
parent5ea1bf556419204511195fa5b2bedbd1318b51ff (diff)
downloadastra-9950fc9bf91073c0168c8847a8f6a8814690f97c.tar.gz
astra-9950fc9bf91073c0168c8847a8f6a8814690f97c.tar.bz2
astra-9950fc9bf91073c0168c8847a8f6a8814690f97c.tar.xz
astra-9950fc9bf91073c0168c8847a8f6a8814690f97c.zip
Small clean up of factors
Diffstat (limited to 'cuda/3d')
-rw-r--r--cuda/3d/cone_bp.cu8
-rw-r--r--cuda/3d/fdk.cu14
2 files changed, 10 insertions, 12 deletions
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index a614c29..fa35442 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -363,10 +363,12 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
bindProjDataTexture(D_projArray);
float fOutputScale;
- if (params.bFDKWeighting)
- fOutputScale = params.fOutputScale / (params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ);
- else
+ if (params.bFDKWeighting) {
+ // NB: assuming cube voxels here
+ fOutputScale = params.fOutputScale / (params.fVolScaleX);
+ } else {
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;
diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu
index 92d3ef6..d765201 100644
--- a/cuda/3d/fdk.cu
+++ b/cuda/3d/fdk.cu
@@ -60,7 +60,7 @@ __constant__ float gC_angle[g_MaxAngles];
// per-detector u/v shifts?
-__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, float fVoxSize, const SDimensions3D dims)
+__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)
{
float* projData = (float*)D_projData;
int angle = startAngle + blockIdx.y * g_anglesPerWeightBlock + threadIdx.y;
@@ -88,13 +88,10 @@ __global__ void devFDK_preweight(void* D_projData, unsigned int projPitch, unsig
// fCentralRayLength / fRayLength : the main FDK preweighting factor
// fSrcOrigin / (fDetUSize * fCentralRayLength)
// : to adjust the filter to the det width
- // || u v s || ^ 2 : see cone_bp.cu, FDKWEIGHT
// pi / (2 * iProjAngles) : scaling of the integral over angles
- // fVoxSize ^ 2 : ...
const float fW2 = fCentralRayLength / (fDetUSize * fSrcOrigin);
- const float fW3 = fVoxSize * fVoxSize;
- const float fW = fCentralRayLength * fW2 * fW3 * (M_PI / 2.0f) / (float)dims.iProjAngles;
+ const float fW = fCentralRayLength * fW2 * (M_PI / 2.0f) / (float)dims.iProjAngles;
for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV)
{
@@ -166,7 +163,7 @@ __global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, un
bool FDK_PreWeight(cudaPitchedPtr D_projData,
float fSrcOrigin, float fDetOrigin,
float fZShift,
- float fDetUSize, float fDetVSize, float fVoxSize,
+ float fDetUSize, float fDetVSize,
bool bShortScan,
const SDimensions3D& dims, const float* angles)
{
@@ -179,7 +176,7 @@ bool FDK_PreWeight(cudaPitchedPtr D_projData,
int projPitch = D_projData.pitch/sizeof(float);
- devFDK_preweight<<<dimGrid, dimBlock>>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fZShift, fDetUSize, fDetVSize, fVoxSize, dims);
+ devFDK_preweight<<<dimGrid, dimBlock>>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fZShift, fDetUSize, fDetVSize, dims);
cudaTextForceKernelsCompletion();
@@ -343,9 +340,8 @@ bool FDK(cudaPitchedPtr D_volumeData,
#if 1
- // NB: assuming cube voxels (params.fVolScaleX)
ok = FDK_PreWeight(D_projData, fSrcOrigin, fDetOrigin,
- fZShift, fDetUSize, fDetVSize, params.fVolScaleX,
+ fZShift, fDetUSize, fDetVSize,
bShortScan, dims, pfAngles);
#else
ok = true;