From 4ec8ad40294acbaec003fb70bbd61f3f35a5db24 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 18 Dec 2013 13:18:00 +0000 Subject: Remove angle limits in par_fp, par_bp --- cuda/2d/par_bp.cu | 25 ++++++++++++++++++++++++- 1 file changed, 24 insertions(+), 1 deletion(-) (limited to 'cuda/2d/par_bp.cu') diff --git a/cuda/2d/par_bp.cu b/cuda/2d/par_bp.cu index 1057879..d202341 100644 --- a/cuda/2d/par_bp.cu +++ b/cuda/2d/par_bp.cu @@ -222,7 +222,7 @@ __global__ void devBP_SART(float* D_volData, unsigned int volPitch, float offset } -bool BP(float* D_volumeData, unsigned int volumePitch, +bool BP_internal(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const float* angles, const float* TOffsets) { @@ -274,6 +274,29 @@ bool BP(float* D_volumeData, unsigned int volumePitch, return true; } +bool BP(float* D_volumeData, unsigned int volumePitch, + float* D_projData, unsigned int projPitch, + const SDimensions& dims, const float* angles, const float* TOffsets) +{ + for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) { + SDimensions subdims = dims; + unsigned int iEndAngle = iAngle + g_MaxAngles; + if (iEndAngle >= dims.iProjAngles) + iEndAngle = dims.iProjAngles; + subdims.iProjAngles = iEndAngle - iAngle; + + bool ret; + ret = BP_internal(D_volumeData, volumePitch, + D_projData + iAngle * projPitch, projPitch, + subdims, angles + iAngle, + TOffsets ? TOffsets + iAngle : 0); + if (!ret) + return false; + } + return true; +} + + bool BP_SART(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, unsigned int angle, const SDimensions& dims, -- cgit v1.2.3