diff options
| author | Willem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be> | 2013-12-18 13:18:00 +0000 |
|---|---|---|
| committer | wpalenst <WillemJan.Palenstijn@uantwerpen.be> | 2013-12-18 13:18:00 +0000 |
| commit | 4ec8ad40294acbaec003fb70bbd61f3f35a5db24 (patch) | |
| tree | e1bbed1722ad7d9abf3f17604b8cde4f3a5de195 /cuda/2d | |
| parent | 0ac885fcb1f5c34a5cbd5b6a694a8c3177d60192 (diff) | |
| download | astra-4ec8ad40294acbaec003fb70bbd61f3f35a5db24.tar.gz astra-4ec8ad40294acbaec003fb70bbd61f3f35a5db24.tar.bz2 astra-4ec8ad40294acbaec003fb70bbd61f3f35a5db24.tar.xz astra-4ec8ad40294acbaec003fb70bbd61f3f35a5db24.zip | |
Remove angle limits in par_fp, par_bp
Diffstat (limited to 'cuda/2d')
| -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, |