diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2014-06-19 13:47:36 +0000 |
|---|---|---|
| committer | wpalenst <Willem.Jan.Palenstijn@cwi.nl> | 2014-06-19 13:47:36 +0000 |
| commit | 97ba7288f6d665c4442b3c9873128529c7dcf508 (patch) | |
| tree | 9b82c9478b4bc33538f6f8fd539ef0c672233108 /cuda/2d/fan_bp.cu | |
| parent | 674fd8a9be846434f8a589b989e7350d8764165a (diff) | |
| download | astra-97ba7288f6d665c4442b3c9873128529c7dcf508.tar.gz astra-97ba7288f6d665c4442b3c9873128529c7dcf508.tar.bz2 astra-97ba7288f6d665c4442b3c9873128529c7dcf508.tar.xz astra-97ba7288f6d665c4442b3c9873128529c7dcf508.zip | |
Remove angle limits in fan_fp, fan_bp
Diffstat (limited to 'cuda/2d/fan_bp.cu')
| -rw-r--r-- | cuda/2d/fan_bp.cu | 50 |
1 files changed, 45 insertions, 5 deletions
diff --git a/cuda/2d/fan_bp.cu b/cuda/2d/fan_bp.cu index dd2a7b6..8983a9c 100644 --- a/cuda/2d/fan_bp.cu +++ b/cuda/2d/fan_bp.cu @@ -277,11 +277,10 @@ __global__ void devFanBP_FBPWeighted(float* D_volData, unsigned int volPitch, un } -bool FanBP(float* D_volumeData, unsigned int volumePitch, +bool FanBP_internal(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const SFanProjection* angles) { - // TODO: process angles block by block assert(dims.iProjAngles <= g_MaxAngles); bindProjDataTexture(D_projData, projPitch, dims.iProjDets, dims.iProjAngles); @@ -324,11 +323,10 @@ bool FanBP(float* D_volumeData, unsigned int volumePitch, return true; } -bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch, +bool FanBP_FBPWeighted_internal(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const SFanProjection* angles) { - // TODO: process angles block by block assert(dims.iProjAngles <= g_MaxAngles); bindProjDataTexture(D_projData, projPitch, dims.iProjDets, dims.iProjAngles); @@ -368,7 +366,6 @@ bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch, return true; } - // D_projData is a pointer to one padded sinogram line bool FanBP_SART(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, @@ -402,6 +399,49 @@ bool FanBP_SART(float* D_volumeData, unsigned int volumePitch, return true; } +bool FanBP(float* D_volumeData, unsigned int volumePitch, + float* D_projData, unsigned int projPitch, + const SDimensions& dims, const SFanProjection* angles) +{ + 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 = FanBP_internal(D_volumeData, volumePitch, + D_projData + iAngle * projPitch, projPitch, + subdims, angles + iAngle); + if (!ret) + return false; + } + return true; +} + +bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch, + float* D_projData, unsigned int projPitch, + const SDimensions& dims, const SFanProjection* angles) +{ + 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 = FanBP_FBPWeighted_internal(D_volumeData, volumePitch, + D_projData + iAngle * projPitch, projPitch, + subdims, angles + iAngle); + + if (!ret) + return false; + } + return true; +} + } |