diff options
| -rw-r--r-- | cuda/2d/par_bp.cu | 25 | ||||
| -rw-r--r-- | cuda/2d/par_fp.cu | 25 |
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, |