diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2016-07-29 12:03:38 +0200 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2016-07-29 12:11:52 +0200 |
| commit | b1ffc11d930c19bd73af9837a08bc8dde9fe8e72 (patch) | |
| tree | 40ce622ffdc8d75c885135db1ce4773c9670e2c3 /cuda/2d/fft.cu | |
| parent | b2611a03577c285ddf48edab0d05dafa09ab362c (diff) | |
| download | astra-b1ffc11d930c19bd73af9837a08bc8dde9fe8e72.tar.gz astra-b1ffc11d930c19bd73af9837a08bc8dde9fe8e72.tar.bz2 astra-b1ffc11d930c19bd73af9837a08bc8dde9fe8e72.tar.xz astra-b1ffc11d930c19bd73af9837a08bc8dde9fe8e72.zip | |
Add CUDA parvec support
Diffstat (limited to 'cuda/2d/fft.cu')
| -rw-r--r-- | cuda/2d/fft.cu | 12 |
1 files changed, 9 insertions, 3 deletions
diff --git a/cuda/2d/fft.cu b/cuda/2d/fft.cu index 2d259a9..8a7cc10 100644 --- a/cuda/2d/fft.cu +++ b/cuda/2d/fft.cu @@ -64,6 +64,8 @@ using namespace astra; } } while (0) +namespace astraCUDA { + __global__ static void applyFilter_kernel(int _iProjectionCount, int _iFreqBinCount, cufftComplex * _pSinogram, @@ -276,11 +278,11 @@ bool runCudaIFFT(int _iProjectionCount, const cufftComplex* _pDevSourceComplex, // Because the input is real, the Fourier transform is symmetric. // CUFFT only outputs the first half (ignoring the redundant second half), // and expects the same as input for the IFFT. -int calcFFTFourSize(int _iFFTRealSize) +int calcFFTFourierSize(int _iFFTRealSize) { - int iFFTFourSize = _iFFTRealSize / 2 + 1; + int iFFTFourierSize = _iFFTRealSize / 2 + 1; - return iFFTFourSize; + return iFFTFourierSize; } void genIdenFilter(int _iProjectionCount, cufftComplex * _pFilter, @@ -695,6 +697,10 @@ void genFilter(E_FBPFILTER _eFilter, float _fD, int _iProjectionCount, delete[] pfW; } + +} + + #ifdef STANDALONE __global__ static void doubleFourierOutput_kernel(int _iProjectionCount, |