/*
-----------------------------------------------------------------------
Copyright: 2010-2018, imec Vision Lab, University of Antwerp
2014-2018, CWI, Amsterdam
Contact: astra@astra-toolbox.com
Website: http://www.astra-toolbox.com/
This file is part of the ASTRA Toolbox.
The ASTRA Toolbox is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
The ASTRA Toolbox is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with the ASTRA Toolbox. If not, see .
-----------------------------------------------------------------------
*/
#include "astra/cuda/2d/fbp.h"
#include "astra/cuda/2d/fft.h"
#include "astra/cuda/2d/par_bp.h"
#include "astra/cuda/2d/fan_bp.h"
#include "astra/cuda/2d/util.h"
// For fan-beam preweighting
#include "astra/cuda/3d/fdk.h"
#include "astra/Logging.h"
#include
namespace astraCUDA {
static int calcNextPowerOfTwo(int n)
{
int x = 1;
while (x < n)
x *= 2;
return x;
}
// static
int FBP::calcFourierFilterSize(int _iDetectorCount)
{
int iFFTRealDetCount = calcNextPowerOfTwo(2 * _iDetectorCount);
int iFreqBinCount = calcFFTFourierSize(iFFTRealDetCount);
// CHECKME: Matlab makes this at least 64. Do we also need to?
return iFreqBinCount;
}
FBP::FBP() : ReconAlgo()
{
D_filter = 0;
}
FBP::~FBP()
{
reset();
}
void FBP::reset()
{
if (D_filter) {
freeComplexOnDevice((cufftComplex *)D_filter);
D_filter = 0;
}
}
bool FBP::init()
{
return true;
}
bool FBP::setFilter(const astra::SFilterConfig &_cfg)
{
if (D_filter)
{
freeComplexOnDevice((cufftComplex*)D_filter);
D_filter = 0;
}
if (_cfg.m_eType == astra::FILTER_NONE)
return true; // leave D_filter set to 0
int iFFTRealDetCount = calcNextPowerOfTwo(2 * dims.iProjDets);
int iFreqBinCount = calcFFTFourierSize(iFFTRealDetCount);
cufftComplex * pHostFilter = new cufftComplex[dims.iProjAngles * iFreqBinCount];
memset(pHostFilter, 0, sizeof(cufftComplex) * dims.iProjAngles * iFreqBinCount);
allocateComplexOnDevice(dims.iProjAngles, iFreqBinCount, (cufftComplex**)&D_filter);
switch(_cfg.m_eType)
{
case astra::FILTER_NONE:
// handled above
break;
case astra::FILTER_RAMLAK:
case astra::FILTER_SHEPPLOGAN:
case astra::FILTER_COSINE:
case astra::FILTER_HAMMING:
case astra::FILTER_HANN:
case astra::FILTER_TUKEY:
case astra::FILTER_LANCZOS:
case astra::FILTER_TRIANGULAR:
case astra::FILTER_GAUSSIAN:
case astra::FILTER_BARTLETTHANN:
case astra::FILTER_BLACKMAN:
case astra::FILTER_NUTTALL:
case astra::FILTER_BLACKMANHARRIS:
case astra::FILTER_BLACKMANNUTTALL:
case astra::FILTER_FLATTOP:
case astra::FILTER_PARZEN:
{
genCuFFTFilter(_cfg, dims.iProjAngles, pHostFilter, iFFTRealDetCount, iFreqBinCount);
uploadComplexArrayToDevice(dims.iProjAngles, iFreqBinCount, pHostFilter, (cufftComplex*)D_filter);
break;
}
case astra::FILTER_PROJECTION:
{
// make sure the offered filter has the correct size
assert(_cfg.m_iCustomFilterWidth == iFreqBinCount);
for(int iFreqBinIndex = 0; iFreqBinIndex < iFreqBinCount; iFreqBinIndex++)
{
float fValue = _cfg.m_pfCustomFilter[iFreqBinIndex];
for(int iProjectionIndex = 0; iProjectionIndex < (int)dims.iProjAngles; iProjectionIndex++)
{
pHostFilter[iFreqBinIndex + iProjectionIndex * iFreqBinCount].x = fValue;
pHostFilter[iFreqBinIndex + iProjectionIndex * iFreqBinCount].y = 0.0f;
}
}
uploadComplexArrayToDevice(dims.iProjAngles, iFreqBinCount, pHostFilter, (cufftComplex*)D_filter);
break;
}
case astra::FILTER_SINOGRAM:
{
// make sure the offered filter has the correct size
assert(_cfg.m_iCustomFilterWidth == iFreqBinCount);
for(int iFreqBinIndex = 0; iFreqBinIndex < iFreqBinCount; iFreqBinIndex++)
{
for(int iProjectionIndex = 0; iProjectionIndex < (int)dims.iProjAngles; iProjectionIndex++)
{
float fValue = _cfg.m_pfCustomFilter[iFreqBinIndex + iProjectionIndex * _cfg.m_iCustomFilterWidth];
pHostFilter[iFreqBinIndex + iProjectionIndex * iFreqBinCount].x = fValue;
pHostFilter[iFreqBinIndex + iProjectionIndex * iFreqBinCount].y = 0.0f;
}
}
uploadComplexArrayToDevice(dims.iProjAngles, iFreqBinCount, pHostFilter, (cufftComplex*)D_filter);
break;
}
case astra::FILTER_RPROJECTION:
{
int iProjectionCount = dims.iProjAngles;
int iRealFilterElementCount = iProjectionCount * iFFTRealDetCount;
float * pfHostRealFilter = new float[iRealFilterElementCount];
memset(pfHostRealFilter, 0, sizeof(float) * iRealFilterElementCount);
int iUsedFilterWidth = min(_cfg.m_iCustomFilterWidth, iFFTRealDetCount);
int iStartFilterIndex = (_cfg.m_iCustomFilterWidth - iUsedFilterWidth) / 2;
int iMaxFilterIndex = iStartFilterIndex + iUsedFilterWidth;
int iFilterShiftSize = _cfg.m_iCustomFilterWidth / 2;
for(int iDetectorIndex = iStartFilterIndex; iDetectorIndex < iMaxFilterIndex; iDetectorIndex++)
{
int iFFTInFilterIndex = (iDetectorIndex + iFFTRealDetCount - iFilterShiftSize) % iFFTRealDetCount;
float fValue = _cfg.m_pfCustomFilter[iDetectorIndex];
for(int iProjectionIndex = 0; iProjectionIndex < (int)dims.iProjAngles; iProjectionIndex++)
{
pfHostRealFilter[iFFTInFilterIndex + iProjectionIndex * iFFTRealDetCount] = fValue;
}
}
float* pfDevRealFilter = NULL;
cudaMalloc((void **)&pfDevRealFilter, sizeof(float) * iRealFilterElementCount); // TODO: check for errors
cudaMemcpy(pfDevRealFilter, pfHostRealFilter, sizeof(float) * iRealFilterElementCount, cudaMemcpyHostToDevice);
delete[] pfHostRealFilter;
runCudaFFT(iProjectionCount, pfDevRealFilter, iFFTRealDetCount, iFFTRealDetCount, iFFTRealDetCount, iFreqBinCount, (cufftComplex*)D_filter);
cudaFree(pfDevRealFilter);
break;
}
case astra::FILTER_RSINOGRAM:
{
int iProjectionCount = dims.iProjAngles;
int iRealFilterElementCount = iProjectionCount * iFFTRealDetCount;
float* pfHostRealFilter = new float[iRealFilterElementCount];
memset(pfHostRealFilter, 0, sizeof(float) * iRealFilterElementCount);
int iUsedFilterWidth = min(_cfg.m_iCustomFilterWidth, iFFTRealDetCount);
int iStartFilterIndex = (_cfg.m_iCustomFilterWidth - iUsedFilterWidth) / 2;
int iMaxFilterIndex = iStartFilterIndex + iUsedFilterWidth;
int iFilterShiftSize = _cfg.m_iCustomFilterWidth / 2;
for(int iDetectorIndex = iStartFilterIndex; iDetectorIndex < iMaxFilterIndex; iDetectorIndex++)
{
int iFFTInFilterIndex = (iDetectorIndex + iFFTRealDetCount - iFilterShiftSize) % iFFTRealDetCount;
for(int iProjectionIndex = 0; iProjectionIndex < (int)dims.iProjAngles; iProjectionIndex++)
{
float fValue = _cfg.m_pfCustomFilter[iDetectorIndex + iProjectionIndex * _cfg.m_iCustomFilterWidth];
pfHostRealFilter[iFFTInFilterIndex + iProjectionIndex * iFFTRealDetCount] = fValue;
}
}
float* pfDevRealFilter = NULL;
cudaMalloc((void **)&pfDevRealFilter, sizeof(float) * iRealFilterElementCount); // TODO: check for errors
cudaMemcpy(pfDevRealFilter, pfHostRealFilter, sizeof(float) * iRealFilterElementCount, cudaMemcpyHostToDevice);
delete[] pfHostRealFilter;
runCudaFFT(iProjectionCount, pfDevRealFilter, iFFTRealDetCount, iFFTRealDetCount, iFFTRealDetCount, iFreqBinCount, (cufftComplex*)D_filter);
cudaFree(pfDevRealFilter);
break;
}
default:
{
ASTRA_ERROR("FBP::setFilter: Unknown filter type requested");
delete [] pHostFilter;
return false;
}
}
delete [] pHostFilter;
return true;
}
bool FBP::iterate(unsigned int iterations)
{
zeroVolumeData(D_volumeData, volumePitch, dims);
bool ok = false;
float fFanDetSize = 0.0f;
if (fanProjs) {
// Call FDK_PreWeight to handle fan beam geometry. We treat
// this as a cone beam setup of a single slice:
// TODO: TOffsets affects this preweighting...
// TODO: We take the fan parameters from the last projection here
// without checking if they're the same in all projections
float *pfAngles = new float[dims.iProjAngles];
float fOriginSource, fOriginDetector, fOffset;
for (unsigned int i = 0; i < dims.iProjAngles; ++i) {
bool ok = astra::getFanParameters(fanProjs[i], dims.iProjDets,
pfAngles[i],
fOriginSource, fOriginDetector,
fFanDetSize, fOffset);
if (!ok) {
ASTRA_ERROR("FBP_CUDA: Failed to extract circular fan beam parameters from fan beam geometry");
return false;
}
}
// We create a fake cudaPitchedPtr
cudaPitchedPtr tmp;
tmp.ptr = D_sinoData;
tmp.pitch = sinoPitch * sizeof(float);
tmp.xsize = dims.iProjDets;
tmp.ysize = dims.iProjAngles;
// and a fake Dimensions3D
astraCUDA3d::SDimensions3D dims3d;
dims3d.iVolX = dims.iVolWidth;
dims3d.iVolY = dims.iVolHeight;
dims3d.iVolZ = 1;
dims3d.iProjAngles = dims.iProjAngles;
dims3d.iProjU = dims.iProjDets;
dims3d.iProjV = 1;
astraCUDA3d::FDK_PreWeight(tmp, fOriginSource,
fOriginDetector, 0.0f,
fFanDetSize, 1.0f, /* fPixelSize */ 1.0f,
m_bShortScan, dims3d, pfAngles);
} else {
// TODO: How should different detector pixel size in different
// projections be handled?
}
if (D_filter) {
int iFFTRealDetCount = calcNextPowerOfTwo(2 * dims.iProjDets);
int iFFTFourDetCount = calcFFTFourierSize(iFFTRealDetCount);
cufftComplex * pDevComplexSinogram = NULL;
allocateComplexOnDevice(dims.iProjAngles, iFFTFourDetCount, &pDevComplexSinogram);
runCudaFFT(dims.iProjAngles, D_sinoData, sinoPitch, dims.iProjDets, iFFTRealDetCount, iFFTFourDetCount, pDevComplexSinogram);
applyFilter(dims.iProjAngles, iFFTFourDetCount, pDevComplexSinogram, (cufftComplex*)D_filter);
runCudaIFFT(dims.iProjAngles, pDevComplexSinogram, D_sinoData, sinoPitch, dims.iProjDets, iFFTRealDetCount, iFFTFourDetCount);
freeComplexOnDevice(pDevComplexSinogram);
}
if (fanProjs) {
float fOutputScale = 1.0 / (/*fPixelSize * fPixelSize * fPixelSize * */ fFanDetSize * fFanDetSize);
ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale);
} else {
// scale by number of angles. For the fan-beam case, this is already
// handled by FDK_PreWeight
float fOutputScale = (M_PI / 2.0f) / (float)dims.iProjAngles;
//fOutputScale /= fDetSize * fDetSize;
ok = BP(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, parProjs, fOutputScale);
}
if(!ok)
{
return false;
}
return true;
}
}