diff options
Diffstat (limited to 'cuda/2d/fft.cu')
-rw-r--r-- | cuda/2d/fft.cu | 207 |
1 files changed, 0 insertions, 207 deletions
diff --git a/cuda/2d/fft.cu b/cuda/2d/fft.cu index 2e94b79..8361ad2 100644 --- a/cuda/2d/fft.cu +++ b/cuda/2d/fft.cu @@ -314,210 +314,3 @@ void genCuFFTFilter(const SFilterConfig &_cfg, int _iProjectionCount, } - - -#ifdef STANDALONE - -__global__ static void doubleFourierOutput_kernel(int _iProjectionCount, - int _iDetectorCount, - cufftComplex* _pFourierOutput) -{ - int iIndex = threadIdx.x + blockIdx.x * blockDim.x; - int iProjectionIndex = iIndex / _iDetectorCount; - int iDetectorIndex = iIndex % _iDetectorCount; - - if(iProjectionIndex >= _iProjectionCount) - { - return; - } - - if(iDetectorIndex <= (_iDetectorCount / 2)) - { - return; - } - - int iOtherDetectorIndex = _iDetectorCount - iDetectorIndex; - - _pFourierOutput[iProjectionIndex * _iDetectorCount + iDetectorIndex].x = _pFourierOutput[iProjectionIndex * _iDetectorCount + iOtherDetectorIndex].x; - _pFourierOutput[iProjectionIndex * _iDetectorCount + iDetectorIndex].y = -_pFourierOutput[iProjectionIndex * _iDetectorCount + iOtherDetectorIndex].y; -} - -static void doubleFourierOutput(int _iProjectionCount, int _iDetectorCount, - cufftComplex * _pFourierOutput) -{ - const int iBlockSize = 256; - int iElementCount = _iProjectionCount * _iDetectorCount; - int iBlockCount = (iElementCount + iBlockSize - 1) / iBlockSize; - - doubleFourierOutput_kernel<<< iBlockCount, iBlockSize >>>(_iProjectionCount, - _iDetectorCount, - _pFourierOutput); - CHECK_ERROR("doubleFourierOutput_kernel failed"); -} - - - -static void writeToMatlabFile(const char * _fileName, const float * _pfData, - int _iRowCount, int _iColumnCount) -{ - std::ofstream out(_fileName); - - for(int iRowIndex = 0; iRowIndex < _iRowCount; iRowIndex++) - { - for(int iColumnIndex = 0; iColumnIndex < _iColumnCount; iColumnIndex++) - { - out << _pfData[iColumnIndex + iRowIndex * _iColumnCount] << " "; - } - - out << std::endl; - } -} - -static void convertComplexToRealImg(const cufftComplex * _pComplex, - int _iElementCount, - float * _pfReal, float * _pfImaginary) -{ - for(int iIndex = 0; iIndex < _iElementCount; iIndex++) - { - _pfReal[iIndex] = _pComplex[iIndex].x; - _pfImaginary[iIndex] = _pComplex[iIndex].y; - } -} - -void testCudaFFT() -{ - const int iProjectionCount = 2; - const int iDetectorCount = 1024; - const int iTotalElementCount = iProjectionCount * iDetectorCount; - - float * pfHostProj = new float[iTotalElementCount]; - memset(pfHostProj, 0, sizeof(float) * iTotalElementCount); - - for(int iProjectionIndex = 0; iProjectionIndex < iProjectionCount; iProjectionIndex++) - { - for(int iDetectorIndex = 0; iDetectorIndex < iDetectorCount; iDetectorIndex++) - { -// int - -// pfHostProj[iIndex] = (float)rand() / (float)RAND_MAX; - } - } - - writeToMatlabFile("proj.mat", pfHostProj, iProjectionCount, iDetectorCount); - - float * pfDevProj = NULL; - SAFE_CALL(cudaMalloc((void **)&pfDevProj, sizeof(float) * iTotalElementCount)); - SAFE_CALL(cudaMemcpy(pfDevProj, pfHostProj, sizeof(float) * iTotalElementCount, cudaMemcpyHostToDevice)); - - cufftComplex * pDevFourProj = NULL; - SAFE_CALL(cudaMalloc((void **)&pDevFourProj, sizeof(cufftComplex) * iTotalElementCount)); - - cufftHandle plan; - cufftResult result; - - result = cufftPlan1d(&plan, iDetectorCount, CUFFT_R2C, iProjectionCount); - if(result != CUFFT_SUCCESS) - { - ASTRA_ERROR("Failed to plan 1d r2c fft"); - } - - result = cufftExecR2C(plan, pfDevProj, pDevFourProj); - if(result != CUFFT_SUCCESS) - { - ASTRA_ERROR("Failed to exec 1d r2c fft"); - } - - cufftDestroy(plan); - - doubleFourierOutput(iProjectionCount, iDetectorCount, pDevFourProj); - - cufftComplex * pHostFourProj = new cufftComplex[iTotalElementCount]; - SAFE_CALL(cudaMemcpy(pHostFourProj, pDevFourProj, sizeof(cufftComplex) * iTotalElementCount, cudaMemcpyDeviceToHost)); - - float * pfHostFourProjReal = new float[iTotalElementCount]; - float * pfHostFourProjImaginary = new float[iTotalElementCount]; - - convertComplexToRealImg(pHostFourProj, iTotalElementCount, pfHostFourProjReal, pfHostFourProjImaginary); - - writeToMatlabFile("proj_four_real.mat", pfHostFourProjReal, iProjectionCount, iDetectorCount); - writeToMatlabFile("proj_four_imaginary.mat", pfHostFourProjImaginary, iProjectionCount, iDetectorCount); - - float * pfDevInFourProj = NULL; - SAFE_CALL(cudaMalloc((void **)&pfDevInFourProj, sizeof(float) * iTotalElementCount)); - - result = cufftPlan1d(&plan, iDetectorCount, CUFFT_C2R, iProjectionCount); - if(result != CUFFT_SUCCESS) - { - ASTRA_ERROR("Failed to plan 1d c2r fft"); - } - - result = cufftExecC2R(plan, pDevFourProj, pfDevInFourProj); - if(result != CUFFT_SUCCESS) - { - ASTRA_ERROR("Failed to exec 1d c2r fft"); - } - - cufftDestroy(plan); - - rescaleInverseFourier(iProjectionCount, iDetectorCount, pfDevInFourProj); - - float * pfHostInFourProj = new float[iTotalElementCount]; - SAFE_CALL(cudaMemcpy(pfHostInFourProj, pfDevInFourProj, sizeof(float) * iTotalElementCount, cudaMemcpyDeviceToHost)); - - writeToMatlabFile("in_four.mat", pfHostInFourProj, iProjectionCount, iDetectorCount); - - SAFE_CALL(cudaFree(pDevFourProj)); - SAFE_CALL(cudaFree(pfDevProj)); - - delete [] pfHostInFourProj; - delete [] pfHostFourProjReal; - delete [] pfHostFourProjImaginary; - delete [] pfHostProj; - delete [] pHostFourProj; -} - -void downloadDebugFilterComplex(float * _pfHostSinogram, int _iProjectionCount, - int _iDetectorCount, - cufftComplex * _pDevFilter, - int _iFilterDetCount) -{ - cufftComplex * pHostFilter = NULL; - size_t complMemSize = sizeof(cufftComplex) * _iFilterDetCount * _iProjectionCount; - pHostFilter = (cufftComplex *)malloc(complMemSize); - SAFE_CALL(cudaMemcpy(pHostFilter, _pDevFilter, complMemSize, cudaMemcpyDeviceToHost)); - - for(int iTargetProjIndex = 0; iTargetProjIndex < _iProjectionCount; iTargetProjIndex++) - { - for(int iTargetDetIndex = 0; iTargetDetIndex < min(_iDetectorCount, _iFilterDetCount); iTargetDetIndex++) - { - cufftComplex source = pHostFilter[iTargetDetIndex + iTargetProjIndex * _iFilterDetCount]; - float fReadValue = sqrtf(source.x * source.x + source.y * source.y); - _pfHostSinogram[iTargetDetIndex + iTargetProjIndex * _iDetectorCount] = fReadValue; - } - } - - free(pHostFilter); -} - -void downloadDebugFilterReal(float * _pfHostSinogram, int _iProjectionCount, - int _iDetectorCount, float * _pfDevFilter, - int _iFilterDetCount) -{ - float * pfHostFilter = NULL; - size_t memSize = sizeof(float) * _iFilterDetCount * _iProjectionCount; - pfHostFilter = (float *)malloc(memSize); - SAFE_CALL(cudaMemcpy(pfHostFilter, _pfDevFilter, memSize, cudaMemcpyDeviceToHost)); - - for(int iTargetProjIndex = 0; iTargetProjIndex < _iProjectionCount; iTargetProjIndex++) - { - for(int iTargetDetIndex = 0; iTargetDetIndex < min(_iDetectorCount, _iFilterDetCount); iTargetDetIndex++) - { - float fSource = pfHostFilter[iTargetDetIndex + iTargetProjIndex * _iFilterDetCount]; - _pfHostSinogram[iTargetDetIndex + iTargetProjIndex * _iDetectorCount] = fSource; - } - } - - free(pfHostFilter); -} - -#endif |