summaryrefslogtreecommitdiffstats
path: root/cuda/2d
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/2d')
-rw-r--r--cuda/2d/par_bp.cu25
-rw-r--r--cuda/2d/par_fp.cu25
2 files changed, 48 insertions, 2 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,
diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu
index 585cb06..f811f98 100644
--- a/cuda/2d/par_fp.cu
+++ b/cuda/2d/par_fp.cu
@@ -448,7 +448,7 @@ __global__ void FPvertical_simple(float* D_projData, unsigned int projPitch, uns
-bool FP_simple(float* D_volumeData, unsigned int volumePitch,
+bool FP_simple_internal(float* D_volumeData, unsigned int volumePitch,
float* D_projData, unsigned int projPitch,
const SDimensions& dims, const float* angles,
const float* TOffsets, float outputScale)
@@ -534,6 +534,29 @@ bool FP_simple(float* D_volumeData, unsigned int volumePitch,
return true;
}
+bool FP_simple(float* D_volumeData, unsigned int volumePitch,
+ float* D_projData, unsigned int projPitch,
+ const SDimensions& dims, const float* angles,
+ const float* TOffsets, float outputScale)
+{
+ 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 = FP_simple_internal(D_volumeData, volumePitch,
+ D_projData + iAngle * projPitch, projPitch,
+ subdims, angles + iAngle,
+ TOffsets ? TOffsets + iAngle : 0, outputScale);
+ if (!ret)
+ return false;
+ }
+ return true;
+}
+
bool FP(float* D_volumeData, unsigned int volumePitch,
float* D_projData, unsigned int projPitch,
const SDimensions& dims, const float* angles,