marwan-abdellah/cufftshift

CUDA-based implementation for linear 1D, 2D and 3D FFT-Shift functions.

cufftshift's Issues

cufftShift_2D_impl() : cufftShift_2D_IP_impl.cu

Hi,

Thanks for providing your FFT Shift implementation!

I am using your software to perform 2D FFT shifts on the NVIDIA TX1. I was originally using the out-of-place version (cufftShift_2D_OP_impl.cu). But I found the out-of-place version caused the NVIDIA TX1 to periodically hang. Best I could tell there was a CPU/GPU synchronization which was hanging the GPU.

Having been unsuccessful in finding the root cause, I switched to the in-place FFT shift and found it did not perform the FFT shift correctly. Only two of the quadrants were shifted (as opposed to all four).

As best I can tell, the following line of code:
kernelConf* conf = cufftShift::GenAutoConf_2D(N/2);
needs to be changed to:
kernelConf* conf = cufftShift::GenAutoConf_2D(N);

I came to this conclusion by comparing the in-place version to the out-of-place version (which works correctly) and noticed the out-of-place version used N instead of N/2.

So in summary, the in-place version now correctly FFT shifts the image and the NVIDIA TX1 GPU does not hang.

Just thought I'd pass on this finding in case others run into a similar issue.

Thanks again for providing this implementation!

Thank you.

That's Cool ~

Is there any code for cuFFT + 2D Phase Correlation ?

Thanks.

A typo & An extension

Hi Marwan,

Your cufftShift implementation is really a impressive job. Here I just have some thoughts on your codes.
1-I think there's a typo in the configGPU.cpp, in the GenAutoConf_1D function that

// 1024 should be changed to 512, otherwise number of blocks may become 0
 if(N >= 512)    threadPerBlock_X = 1024; 

2- I think it's easy to extend the code to work not only for square matrix but also for rectangle ones with even elements on each dimension. Take the 2D kernel for example,

template <typename T>
__global__
void cufftShift_2D_kernel(T* data, int width, int height)

{
    // 2D Slice & 1D Line
    int sSlice = width * height;

    // Transformations Equations
    int sEq1 = (sSlice + width) / 2;
    int sEq2 = (sSlice - width) / 2;

    // Thread Index (1D)
    int xThreadIdx = threadIdx.x;
    int yThreadIdx = threadIdx.y;

    // Block Width & Height
    int blockWidth = blockDim.x;
    int blockHeight = blockDim.y;

    // Thread Index (2D)
    int xIndex = blockIdx.x * blockWidth + xThreadIdx;
    int yIndex = blockIdx.y * blockHeight + yThreadIdx;

    // Thread Index Converted into 1D Index
    int index = (yIndex * width) + xIndex;

    T regTemp;

    if (xIndex < width / 2&&yIndex<height)
    {

        regTemp = data[index];

        // First Quad
        data[index] = data[index + sEq1];

        // Third Quad
        data[index + sEq1] = regTemp;       
    }
    else
    {
        if (yIndex < height / 2&&xIndex>=width/2&&xIndex<width)
        {
            regTemp = data[index];

            // Second Quad
            data[index] = data[index + sEq2];

            // Fourth Quad
            data[index + sEq2] = regTemp;
        }
    }
}

Thanks for your codes again!

Best,
Teller

.