From 0f4ceb4c7f3f63fddf8fbf44c59fcd8f415e3847 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 29 Mar 2019 14:42:53 +0100 Subject: Adjust line kernels to line integral scaling --- include/astra/FanFlatBeamLineKernelProjector2D.inl | 6 ++---- include/astra/ParallelBeamLineKernelProjector2D.inl | 6 ++---- 2 files changed, 4 insertions(+), 8 deletions(-) (limited to 'include') diff --git a/include/astra/FanFlatBeamLineKernelProjector2D.inl b/include/astra/FanFlatBeamLineKernelProjector2D.inl index 58dec61..8f2e673 100644 --- a/include/astra/FanFlatBeamLineKernelProjector2D.inl +++ b/include/astra/FanFlatBeamLineKernelProjector2D.inl @@ -82,8 +82,6 @@ void CFanFlatBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, in const SFanProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; - float32 detSize = sqrt(proj->fDetUX * proj->fDetUX + proj->fDetUY * proj->fDetUY); - // loop detectors for (iDetector = _iDetFrom; iDetector < _iDetTo; ++iDetector) { @@ -104,7 +102,7 @@ void CFanFlatBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, in // vertically if (vertical) { RxOverRy = Rx/Ry; - lengthPerRow = detSize * pixelLengthX * sqrt(Rx*Rx + Ry*Ry) / abs(Ry); + lengthPerRow = pixelLengthX * sqrt(Rx*Rx + Ry*Ry) / abs(Ry); deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; S = 0.5f - 0.5f*fabs(RxOverRy); T = 0.5f + 0.5f*fabs(RxOverRy); @@ -154,7 +152,7 @@ void CFanFlatBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, in // horizontally else { RyOverRx = Ry/Rx; - lengthPerCol = detSize * pixelLengthY * sqrt(Rx*Rx + Ry*Ry) / abs(Rx); + lengthPerCol = pixelLengthY * sqrt(Rx*Rx + Ry*Ry) / abs(Rx); deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; S = 0.5f - 0.5f*fabs(RyOverRx); T = 0.5f + 0.5f*fabs(RyOverRx); diff --git a/include/astra/ParallelBeamLineKernelProjector2D.inl b/include/astra/ParallelBeamLineKernelProjector2D.inl index a9f1aa0..1fc67a8 100644 --- a/include/astra/ParallelBeamLineKernelProjector2D.inl +++ b/include/astra/ParallelBeamLineKernelProjector2D.inl @@ -166,19 +166,17 @@ void CParallelBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, i const SParProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; - float32 detSize = sqrt(proj->fDetUX * proj->fDetUX + proj->fDetUY * proj->fDetUY); - bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); if (vertical) { RxOverRy = proj->fRayX/proj->fRayY; - lengthPerRow = detSize * pixelLengthX * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); + lengthPerRow = pixelLengthX * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; S = 0.5f - 0.5f*fabs(RxOverRy); T = 0.5f + 0.5f*fabs(RxOverRy); invTminSTimesLengthPerRow = lengthPerRow / (T - S); } else { RyOverRx = proj->fRayY/proj->fRayX; - lengthPerCol = detSize * pixelLengthY * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); + lengthPerCol = pixelLengthY * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; S = 0.5f - 0.5f*fabs(RyOverRx); T = 0.5f + 0.5f*fabs(RyOverRx); -- cgit v1.2.3 From 35957b6ef72749cdc520ded67a0eb8cdfd7ea655 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 29 Mar 2019 15:03:57 +0100 Subject: Adjust linear/cuda kernels to line integral scaling --- cuda/2d/astra.cu | 3 ++- cuda/2d/par_fp.cu | 10 ++++------ include/astra/ParallelBeamLinearKernelProjector2D.inl | 6 ++---- tests/python/test_line2d.py | 16 ++++------------ 4 files changed, 12 insertions(+), 23 deletions(-) (limited to 'include') diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu index ec03517..7ff1c95 100644 --- a/cuda/2d/astra.cu +++ b/cuda/2d/astra.cu @@ -302,7 +302,8 @@ static bool convertAstraGeometry_internal(const CVolumeGeometry2D* pVolGeom, pProjs[i].scale(factor); } // CHECKME: Check factor - fOutputScale *= pVolGeom->getPixelLengthX() * pVolGeom->getPixelLengthY(); + // NB: Only valid for square pixels + fOutputScale *= pVolGeom->getPixelLengthX(); return true; } diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu index 0835301..e03381c 100644 --- a/cuda/2d/par_fp.cu +++ b/cuda/2d/par_fp.cu @@ -115,10 +115,9 @@ __global__ void FPhorizontal_simple(float* D_projData, unsigned int projPitch, u float fSliceStep = cos_theta / sin_theta; float fDistCorr; if (sin_theta > 0.0f) - fDistCorr = -fDetStep; + fDistCorr = outputScale / sin_theta; else - fDistCorr = fDetStep; - fDistCorr *= outputScale; + fDistCorr = -outputScale / sin_theta; float fVal = 0.0f; // project detector on slice @@ -193,10 +192,9 @@ __global__ void FPvertical_simple(float* D_projData, unsigned int projPitch, uns float fSliceStep = sin_theta / cos_theta; float fDistCorr; if (cos_theta < 0.0f) - fDistCorr = -fDetStep; + fDistCorr = -outputScale / cos_theta; else - fDistCorr = fDetStep; - fDistCorr *= outputScale; + fDistCorr = outputScale / cos_theta; float fVal = 0.0f; float fP = (detector - 0.5f*dims.iProjDets + 0.5f - gC_angle_offset[angle]) * fDetStep + (startSlice - 0.5f*dims.iVolHeight + 0.5f) * fSliceStep + 0.5f*dims.iVolWidth - 0.5f + 0.5f; diff --git a/include/astra/ParallelBeamLinearKernelProjector2D.inl b/include/astra/ParallelBeamLinearKernelProjector2D.inl index 10d4892..1acd422 100644 --- a/include/astra/ParallelBeamLinearKernelProjector2D.inl +++ b/include/astra/ParallelBeamLinearKernelProjector2D.inl @@ -154,16 +154,14 @@ void CParallelBeamLinearKernelProjector2D::projectBlock_internal(int _iProjFrom, const SParProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; - float32 detSize = sqrt(proj->fDetUX * proj->fDetUX + proj->fDetUY * proj->fDetUY); - const bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); if (vertical) { RxOverRy = proj->fRayX/proj->fRayY; - lengthPerRow = detSize * m_pVolumeGeometry->getPixelLengthX() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); + lengthPerRow = m_pVolumeGeometry->getPixelLengthX() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; } else { RyOverRx = proj->fRayY/proj->fRayX; - lengthPerCol = detSize * m_pVolumeGeometry->getPixelLengthY() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); + lengthPerCol = m_pVolumeGeometry->getPixelLengthY() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; } diff --git a/tests/python/test_line2d.py b/tests/python/test_line2d.py index 755bd27..5647053 100644 --- a/tests/python/test_line2d.py +++ b/tests/python/test_line2d.py @@ -486,17 +486,9 @@ class Test2DKernel(unittest.TestCase): for i, (center, edge1, edge2) in enumerate(gen_lines(pg)): (src, det) = center (xd, yd) = det - src - try: - detweight = pg['DetectorWidth'] - except KeyError: - if 'fan' not in type: - detweight = effective_detweight(src, det, pg['Vectors'][i//pg['DetectorCount'],4:6]) - else: - detweight = np.linalg.norm(pg['Vectors'][i//pg['DetectorCount'],4:6], ord=2) - l = 0.0 if np.abs(xd) > np.abs(yd): # horizontal ray - length = math.sqrt(1.0 + abs(yd/xd)**2) + length = math.sqrt(1.0 + abs(yd/xd)**2) * pixsize[0] y_seg = (ymin, ymax) for j in range(rect_min[0], rect_max[0]): x = origin[0] + (-0.5 * shape[0] + j + 0.5) * pixsize[0] @@ -504,9 +496,9 @@ class Test2DKernel(unittest.TestCase): # limited interpolation precision with cuda if CUDA_8BIT_LINEAR and proj_type == 'cuda': w = np.round(w * 256.0) / 256.0 - l += w * length * pixsize[0] * detweight + l += w * length else: - length = math.sqrt(1.0 + abs(xd/yd)**2) + length = math.sqrt(1.0 + abs(xd/yd)**2) * pixsize[1] x_seg = (xmin, xmax) for j in range(rect_min[1], rect_max[1]): y = origin[1] + (+0.5 * shape[1] - j - 0.5) * pixsize[1] @@ -514,7 +506,7 @@ class Test2DKernel(unittest.TestCase): # limited interpolation precision with cuda if CUDA_8BIT_LINEAR and proj_type == 'cuda': w = np.round(w * 256.0) / 256.0 - l += w * length * pixsize[1] * detweight + l += w * length a[i] = l a = a.reshape(astra.functions.geom_size(pg)) if not np.all(np.isfinite(a)): -- cgit v1.2.3 From 15f53527e2f38e8eeb8bece79376b5e4686f3dd4 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 29 Mar 2019 16:01:02 +0100 Subject: Adjust distance driven kernels to line integral scaling --- include/astra/ParallelBeamDistanceDrivenProjector2D.inl | 9 ++++++--- tests/python/test_line2d.py | 15 +++++++++++---- 2 files changed, 17 insertions(+), 7 deletions(-) (limited to 'include') diff --git a/include/astra/ParallelBeamDistanceDrivenProjector2D.inl b/include/astra/ParallelBeamDistanceDrivenProjector2D.inl index aedcee9..3a18c6f 100644 --- a/include/astra/ParallelBeamDistanceDrivenProjector2D.inl +++ b/include/astra/ParallelBeamDistanceDrivenProjector2D.inl @@ -72,7 +72,7 @@ void CParallelBeamDistanceDrivenProjector2D::projectBlock_internal(int _iProjFro const int rowCount = m_pVolumeGeometry->getGridRowCount(); // Performance note: - // This is not a very well optimizated version of the distance driven + // This is not a very well optimized version of the distance driven // projector. The CPU projector model in ASTRA requires ray-driven iteration, // which limits re-use of intermediate computations. @@ -86,6 +86,9 @@ void CParallelBeamDistanceDrivenProjector2D::projectBlock_internal(int _iProjFro const float32 Ex = m_pVolumeGeometry->getWindowMinX() + pixelLengthX*0.5f; const float32 Ey = m_pVolumeGeometry->getWindowMaxY() - pixelLengthY*0.5f; + const float32 rayWidth = fabs(proj->fDetUX * proj->fRayY - proj->fDetUY * proj->fRayX) / + sqrt(proj->fRayX * proj->fRayX + proj->fRayY * proj->fRayY); + // loop detectors for (int iDetector = _iDetFrom; iDetector < _iDetTo; ++iDetector) { @@ -100,7 +103,7 @@ void CParallelBeamDistanceDrivenProjector2D::projectBlock_internal(int _iProjFro if (vertical) { const float32 RxOverRy = proj->fRayX/proj->fRayY; - const float32 lengthPerRow = m_pVolumeGeometry->getPixelLengthX() * m_pVolumeGeometry->getPixelLengthY(); + const float32 lengthPerRow = m_pVolumeGeometry->getPixelLengthX() * m_pVolumeGeometry->getPixelLengthY() / rayWidth; const float32 deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; const float32 deltad = 0.5f * fabs((proj->fDetUX - proj->fDetUY * RxOverRy) * inv_pixelLengthX); @@ -157,7 +160,7 @@ void CParallelBeamDistanceDrivenProjector2D::projectBlock_internal(int _iProjFro } else { const float32 RyOverRx = proj->fRayY/proj->fRayX; - const float32 lengthPerCol = m_pVolumeGeometry->getPixelLengthX() * m_pVolumeGeometry->getPixelLengthY(); + const float32 lengthPerCol = m_pVolumeGeometry->getPixelLengthX() * m_pVolumeGeometry->getPixelLengthY() / rayWidth; const float32 deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; const float32 deltad = 0.5f * fabs((proj->fDetUY - proj->fDetUX * RyOverRx) * inv_pixelLengthY); diff --git a/tests/python/test_line2d.py b/tests/python/test_line2d.py index 5647053..3019277 100644 --- a/tests/python/test_line2d.py +++ b/tests/python/test_line2d.py @@ -516,21 +516,26 @@ class Test2DKernel(unittest.TestCase): if DISPLAY and x > TOL: display_mismatch(data, sinogram, a) self.assertFalse(x > TOL) - elif proj_type == 'distance_driven': + elif proj_type == 'distance_driven' and 'par' in type: a = np.zeros(np.prod(astra.functions.geom_size(pg)), dtype=np.float32) for i, (center, edge1, edge2) in enumerate(gen_lines(pg)): - (xd, yd) = center[1] - center[0] + (src, det) = center + try: + detweight = pg['DetectorWidth'] + except KeyError: + detweight = effective_detweight(src, det, pg['Vectors'][i//pg['DetectorCount'],4:6]) + (xd, yd) = det - src l = 0.0 if np.abs(xd) > np.abs(yd): # horizontal ray y_seg = (ymin, ymax) for j in range(rect_min[0], rect_max[0]): x = origin[0] + (-0.5 * shape[0] + j + 0.5) * pixsize[0] - l += intersect_ray_vertical_segment(edge1, edge2, x, y_seg) * pixsize[0] + l += intersect_ray_vertical_segment(edge1, edge2, x, y_seg) * pixsize[0] / detweight else: x_seg = (xmin, xmax) for j in range(rect_min[1], rect_max[1]): y = origin[1] + (+0.5 * shape[1] - j - 0.5) * pixsize[1] - l += intersect_ray_horizontal_segment(edge1, edge2, y, x_seg) * pixsize[1] + l += intersect_ray_horizontal_segment(edge1, edge2, y, x_seg) * pixsize[1] / detweight a[i] = l a = a.reshape(astra.functions.geom_size(pg)) if not np.all(np.isfinite(a)): @@ -578,6 +583,8 @@ class Test2DKernel(unittest.TestCase): if DISPLAY and x > TOL: display_mismatch(data, sinogram, a) self.assertFalse(x > TOL) + else: + raise RuntimeError("Unsupported projector") def multi_test(self, type, proj_type): np.random.seed(seed) -- cgit v1.2.3 From 9c869d834ddc681299df9999787eb92bd6010dac Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 29 Mar 2019 19:08:11 +0100 Subject: Adjust strip kernels to line integral scaling --- .../astra/FanFlatBeamStripKernelProjector2D.inl | 12 +++++++---- .../astra/ParallelBeamStripKernelProjector2D.inl | 8 ++++++-- tests/python/test_line2d.py | 24 +++++++++++----------- 3 files changed, 26 insertions(+), 18 deletions(-) (limited to 'include') diff --git a/include/astra/FanFlatBeamStripKernelProjector2D.inl b/include/astra/FanFlatBeamStripKernelProjector2D.inl index 889d0a3..f5a688c 100644 --- a/include/astra/FanFlatBeamStripKernelProjector2D.inl +++ b/include/astra/FanFlatBeamStripKernelProjector2D.inl @@ -109,8 +109,12 @@ void CFanFlatBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, i // POLICY: RAY PRIOR if (!p.rayPrior(iRayIndex)) continue; - float32 dist_srcDetPixSquared = projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() + - (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * DW * DW; + float32 dist_srcDetPixSquared = projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() + (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * DW * DW; + dist_srcDetPixSquared = dist_srcDetPixSquared * dist_srcDetPixSquared / (projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() * DW * DW); + + + + //float32 InvRayWidthSquared = (projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance()) / dist_srcDetPixSquared; float32 sin_theta_left, cos_theta_left; float32 sin_theta_right, cos_theta_right; @@ -257,8 +261,8 @@ void CFanFlatBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, i // POLICY: RAY PRIOR if (!p.rayPrior(iRayIndex)) continue; - float32 dist_srcDetPixSquared = projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() + - (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * DW * DW; + float32 dist_srcDetPixSquared = projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() + (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * (iDetector + 0.5f - m_pProjectionGeometry->getDetectorCount()*0.5f) * DW * DW; + dist_srcDetPixSquared = dist_srcDetPixSquared * dist_srcDetPixSquared / (projgeom->getSourceDetectorDistance() * projgeom->getSourceDetectorDistance() * DW * DW); // get theta_l = alpha_left + theta and theta_r = alpha_right + theta float32 sin_theta_left, cos_theta_left; diff --git a/include/astra/ParallelBeamStripKernelProjector2D.inl b/include/astra/ParallelBeamStripKernelProjector2D.inl index 0d775b3..bed02ab 100644 --- a/include/astra/ParallelBeamStripKernelProjector2D.inl +++ b/include/astra/ParallelBeamStripKernelProjector2D.inl @@ -142,6 +142,10 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, const SParProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; + const float32 rayWidth = fabs(proj->fDetUX * proj->fRayY - proj->fDetUY * proj->fRayX) / + sqrt(proj->fRayX * proj->fRayX + proj->fRayY * proj->fRayY); + const float32 relPixelArea = pixelArea / rayWidth; + bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); if (vertical) { RxOverRy = proj->fRayX/proj->fRayY; @@ -219,7 +223,7 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, else if (-S < offsetL) res -= 0.5f + offsetL; else if (-T < offsetL) res -= 0.5f*(offsetL+T)*(offsetL+T)*invTminS; - p.addWeight(iRayIndex, iVolumeIndex, pixelArea*res); + p.addWeight(iRayIndex, iVolumeIndex, relPixelArea*res); p.pixelPosterior(iVolumeIndex); } } @@ -272,7 +276,7 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, else if (-S < offsetL) res -= 0.5f + offsetL; else if (-T < offsetL) res -= 0.5f*(offsetL+T)*(offsetL+T)*invTminS; - p.addWeight(iRayIndex, iVolumeIndex, pixelArea*res); + p.addWeight(iRayIndex, iVolumeIndex, relPixelArea*res); p.pixelPosterior(iVolumeIndex); } } diff --git a/tests/python/test_line2d.py b/tests/python/test_line2d.py index 3019277..ba3f7ea 100644 --- a/tests/python/test_line2d.py +++ b/tests/python/test_line2d.py @@ -20,15 +20,8 @@ nloops = 50 seed = 123 -# FAILURES: -# fan/cuda with flexible volume -# detweight for fan/cuda -# fan/strip relatively high numerical errors? -# parvec/line+linear for oblique - -# INCONSISTENCY: -# effective_detweight vs norm(detu) in line/linear (oblique) - +# KNOWN FAILURES: +# fan/strip relatively high numerical errors around 45 degrees # return length of intersection of the line through points src = (x,y) @@ -549,6 +542,7 @@ class Test2DKernel(unittest.TestCase): a = np.zeros(np.prod(astra.functions.geom_size(pg)), dtype=np.float32) for i, (center, edge1, edge2) in enumerate(gen_lines(pg)): (src, det) = center + detweight = effective_detweight(src, det, edge2[1] - edge1[1]) det_dist = np.linalg.norm(src-det, ord=2) l = 0.0 for j in range(rect_min[0], rect_max[0]): @@ -559,7 +553,7 @@ class Test2DKernel(unittest.TestCase): ymin = origin[1] + (+0.5 * shape[1] - k - 1) * pixsize[1] ymax = origin[1] + (+0.5 * shape[1] - k) * pixsize[1] ycen = 0.5 * (ymin + ymax) - scale = det_dist / np.linalg.norm( src - np.array((xcen,ycen)), ord=2 ) + scale = det_dist / (np.linalg.norm( src - np.array((xcen,ycen)), ord=2 ) * detweight) w = intersect_ray_rect(edge1, edge2, xmin, xmax, ymin, ymax) l += w * scale a[i] = l @@ -567,14 +561,20 @@ class Test2DKernel(unittest.TestCase): if not np.all(np.isfinite(a)): raise RuntimeError("Invalid value in reference sinogram") x = np.max(np.abs(sinogram-a)) - TOL = 8e-3 + # BUG: Known bug in fan/strip code around 45 degree projections causing larger errors than desirable + TOL = 4e-2 if DISPLAY and x > TOL: display_mismatch(data, sinogram, a) self.assertFalse(x > TOL) elif proj_type == 'strip': a = np.zeros(np.prod(astra.functions.geom_size(pg)), dtype=np.float32) for i, (center, edge1, edge2) in enumerate(gen_lines(pg)): - a[i] = intersect_ray_rect(edge1, edge2, xmin, xmax, ymin, ymax) + (src, det) = center + try: + detweight = pg['DetectorWidth'] + except KeyError: + detweight = effective_detweight(src, det, pg['Vectors'][i//pg['DetectorCount'],4:6]) + a[i] = intersect_ray_rect(edge1, edge2, xmin, xmax, ymin, ymax) / detweight a = a.reshape(astra.functions.geom_size(pg)) if not np.all(np.isfinite(a)): raise RuntimeError("Invalid value in reference sinogram") -- cgit v1.2.3 From 48f4e7b165404a0375db300b9fe59da92edf1cce Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Sat, 30 Mar 2019 20:49:42 +0100 Subject: Adjust FBP to line integral scaling --- cuda/2d/algo.cu | 13 +++++++------ cuda/2d/cgls.cu | 4 ++-- cuda/2d/fbp.cu | 5 ++--- include/astra/cuda/2d/algo.h | 7 +++++-- include/astra/cuda/2d/cgls.h | 2 +- src/CudaFilteredBackProjectionAlgorithm.cpp | 7 +++++++ src/CudaReconstructionAlgorithm2D.cpp | 4 +--- src/FilteredBackProjectionAlgorithm.cpp | 8 ++++++-- 8 files changed, 31 insertions(+), 19 deletions(-) (limited to 'include') diff --git a/cuda/2d/algo.cu b/cuda/2d/algo.cu index 11422ff..59ea984 100644 --- a/cuda/2d/algo.cu +++ b/cuda/2d/algo.cu @@ -153,6 +153,12 @@ bool ReconAlgo::setSuperSampling(int raysPerDet, int raysPerPixelDim) return true; } +bool ReconAlgo::setReconstructionScale(float fScale) +{ + fOutputScale *= fScale; + return true; +} + bool ReconAlgo::setVolumeMask(float* _D_maskData, unsigned int _maskPitch) { assert(useVolumeMask); @@ -242,7 +248,7 @@ bool ReconAlgo::allocateBuffers() return true; } -bool ReconAlgo::copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, float fSinogramScale, +bool ReconAlgo::copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, const float* pfReconstruction, unsigned int iReconstructionPitch, const float* pfVolMask, unsigned int iVolMaskPitch, const float* pfSinoMask, unsigned int iSinoMaskPitch) @@ -258,11 +264,6 @@ bool ReconAlgo::copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPit if (!ok) return false; - // rescale sinogram - if (fSinogramScale != 1.0f) - processSino(D_sinoData, fSinogramScale, - sinoPitch, dims); - ok = copyVolumeToDevice(pfReconstruction, iReconstructionPitch, dims, D_volumeData, volumePitch); diff --git a/cuda/2d/cgls.cu b/cuda/2d/cgls.cu index b6a9fae..905b960 100644 --- a/cuda/2d/cgls.cu +++ b/cuda/2d/cgls.cu @@ -102,14 +102,14 @@ bool CGLS::setBuffers(float* _D_volumeData, unsigned int _volumePitch, return true; } -bool CGLS::copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, float fSinogramScale, +bool CGLS::copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, const float* pfReconstruction, unsigned int iReconstructionPitch, const float* pfVolMask, unsigned int iVolMaskPitch, const float* pfSinoMask, unsigned int iSinoMaskPitch) { sliceInitialized = false; - return ReconAlgo::copyDataToGPU(pfSinogram, iSinogramPitch, fSinogramScale, pfReconstruction, iReconstructionPitch, pfVolMask, iVolMaskPitch, pfSinoMask, iSinoMaskPitch); + return ReconAlgo::copyDataToGPU(pfSinogram, iSinogramPitch, pfReconstruction, iReconstructionPitch, pfVolMask, iVolMaskPitch, pfSinoMask, iSinoMaskPitch); } bool CGLS::iterate(unsigned int iterations) diff --git a/cuda/2d/fbp.cu b/cuda/2d/fbp.cu index a5b8a7a..65f1d68 100644 --- a/cuda/2d/fbp.cu +++ b/cuda/2d/fbp.cu @@ -321,15 +321,14 @@ bool FBP::iterate(unsigned int iterations) if (fanProjs) { float fOutputScale = 1.0 / (/*fPixelSize * fPixelSize * fPixelSize * */ fFanDetSize * fFanDetSize); - ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale); + ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale * this->fOutputScale); } else { // scale by number of angles. For the fan-beam case, this is already // handled by FDK_PreWeight float fOutputScale = (M_PI / 2.0f) / (float)dims.iProjAngles; - //fOutputScale /= fDetSize * fDetSize; - ok = BP(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, parProjs, fOutputScale); + ok = BP(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, parProjs, fOutputScale * this->fOutputScale); } if(!ok) { diff --git a/include/astra/cuda/2d/algo.h b/include/astra/cuda/2d/algo.h index 3fccdef..b6ec231 100644 --- a/include/astra/cuda/2d/algo.h +++ b/include/astra/cuda/2d/algo.h @@ -56,6 +56,10 @@ public: bool setSuperSampling(int raysPerDet, int raysPerPixelDim); + // Scale the final reconstruction. + // May be called at any time after setGeometry and before iterate(). Multiple calls stack. + bool setReconstructionScale(float fScale); + virtual bool enableVolumeMask(); virtual bool enableSinogramMask(); @@ -88,8 +92,7 @@ public: // sinogram, reconstruction, volume mask and sinogram mask in system RAM, // respectively. The corresponding pitch variables give the pitches // of these buffers, measured in floats. - // The sinogram is multiplied by fSinogramScale after uploading it. - virtual bool copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, float fSinogramScale, + virtual bool copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, const float* pfReconstruction, unsigned int iReconstructionPitch, const float* pfVolMask, unsigned int iVolMaskPitch, const float* pfSinoMask, unsigned int iSinoMaskPitch); diff --git a/include/astra/cuda/2d/cgls.h b/include/astra/cuda/2d/cgls.h index 375a425..a854a74 100644 --- a/include/astra/cuda/2d/cgls.h +++ b/include/astra/cuda/2d/cgls.h @@ -47,7 +47,7 @@ public: virtual bool setBuffers(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch); - virtual bool copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, float fSinogramScale, + virtual bool copyDataToGPU(const float* pfSinogram, unsigned int iSinogramPitch, const float* pfReconstruction, unsigned int iReconstructionPitch, const float* pfVolMask, unsigned int iVolMaskPitch, const float* pfSinoMask, unsigned int iSinoMaskPitch); diff --git a/src/CudaFilteredBackProjectionAlgorithm.cpp b/src/CudaFilteredBackProjectionAlgorithm.cpp index 88e235b..69140fc 100644 --- a/src/CudaFilteredBackProjectionAlgorithm.cpp +++ b/src/CudaFilteredBackProjectionAlgorithm.cpp @@ -151,6 +151,13 @@ void CCudaFilteredBackProjectionAlgorithm::initCUDAAlgorithm() if (!ok) { ASTRA_ERROR("CCudaFilteredBackProjectionAlgorithm: Failed to set short-scan mode"); } + + const CVolumeGeometry2D& volGeom = *m_pProjector->getVolumeGeometry(); + float fPixelArea = volGeom.getPixelArea(); + ok &= pFBP->setReconstructionScale(1.0f/fPixelArea); + if (!ok) { + ASTRA_ERROR("CCudaFilteredBackProjectionAlgorithm: Failed to set reconstruction scale"); + } } diff --git a/src/CudaReconstructionAlgorithm2D.cpp b/src/CudaReconstructionAlgorithm2D.cpp index 939a026..6730cea 100644 --- a/src/CudaReconstructionAlgorithm2D.cpp +++ b/src/CudaReconstructionAlgorithm2D.cpp @@ -309,9 +309,7 @@ void CCudaReconstructionAlgorithm2D::run(int _iNrIterations) m_bAlgoInit = true; } - float fSinogramScale = 1.0f; - - ok = m_pAlgo->copyDataToGPU(m_pSinogram->getDataConst(), m_pSinogram->getGeometry()->getDetectorCount(), fSinogramScale, + ok = m_pAlgo->copyDataToGPU(m_pSinogram->getDataConst(), m_pSinogram->getGeometry()->getDetectorCount(), m_pReconstruction->getDataConst(), volgeom.getGridColCount(), m_bUseReconstructionMask ? m_pReconstructionMask->getDataConst() : 0, volgeom.getGridColCount(), m_bUseSinogramMask ? m_pSinogramMask->getDataConst() : 0, m_pSinogram->getGeometry()->getDetectorCount()); diff --git a/src/FilteredBackProjectionAlgorithm.cpp b/src/FilteredBackProjectionAlgorithm.cpp index 423dc6c..95bef3c 100644 --- a/src/FilteredBackProjectionAlgorithm.cpp +++ b/src/FilteredBackProjectionAlgorithm.cpp @@ -264,8 +264,12 @@ void CFilteredBackProjectionAlgorithm::run(int _iNrIterations) DefaultBPPolicy(m_pReconstruction, &filteredSinogram)); // Scale data - int iAngleCount = m_pProjector->getProjectionGeometry()->getProjectionAngleCount(); - (*m_pReconstruction) *= (PI/2)/iAngleCount; + const CVolumeGeometry2D& volGeom = *m_pProjector->getVolumeGeometry(); + const CProjectionGeometry2D& projGeom = *m_pProjector->getProjectionGeometry(); + + int iAngleCount = projGeom.getProjectionAngleCount(); + float fPixelArea = volGeom.getPixelArea(); + (*m_pReconstruction) *= PI/(2*iAngleCount*fPixelArea); m_pReconstruction->updateStatistics(); } -- cgit v1.2.3 From 11114e33fb504b0b74f3d239e77fe248a027cc23 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Sat, 30 Mar 2019 21:12:01 +0100 Subject: Clean up outputscale naming confusion in cuda::algo --- cuda/2d/algo.cu | 18 ++++++------------ cuda/2d/fbp.cu | 15 ++++++++++++--- cuda/2d/sart.cu | 8 ++++---- include/astra/cuda/2d/algo.h | 2 +- include/astra/cuda/2d/fbp.h | 6 ++++++ 5 files changed, 29 insertions(+), 20 deletions(-) (limited to 'include') diff --git a/cuda/2d/algo.cu b/cuda/2d/algo.cu index 59ea984..be15b25 100644 --- a/cuda/2d/algo.cu +++ b/cuda/2d/algo.cu @@ -134,8 +134,8 @@ bool ReconAlgo::setGeometry(const astra::CVolumeGeometry2D* pVolGeom, delete[] fanProjs; fanProjs = 0; - fOutputScale = 1.0f; - ok = convertAstraGeometry(pVolGeom, pProjGeom, parProjs, fanProjs, fOutputScale); + fProjectorScale = 1.0f; + ok = convertAstraGeometry(pVolGeom, pProjGeom, parProjs, fanProjs, fProjectorScale); if (!ok) return false; @@ -153,12 +153,6 @@ bool ReconAlgo::setSuperSampling(int raysPerDet, int raysPerPixelDim) return true; } -bool ReconAlgo::setReconstructionScale(float fScale) -{ - fOutputScale *= fScale; - return true; -} - bool ReconAlgo::setVolumeMask(float* _D_maskData, unsigned int _maskPitch) { assert(useVolumeMask); @@ -317,11 +311,11 @@ bool ReconAlgo::callFP(float* D_volumeData, unsigned int volumePitch, if (parProjs) { assert(!fanProjs); return FP(D_volumeData, volumePitch, D_projData, projPitch, - dims, parProjs, fOutputScale * outputScale); + dims, parProjs, fProjectorScale * outputScale); } else { assert(fanProjs); return FanFP(D_volumeData, volumePitch, D_projData, projPitch, - dims, fanProjs, fOutputScale * outputScale); + dims, fanProjs, fProjectorScale * outputScale); } } @@ -332,11 +326,11 @@ bool ReconAlgo::callBP(float* D_volumeData, unsigned int volumePitch, if (parProjs) { assert(!fanProjs); return BP(D_volumeData, volumePitch, D_projData, projPitch, - dims, parProjs, fOutputScale * outputScale); + dims, parProjs, fProjectorScale * outputScale); } else { assert(fanProjs); return FanBP(D_volumeData, volumePitch, D_projData, projPitch, - dims, fanProjs, fOutputScale * outputScale); + dims, fanProjs, fProjectorScale * outputScale); } } diff --git a/cuda/2d/fbp.cu b/cuda/2d/fbp.cu index 65f1d68..f0edc19 100644 --- a/cuda/2d/fbp.cu +++ b/cuda/2d/fbp.cu @@ -58,7 +58,8 @@ int FBP::calcFourierFilterSize(int _iDetectorCount) FBP::FBP() : ReconAlgo() { D_filter = 0; - + m_bShortScan = false; + fReconstructionScale = 1.0f; } FBP::~FBP() @@ -72,6 +73,8 @@ void FBP::reset() freeComplexOnDevice((cufftComplex *)D_filter); D_filter = 0; } + m_bShortScan = false; + fReconstructionScale = 1.0f; } bool FBP::init() @@ -79,6 +82,12 @@ bool FBP::init() return true; } +bool FBP::setReconstructionScale(float fScale) +{ + fReconstructionScale = fScale; + return true; +} + bool FBP::setFilter(const astra::SFilterConfig &_cfg) { if (D_filter) @@ -321,14 +330,14 @@ bool FBP::iterate(unsigned int iterations) if (fanProjs) { float fOutputScale = 1.0 / (/*fPixelSize * fPixelSize * fPixelSize * */ fFanDetSize * fFanDetSize); - ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale * this->fOutputScale); + ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale * fProjectorScale * fReconstructionScale); } else { // scale by number of angles. For the fan-beam case, this is already // handled by FDK_PreWeight float fOutputScale = (M_PI / 2.0f) / (float)dims.iProjAngles; - ok = BP(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, parProjs, fOutputScale * this->fOutputScale); + ok = BP(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, parProjs, fOutputScale * fProjectorScale * fReconstructionScale); } if(!ok) { diff --git a/cuda/2d/sart.cu b/cuda/2d/sart.cu index 64973ba..aff77a3 100644 --- a/cuda/2d/sart.cu +++ b/cuda/2d/sart.cu @@ -254,11 +254,11 @@ bool SART::callFP_SART(float* D_volumeData, unsigned int volumePitch, if (parProjs) { assert(!fanProjs); return FP(D_volumeData, volumePitch, D_projData, projPitch, - d, &parProjs[angle], outputScale); + d, &parProjs[angle], outputScale * fProjectorScale); } else { assert(fanProjs); return FanFP(D_volumeData, volumePitch, D_projData, projPitch, - d, &fanProjs[angle], outputScale); + d, &fanProjs[angle], outputScale * fProjectorScale); } } @@ -269,11 +269,11 @@ bool SART::callBP_SART(float* D_volumeData, unsigned int volumePitch, if (parProjs) { assert(!fanProjs); return BP_SART(D_volumeData, volumePitch, D_projData, projPitch, - angle, dims, parProjs, outputScale); + angle, dims, parProjs, outputScale * fProjectorScale); } else { assert(fanProjs); return FanBP_SART(D_volumeData, volumePitch, D_projData, projPitch, - angle, dims, fanProjs, outputScale); + angle, dims, fanProjs, outputScale * fProjectorScale); } } diff --git a/include/astra/cuda/2d/algo.h b/include/astra/cuda/2d/algo.h index b6ec231..b670b8b 100644 --- a/include/astra/cuda/2d/algo.h +++ b/include/astra/cuda/2d/algo.h @@ -136,7 +136,7 @@ protected: SDimensions dims; SParProjection* parProjs; SFanProjection* fanProjs; - float fOutputScale; + float fProjectorScale; bool freeGPUMemory; diff --git a/include/astra/cuda/2d/fbp.h b/include/astra/cuda/2d/fbp.h index 1adf3b1..3aa4741 100644 --- a/include/astra/cuda/2d/fbp.h +++ b/include/astra/cuda/2d/fbp.h @@ -79,6 +79,11 @@ public: bool setShortScan(bool ss) { m_bShortScan = ss; return true; } + // Scale the final reconstruction. + // May be called at any time before iterate(). + bool setReconstructionScale(float fScale); + + virtual bool init(); virtual bool iterate(unsigned int iterations); @@ -90,6 +95,7 @@ protected: void* D_filter; // cufftComplex* bool m_bShortScan; + float fReconstructionScale; }; } -- cgit v1.2.3 From 027dc56ffe325652314b4f9e5d804b3b20fd0db3 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Sat, 30 Mar 2019 22:43:26 +0100 Subject: Work around some warnings --- .../astra/ParallelBeamLineKernelProjector2D.inl | 29 +++++++++++----------- .../astra/ParallelBeamLinearKernelProjector2D.inl | 17 ++++++------- .../astra/ParallelBeamStripKernelProjector2D.inl | 25 +++++++++---------- 3 files changed, 34 insertions(+), 37 deletions(-) (limited to 'include') diff --git a/include/astra/ParallelBeamLineKernelProjector2D.inl b/include/astra/ParallelBeamLineKernelProjector2D.inl index 1fc67a8..903ebb6 100644 --- a/include/astra/ParallelBeamLineKernelProjector2D.inl +++ b/include/astra/ParallelBeamLineKernelProjector2D.inl @@ -167,21 +167,6 @@ void CParallelBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, i const SParProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); - if (vertical) { - RxOverRy = proj->fRayX/proj->fRayY; - lengthPerRow = pixelLengthX * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); - deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; - S = 0.5f - 0.5f*fabs(RxOverRy); - T = 0.5f + 0.5f*fabs(RxOverRy); - invTminSTimesLengthPerRow = lengthPerRow / (T - S); - } else { - RyOverRx = proj->fRayY/proj->fRayX; - lengthPerCol = pixelLengthY * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); - deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; - S = 0.5f - 0.5f*fabs(RyOverRx); - T = 0.5f + 0.5f*fabs(RyOverRx); - invTminSTimesLengthPerCol = lengthPerCol / (T - S); - } Ex = m_pVolumeGeometry->getWindowMinX() + pixelLengthX*0.5f; Ey = m_pVolumeGeometry->getWindowMaxY() - pixelLengthY*0.5f; @@ -202,6 +187,13 @@ void CParallelBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, i // vertically if (vertical) { + RxOverRy = proj->fRayX/proj->fRayY; + lengthPerRow = pixelLengthX * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); + deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; + S = 0.5f - 0.5f*fabs(RxOverRy); + T = 0.5f + 0.5f*fabs(RxOverRy); + invTminSTimesLengthPerRow = lengthPerRow / (T - S); + // calculate c for row 0 c = (Dx + (Ey - Dy)*RxOverRy - Ex) * inv_pixelLengthX; @@ -246,6 +238,13 @@ void CParallelBeamLineKernelProjector2D::projectBlock_internal(int _iProjFrom, i // horizontally else { + RyOverRx = proj->fRayY/proj->fRayX; + lengthPerCol = pixelLengthY * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); + deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; + S = 0.5f - 0.5f*fabs(RyOverRx); + T = 0.5f + 0.5f*fabs(RyOverRx); + invTminSTimesLengthPerCol = lengthPerCol / (T - S); + // calculate r for col 0 r = -(Dy + (Ex - Dx)*RyOverRx - Ey) * inv_pixelLengthY; diff --git a/include/astra/ParallelBeamLinearKernelProjector2D.inl b/include/astra/ParallelBeamLinearKernelProjector2D.inl index 1acd422..53451e5 100644 --- a/include/astra/ParallelBeamLinearKernelProjector2D.inl +++ b/include/astra/ParallelBeamLinearKernelProjector2D.inl @@ -155,15 +155,6 @@ void CParallelBeamLinearKernelProjector2D::projectBlock_internal(int _iProjFrom, const SParProjection * proj = &pVecProjectionGeometry->getProjectionVectors()[iAngle]; const bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); - if (vertical) { - RxOverRy = proj->fRayX/proj->fRayY; - lengthPerRow = m_pVolumeGeometry->getPixelLengthX() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); - deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; - } else { - RyOverRx = proj->fRayY/proj->fRayX; - lengthPerCol = m_pVolumeGeometry->getPixelLengthY() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); - deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; - } Ex = m_pVolumeGeometry->getWindowMinX() + pixelLengthX*0.5f; Ey = m_pVolumeGeometry->getWindowMaxY() - pixelLengthY*0.5f; @@ -184,6 +175,10 @@ void CParallelBeamLinearKernelProjector2D::projectBlock_internal(int _iProjFrom, // vertically if (vertical) { + RxOverRy = proj->fRayX/proj->fRayY; + lengthPerRow = m_pVolumeGeometry->getPixelLengthX() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayY); + deltac = -pixelLengthY * RxOverRy * inv_pixelLengthX; + // calculate c for row 0 c = (Dx + (Ey - Dy)*RxOverRy - Ex) * inv_pixelLengthX; @@ -207,6 +202,10 @@ void CParallelBeamLinearKernelProjector2D::projectBlock_internal(int _iProjFrom, // horizontally else { + RyOverRx = proj->fRayY/proj->fRayX; + lengthPerCol = m_pVolumeGeometry->getPixelLengthY() * sqrt(proj->fRayY*proj->fRayY + proj->fRayX*proj->fRayX) / abs(proj->fRayX); + deltar = -pixelLengthX * RyOverRx * inv_pixelLengthY; + // calculate r for col 0 r = -(Dy + (Ex - Dx)*RyOverRx - Ey) * inv_pixelLengthY; diff --git a/include/astra/ParallelBeamStripKernelProjector2D.inl b/include/astra/ParallelBeamStripKernelProjector2D.inl index bed02ab..2031560 100644 --- a/include/astra/ParallelBeamStripKernelProjector2D.inl +++ b/include/astra/ParallelBeamStripKernelProjector2D.inl @@ -147,19 +147,6 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, const float32 relPixelArea = pixelArea / rayWidth; bool vertical = fabs(proj->fRayX) < fabs(proj->fRayY); - if (vertical) { - RxOverRy = proj->fRayX/proj->fRayY; - deltac = -m_pVolumeGeometry->getPixelLengthY() * RxOverRy * inv_pixelLengthX; - S = 0.5f - 0.5f*fabs(RxOverRy); - T = 0.5f + 0.5f*fabs(RxOverRy); - invTminS = 1.0f / (T-S); - } else { - RyOverRx = proj->fRayY/proj->fRayX; - deltar = -m_pVolumeGeometry->getPixelLengthX() * RyOverRx * inv_pixelLengthY; - S = 0.5f - 0.5f*fabs(RyOverRx); - T = 0.5f + 0.5f*fabs(RyOverRx); - invTminS = 1.0f / (T-S); - } Ex = m_pVolumeGeometry->getWindowMinX() + pixelLengthX*0.5f; Ey = m_pVolumeGeometry->getWindowMaxY() - pixelLengthY*0.5f; @@ -180,6 +167,12 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, // vertically if (vertical) { + RxOverRy = proj->fRayX/proj->fRayY; + deltac = -m_pVolumeGeometry->getPixelLengthY() * RxOverRy * inv_pixelLengthX; + S = 0.5f - 0.5f*fabs(RxOverRy); + T = 0.5f + 0.5f*fabs(RxOverRy); + invTminS = 1.0f / (T-S); + // calculate cL and cR for row 0 cL = (DLx + (Ey - DLy)*RxOverRy - Ex) * inv_pixelLengthX; cR = (DRx + (Ey - DRy)*RxOverRy - Ex) * inv_pixelLengthX; @@ -233,6 +226,12 @@ void CParallelBeamStripKernelProjector2D::projectBlock_internal(int _iProjFrom, // horizontally else { + RyOverRx = proj->fRayY/proj->fRayX; + deltar = -m_pVolumeGeometry->getPixelLengthX() * RyOverRx * inv_pixelLengthY; + S = 0.5f - 0.5f*fabs(RyOverRx); + T = 0.5f + 0.5f*fabs(RyOverRx); + invTminS = 1.0f / (T-S); + // calculate rL and rR for row 0 rL = -(DLy + (Ex - DLx)*RyOverRx - Ey) * inv_pixelLengthY; rR = -(DRy + (Ex - DRx)*RyOverRx - Ey) * inv_pixelLengthY; -- cgit v1.2.3 From 4b576ee6ad461b162dd64538f54df47b65b37972 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 3 Apr 2019 22:31:55 +0200 Subject: Remove obsolete DensityWeighting option --- cuda/3d/mem3d.cu | 4 ++-- include/astra/CudaProjector3D.h | 2 -- src/CompositeGeometryManager.cpp | 4 +--- src/CudaProjector3D.cpp | 8 -------- 4 files changed, 3 insertions(+), 15 deletions(-) (limited to 'include') diff --git a/cuda/3d/mem3d.cu b/cuda/3d/mem3d.cu index 697d2d2..50cfe75 100644 --- a/cuda/3d/mem3d.cu +++ b/cuda/3d/mem3d.cu @@ -268,7 +268,7 @@ bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con return ok; } -bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling, bool bFDKWeighting) +bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling) { assert(!volData.d->arr); SDimensions3D dims; @@ -289,7 +289,7 @@ bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, con pParProjs, pConeProjs, params); - params.bFDKWeighting = bFDKWeighting; + params.bFDKWeighting = false; if (pParProjs) { if (projData.d->arr) diff --git a/include/astra/CudaProjector3D.h b/include/astra/CudaProjector3D.h index 60df7bb..9b4ff1f 100644 --- a/include/astra/CudaProjector3D.h +++ b/include/astra/CudaProjector3D.h @@ -117,7 +117,6 @@ public: int getVoxelSuperSampling() const { return m_iVoxelSuperSampling; } int getDetectorSuperSampling() const { return m_iDetectorSuperSampling; } int getGPUIndex() const { return m_iGPUIndex; } - bool getDensityWeighting() const { return m_bDensityWeighting; } protected: @@ -125,7 +124,6 @@ protected: int m_iVoxelSuperSampling; int m_iDetectorSuperSampling; int m_iGPUIndex; - bool m_bDensityWeighting; }; diff --git a/src/CompositeGeometryManager.cpp b/src/CompositeGeometryManager.cpp index 1319a87..822f746 100644 --- a/src/CompositeGeometryManager.cpp +++ b/src/CompositeGeometryManager.cpp @@ -1462,12 +1462,10 @@ static bool doJob(const CCompositeGeometryManager::TJobSet::const_iterator& iter Cuda3DProjectionKernel projKernel = ker3d_default; int detectorSuperSampling = 1; int voxelSuperSampling = 1; - bool densityWeighting = false; if (projector) { projKernel = projector->getProjectionKernel(); detectorSuperSampling = projector->getDetectorSuperSampling(); voxelSuperSampling = projector->getVoxelSuperSampling(); - densityWeighting = projector->getDensityWeighting(); } size_t inx, iny, inz; @@ -1513,7 +1511,7 @@ static bool doJob(const CCompositeGeometryManager::TJobSet::const_iterator& iter ASTRA_DEBUG("CCompositeGeometryManager::doJobs: doing BP"); - ok = astraCUDA3d::BP(((CCompositeGeometryManager::CProjectionPart*)j.pInput.get())->pGeom, srcMem->hnd, ((CCompositeGeometryManager::CVolumePart*)j.pOutput.get())->pGeom, dstMem->hnd, voxelSuperSampling, densityWeighting); + ok = astraCUDA3d::BP(((CCompositeGeometryManager::CProjectionPart*)j.pInput.get())->pGeom, srcMem->hnd, ((CCompositeGeometryManager::CVolumePart*)j.pOutput.get())->pGeom, dstMem->hnd, voxelSuperSampling); if (!ok) ASTRA_ERROR("Error performing sub-BP"); ASTRA_DEBUG("CCompositeGeometryManager::doJobs: BP done"); } diff --git a/src/CudaProjector3D.cpp b/src/CudaProjector3D.cpp index 3ea7043..e5c55cc 100644 --- a/src/CudaProjector3D.cpp +++ b/src/CudaProjector3D.cpp @@ -67,7 +67,6 @@ void CCudaProjector3D::_clear() m_iVoxelSuperSampling = 1; m_iDetectorSuperSampling = 1; m_iGPUIndex = -1; - m_bDensityWeighting = false; } //---------------------------------------------------------------------------------------- @@ -132,13 +131,6 @@ bool CCudaProjector3D::initialize(const Config& _cfg) m_iDetectorSuperSampling = (int)_cfg.self.getOptionNumerical("DetectorSuperSampling", 1); CC.markOptionParsed("DetectorSuperSampling"); - if (dynamic_cast(m_pProjectionGeometry) || - dynamic_cast(m_pProjectionGeometry)) - { - m_bDensityWeighting = _cfg.self.getOptionBool("DensityWeighting", false); - CC.markOptionParsed("DensityWeighting"); - } - m_iGPUIndex = (int)_cfg.self.getOptionNumerical("GPUindex", -1); m_iGPUIndex = (int)_cfg.self.getOptionNumerical("GPUIndex", m_iGPUIndex); CC.markOptionParsed("GPUIndex"); -- cgit v1.2.3 From 744605ed826196aa6e0e3e3b5e945e50d830ed3a Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 3 Apr 2019 22:32:11 +0200 Subject: Add feature flags for changed scaling behaviour --- include/astra/Features.h | 16 ++++++++++++++-- src/Features.cpp | 6 ++++++ 2 files changed, 20 insertions(+), 2 deletions(-) (limited to 'include') diff --git a/include/astra/Features.h b/include/astra/Features.h index d88ae71..c7ef98c 100644 --- a/include/astra/Features.h +++ b/include/astra/Features.h @@ -38,10 +38,22 @@ _AstraExport bool hasFeature(const std::string &feature); FEATURES: -cuda: is cuda support compiled in? +cuda + is cuda support compiled in? NB: To check if there is also actually a usable GPU, use cudaAvailable() -mex_link: is there support for the matlab command astra_mex_data3d('link')? +mex_link + is there support for the matlab command astra_mex_data3d('link')? + +projectors_scaled_as_line_integrals + This is set since all 2D and 3D, CPU and GPU projectors scale their outputs + to approximate line integrals. (Previously, some 2D projectors were scaled + as area integrals.) + +fan_cone_BP_density_weighting_by_default + This is set since fan beam and cone beam BP operations perform ray density + weighting by default to more closely approximate the true mathematical adjoint. + The DensityWeighting cuda3d projector option is removed. For future backward-incompatible changes, extra features will be added here diff --git a/src/Features.cpp b/src/Features.cpp index 9114131..09a3499 100644 --- a/src/Features.cpp +++ b/src/Features.cpp @@ -34,6 +34,12 @@ _AstraExport bool hasFeature(const std::string &flag) { if (flag == "cuda") { return cudaEnabled(); } + if (flag == "projectors_scaled_as_line_integrals") { + return true; + } + if (flag == "fan_cone_BP_density_weighting_by_default") { + return true; + } return false; } -- cgit v1.2.3 From 9bf1863c585506fda6b2fd23de08164bb4fb530d Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 3 Apr 2019 23:10:26 +0200 Subject: Add missing header --- include/astra/cuda/3d/mem3d.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'include') diff --git a/include/astra/cuda/3d/mem3d.h b/include/astra/cuda/3d/mem3d.h index 8c3956e..fff1490 100644 --- a/include/astra/cuda/3d/mem3d.h +++ b/include/astra/cuda/3d/mem3d.h @@ -110,7 +110,7 @@ bool copyIntoArray(MemHandle3D handle, MemHandle3D subdata, const SSubDimensions bool FP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iDetectorSuperSampling, astra::Cuda3DProjectionKernel projKernel); -bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling, bool bFDKWeighting); +bool BP(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, int iVoxelSuperSampling); bool FDK(const astra::CProjectionGeometry3D* pProjGeom, MemHandle3D projData, const astra::CVolumeGeometry3D* pVolGeom, MemHandle3D volData, bool bShortScan, const float *pfFilter = 0); -- cgit v1.2.3 From 9950fc9bf91073c0168c8847a8f6a8814690f97c Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Thu, 4 Apr 2019 12:09:07 +0200 Subject: Small clean up of factors --- cuda/2d/fbp.cu | 2 +- cuda/3d/cone_bp.cu | 8 +++++--- cuda/3d/fdk.cu | 14 +++++--------- include/astra/cuda/3d/fdk.h | 2 +- 4 files changed, 12 insertions(+), 14 deletions(-) (limited to 'include') diff --git a/cuda/2d/fbp.cu b/cuda/2d/fbp.cu index 28cdd92..4fc3983 100644 --- a/cuda/2d/fbp.cu +++ b/cuda/2d/fbp.cu @@ -301,7 +301,7 @@ bool FBP::iterate(unsigned int iterations) astraCUDA3d::FDK_PreWeight(tmp, fOriginSource, fOriginDetector, 0.0f, - fFanDetSize, 1.0f, /* fPixelSize, but is normalized */ 1.0f, + fFanDetSize, 1.0f, m_bShortScan, dims3d, pfAngles); } else { // TODO: How should different detector pixel size in different 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<<>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fZShift, fDetUSize, fDetVSize, fVoxSize, dims); + devFDK_preweight<<>>(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; diff --git a/include/astra/cuda/3d/fdk.h b/include/astra/cuda/3d/fdk.h index 3b1a9e1..6817154 100644 --- a/include/astra/cuda/3d/fdk.h +++ b/include/astra/cuda/3d/fdk.h @@ -35,7 +35,7 @@ namespace astraCUDA3d { 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); -- cgit v1.2.3 From 2e334ee443a8fd3ce0a6218dd877117a67163533 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Fri, 5 Apr 2019 16:07:21 +0200 Subject: Adjust par3d adjoint scaling, and clean up --- cuda/3d/cone_bp.cu | 46 --------------------- cuda/3d/par3d_bp.cu | 91 +++++++++++++++++++++++------------------- include/astra/cuda/3d/util3d.h | 47 ++++++++++++++++++++++ 3 files changed, 96 insertions(+), 88 deletions(-) (limited to 'include') diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu index fa35442..024c791 100644 --- a/cuda/3d/cone_bp.cu +++ b/cuda/3d/cone_bp.cu @@ -241,52 +241,6 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start } } -struct Vec3 { - double x; - double y; - double z; - Vec3(double x_, double y_, double z_) : x(x_), y(y_), z(z_) { } - Vec3 operator+(const Vec3 &b) const { - return Vec3(x + b.x, y + b.y, z + b.z); - } - Vec3 operator-(const Vec3 &b) const { - return Vec3(x - b.x, y - b.y, z - b.z); - } - Vec3 operator-() const { - return Vec3(-x, -y, -z); - } - double norm() const { - return sqrt(x*x + y*y + z*z); - } -}; - -double det3x(const Vec3 &b, const Vec3 &c) { - return (b.y * c.z - b.z * c.y); -} -double det3y(const Vec3 &b, const Vec3 &c) { - return -(b.x * c.z - b.z * c.x); -} - -double det3z(const Vec3 &b, const Vec3 &c) { - return (b.x * c.y - b.y * c.x); -} - -double det3(const Vec3 &a, const Vec3 &b, const Vec3 &c) { - return a.x * det3x(b,c) + a.y * det3y(b,c) + a.z * det3z(b,c); -} - -Vec3 cross3(const Vec3 &a, const Vec3 &b) { - return Vec3(det3x(a,b), det3y(a,b), det3z(a,b)); -} - -Vec3 scaled_cross3(const Vec3 &a, const Vec3 &b, const Vec3 &sc) { - Vec3 ret = cross3(a, b); - ret.x *= sc.y * sc.z; - ret.y *= sc.x * sc.z; - ret.z *= sc.x * sc.y; - return ret; -} - bool transferConstants(const SConeProjection* angles, unsigned int iProjAngles, const SProjectorParams3D& params) { diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 3656f78..3673fa8 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -55,7 +55,13 @@ static const unsigned int g_volBlockY = 32; static const unsigned g_MaxAngles = 1024; -__constant__ float gC_C[8*g_MaxAngles]; +struct DevPar3DParams { + float4 fNumU; + float4 fNumV; +}; + +__constant__ DevPar3DParams gC_C[g_MaxAngles]; +__constant__ float gC_scale[g_MaxAngles]; static bool bindProjDataTexture(const cudaArray* array) @@ -115,8 +121,9 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn for (int angle = startAngle; angle < endAngle; ++angle, fAngle += 1.0f) { - float4 fCu = make_float4(gC_C[8*angle+0], gC_C[8*angle+1], gC_C[8*angle+2], gC_C[8*angle+3]); - float4 fCv = make_float4(gC_C[8*angle+4], gC_C[8*angle+5], gC_C[8*angle+6], gC_C[8*angle+7]); + float4 fCu = gC_C[angle].fNumU; + float4 fCv = gC_C[angle].fNumV; + float fS = gC_scale[angle]; float fU = fCu.w + fX * fCu.x + fY * fCu.y + fZ * fCu.z; float fV = fCv.w + fX * fCv.x + fY * fCv.y + fZ * fCv.z; @@ -124,7 +131,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn for (int idx = 0; idx < ZSIZE; ++idx) { float fVal = tex3D(gT_par3DProjTexture, fU, fAngle, fV); - Z[idx] += fVal; + Z[idx] += fVal * fS; fU += fCu.z; fV += fCv.z; @@ -190,14 +197,9 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star for (int angle = startAngle; angle < endAngle; ++angle, fAngle += 1.0f) { - const float fCux = gC_C[8*angle+0]; - const float fCuy = gC_C[8*angle+1]; - const float fCuz = gC_C[8*angle+2]; - const float fCuc = gC_C[8*angle+3]; - const float fCvx = gC_C[8*angle+4]; - const float fCvy = gC_C[8*angle+5]; - const float fCvz = gC_C[8*angle+6]; - const float fCvc = gC_C[8*angle+7]; + float4 fCu = gC_C[angle].fNumU; + float4 fCv = gC_C[angle].fNumV; + float fS = gC_scale[angle]; float fXs = fX; for (int iSubX = 0; iSubX < iRaysPerVoxelDim; ++iSubX) { @@ -206,10 +208,10 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star float fZs = fZ; 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; + const float fU = fCu.w + fXs * fCu.x + fYs * fCu.y + fZs * fCu.z; + const float fV = fCv.w + fXs * fCv.x + fYs * fCv.y + fZs * fCv.z; - fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV); + fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV) * fS; fZs += fSubStep; } fYs += fSubStep; @@ -224,6 +226,35 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star } +bool transferConstants(const SPar3DProjection* angles, unsigned int iProjAngles, const SProjectorParams3D& params) +{ + DevPar3DParams *p = new DevPar3DParams[iProjAngles]; + float *s = new float[iProjAngles]; + + for (unsigned int i = 0; i < iProjAngles; ++i) { + Vec3 u(angles[i].fDetUX, angles[i].fDetUY, angles[i].fDetUZ); + Vec3 v(angles[i].fDetVX, angles[i].fDetVY, angles[i].fDetVZ); + Vec3 r(angles[i].fRayX, angles[i].fRayY, angles[i].fRayZ); + Vec3 d(angles[i].fDetSX, angles[i].fDetSY, angles[i].fDetSZ); + + double fDen = det3(r,u,v); + p[i].fNumU.x = -det3x(r,v) / fDen; + p[i].fNumU.y = -det3y(r,v) / fDen; + p[i].fNumU.z = -det3z(r,v) / fDen; + p[i].fNumU.w = -det3(r,d,v) / fDen; + p[i].fNumV.x = det3x(r,u) / fDen; + p[i].fNumV.y = det3y(r,u) / fDen; + p[i].fNumV.z = det3z(r,u) / fDen; + p[i].fNumV.w = det3(r,d,u) / fDen; + + s[i] = 1.0 / scaled_cross3(u,v,Vec3(params.fVolScaleX,params.fVolScaleY,params.fVolScaleZ)).norm(); + } + + cudaMemcpyToSymbol(gC_C, p, iProjAngles*sizeof(DevPar3DParams), 0, cudaMemcpyHostToDevice); + cudaMemcpyToSymbol(gC_scale, s, iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); + return true; +} + bool Par3DBP_Array(cudaPitchedPtr D_volumeData, cudaArray *D_projArray, const SDimensions3D& dims, const SPar3DProjection* angles, @@ -238,33 +269,9 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData, if (th + angleCount > dims.iProjAngles) angleCount = dims.iProjAngles - th; - // transfer angles to constant memory - float* tmp = new float[8*dims.iProjAngles]; - - // 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[8*i + name] = (expr) ; } 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) - - TRANSFER_TO_CONSTANT( ( - (angles[i].fRayY*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVY)) / DENOM , 0 ); - TRANSFER_TO_CONSTANT( ( (angles[i].fRayX*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVX)) / DENOM , 1 ); - TRANSFER_TO_CONSTANT( (- (angles[i].fRayX*angles[i].fDetVY - angles[i].fRayY*angles[i].fDetVX) ) / DENOM , 2 ); - TRANSFER_TO_CONSTANT( (-(angles[i].fDetSY*angles[i].fDetVZ - angles[i].fDetSZ*angles[i].fDetVY)*angles[i].fRayX + (angles[i].fRayY*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVY)*angles[i].fDetSX - (angles[i].fRayY*angles[i].fDetSZ - angles[i].fRayZ*angles[i].fDetSY)*angles[i].fDetVX) / DENOM , 3 ); - - TRANSFER_TO_CONSTANT( ((angles[i].fRayY*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUY) ) / DENOM , 4 ); - TRANSFER_TO_CONSTANT( (- (angles[i].fRayX*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUX) ) / DENOM , 5 ); - TRANSFER_TO_CONSTANT( ((angles[i].fRayX*angles[i].fDetUY - angles[i].fRayY*angles[i].fDetUX) ) / DENOM , 6 ); - TRANSFER_TO_CONSTANT( ((angles[i].fDetSY*angles[i].fDetUZ - angles[i].fDetSZ*angles[i].fDetUY)*angles[i].fRayX - (angles[i].fRayY*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUY)*angles[i].fDetSX + (angles[i].fRayY*angles[i].fDetSZ - angles[i].fRayZ*angles[i].fDetSY)*angles[i].fDetUX ) / DENOM , 7 ); - -#undef TRANSFER_TO_CONSTANT -#undef DENOM - cudaMemcpyToSymbol(gC_C, tmp, angleCount*8*sizeof(float), 0, cudaMemcpyHostToDevice); - - delete[] tmp; + bool ok = transferConstants(angles, dims.iProjAngles, params); + if (!ok) + return false; dim3 dimBlock(g_volBlockX, g_volBlockY); diff --git a/include/astra/cuda/3d/util3d.h b/include/astra/cuda/3d/util3d.h index 0146d2d..200abfc 100644 --- a/include/astra/cuda/3d/util3d.h +++ b/include/astra/cuda/3d/util3d.h @@ -64,6 +64,53 @@ float dotProduct3D(cudaPitchedPtr data, unsigned int x, unsigned int y, unsigned int calcNextPowerOfTwo(int _iValue); +struct Vec3 { + double x; + double y; + double z; + Vec3(double x_, double y_, double z_) : x(x_), y(y_), z(z_) { } + Vec3 operator+(const Vec3 &b) const { + return Vec3(x + b.x, y + b.y, z + b.z); + } + Vec3 operator-(const Vec3 &b) const { + return Vec3(x - b.x, y - b.y, z - b.z); + } + Vec3 operator-() const { + return Vec3(-x, -y, -z); + } + double norm() const { + return sqrt(x*x + y*y + z*z); + } +}; + +static double det3x(const Vec3 &b, const Vec3 &c) { + return (b.y * c.z - b.z * c.y); +} +static double det3y(const Vec3 &b, const Vec3 &c) { + return -(b.x * c.z - b.z * c.x); +} + +static double det3z(const Vec3 &b, const Vec3 &c) { + return (b.x * c.y - b.y * c.x); +} + +static double det3(const Vec3 &a, const Vec3 &b, const Vec3 &c) { + return a.x * det3x(b,c) + a.y * det3y(b,c) + a.z * det3z(b,c); +} + +static Vec3 cross3(const Vec3 &a, const Vec3 &b) { + return Vec3(det3x(a,b), det3y(a,b), det3z(a,b)); +} + +static Vec3 scaled_cross3(const Vec3 &a, const Vec3 &b, const Vec3 &sc) { + Vec3 ret = cross3(a, b); + ret.x *= sc.y * sc.z; + ret.y *= sc.x * sc.z; + ret.z *= sc.x * sc.y; + return ret; +} + + } #endif -- cgit v1.2.3