summaryrefslogtreecommitdiffstats
path: root/cuda/2d/par_bp.cu
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/2d/par_bp.cu')
-rw-r--r--cuda/2d/par_bp.cu25
1 files changed, 24 insertions, 1 deletions
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,