GeoShiftTransformer< T > Class Template Reference

#include <cuda_gpu_geo_shift_transformer.h>

Public Member Functions
	GeoShiftTransformer ()
	~GeoShiftTransformer ()
void	init (const GPU &gpu, size_t x, size_t y, size_t n, int device, myStreamHandle *stream)
void	initLazy (const GPU &gpu, size_t x, size_t y, size_t n, int device, myStreamHandle *stream=NULL)
void	release ()
template<typename T_IN >
void	applyShift (MultidimArray< T > &output, const MultidimArray< T_IN > &input, T shiftX, T shiftY)
void	test ()

Detailed Description

template<typename T>
class GeoShiftTransformer< T >

Definition at line 51 of file cuda_gpu_geo_shift_transformer.h.

Constructor & Destructor Documentation

GeoShiftTransformer()

template<typename T>

GeoShiftTransformer< T >::GeoShiftTransformer ( )

inline

Constructor

Definition at line 55 of file cuda_gpu_geo_shift_transformer.h.

                          :
        device(-1), imgs(NULL), ffts(NULL), stream(NULL), isReady(false) {
    };

~GeoShiftTransformer()

template<typename T>

GeoShiftTransformer< T >::~GeoShiftTransformer ( )

inline

Definition at line 59 of file cuda_gpu_geo_shift_transformer.h.

                           {
        release();
    }

Member Function Documentation

applyShift()

template<typename T >

template<typename T_IN >

void GeoShiftTransformer< T >::applyShift	(	MultidimArray< T > &	output,
		const MultidimArray< T_IN > &	input,
		T	shiftX,
		T	shiftY
	)

Apply 2D shift. Image will be repeated at the border.

Parameters

output	where resulting images will be stored. Does not have to be initialized
input	to process
shiftX	to apply
shiftY	to apply

Definition at line 87 of file cuda_gpu_geo_shift_transformer.cpp.

                                                              {
    checkRestrictions(output, input);
    if (output.xdim == 0) {
        output.resizeNoCopy(input.ndim, input.zdim, input.ydim, input.xdim);
    }
    if (shiftX == 0 && shiftY == 0) {
        typeCast(input, output);
        return;
    }

    GpuMultidimArrayAtGpu<std::complex<float> > mask;

    MultidimArray<T> tmp;
    typeCast(input, tmp);
    imgs->copyToGpu(tmp.data);
    if(stream) {
        imgs->fftStream(*ffts, fftHandle, *stream, false, mask);
    } else {
        imgs->fft(*ffts, fftHandle);
    }

    dim3 dimBlock(BLOCK_DIM_X, BLOCK_DIM_X);
    dim3 dimGrid(ceil(ffts->Xdim / (T) dimBlock.x), ceil(ffts->Ydim / (T) dimBlock.y));

    if (stream) {
        shiftFFT2D<true><<<dimGrid, dimBlock, 0, *(cudaStream_t*)stream->ptr>>>(
            (float2*)ffts->d_data, ffts->Ndim, ffts->Xdim, imgs->Xdim,
            imgs->Ydim, shiftX, shiftY);
    } else {
        shiftFFT2D<true><<<dimGrid, dimBlock>>>((float2*)ffts->d_data,
            ffts->Ndim, ffts->Xdim, imgs->Xdim, imgs->Ydim, shiftX, shiftY);
    }
    gpuErrchk(cudaPeekAtLastError());

    if (stream) {
        ffts->ifftStream(*imgs, ifftHandle, *stream, false, mask);
    } else {
        ffts->ifft(*imgs, ifftHandle);
    }
    imgs->copyToCpu(output.data);
}

init()

template<typename T >

void GeoShiftTransformer< T >::init	(	const GPU &	gpu,
		size_t	x,
		size_t	y,
		size_t	n,
		int	device,
		myStreamHandle *	stream
	)

Release previously obtained resources and initialize the transformer for processing images of given size. It also allocates all resources on GPU.

Parameters

gpu	to use
x	dim (inner-most) of the resulting image
y	dim (outer-most) of the resulting image
n	no. of images to process in a single batch
device	to be used
stream	to be used. NULL for default

Definition at line 53 of file cuda_gpu_geo_shift_transformer.cpp.

                                {
    release();

    this->device = device;
    this->stream = stream;

    m_gpu = &gpu; // FIXME DS implement set() check
    this->imgs = new GpuMultidimArrayAtGpu<T>(x, y, 1, n);
    this->ffts = new GpuMultidimArrayAtGpu<std::complex<T> >(x / 2 + 1, y, 1, n);

    fftHandle.ptr = new cufftHandle;
    ifftHandle.ptr = new cufftHandle;
    if (this->stream) {
        createPlanFFTStream(x, y, n, 1, true, (cufftHandle*)fftHandle.ptr, *stream);
        createPlanFFTStream(x, y, n, 1, false, (cufftHandle*)ifftHandle.ptr, *stream);
    } else {
        createPlanFFT(x, y, n, 1, true, (cufftHandle*)fftHandle.ptr);
        createPlanFFT(x, y, n, 1, false, (cufftHandle*)ifftHandle.ptr);
    }

    isReady = true;
}

initLazy()

template<typename T >

void GeoShiftTransformer< T >::initLazy	(	const GPU &	gpu,
		size_t	x,
		size_t	y,
		size_t	n,
		int	device,
		myStreamHandle *	stream = NULL
	)

Similar as init(), except this method has no effect should the instance be already initialized. It is useful for example in a for loop, where first call will initialize resources and following calls will be ignored Do NOT use it for reinitialization.

Definition at line 78 of file cuda_gpu_geo_shift_transformer.cpp.

                                {
    if (!isReady) {
        init(gpu, x, y, n, device, stream);
    }
}

release()

template<typename T >

void GeoShiftTransformer< T >::release

(

)

Release all resources hold by this instance

Definition at line 34 of file cuda_gpu_geo_shift_transformer.cpp.

                                     {
    delete imgs;
    imgs = nullptr;
    delete ffts;
    ffts = nullptr;

    fftHandle.clear();
    ifftHandle.clear();

    // do not release stream, it has been passed by pointer, so we don't own it
    stream = nullptr;
    device = -1;

    m_gpu = nullptr;

    isReady = false;
}

test()

template<typename T >

void GeoShiftTransformer< T >::test

(

)

Definition at line 131 of file cuda_gpu_geo_shift_transformer.cpp.

                                  {
    int offsetX = 5;
    int offsetY = -7;
    size_t xSize = 2048;
    size_t ySize = 3072;
    MultidimArray<float> resGpu(ySize, xSize);
    MultidimArray<float> expected(ySize, xSize);
    MultidimArray<float> input(ySize, xSize);
    for (int y = 10; y < 15; ++y) {
        for (int x = 10; x < 15; ++x) {
            int indexIn = (y * input.xdim) + x;
            int indexExp = ((y + offsetY) * input.xdim) + (x + offsetX);
            input.data[indexIn] = 10;
            expected.data[indexExp] = 10;
        }
    }

    GeoShiftTransformer<float> tr;
    auto gpu = GPU();
    tr.initLazy(gpu, input.xdim, input.ydim, 1, 0);
    tr.applyShift(resGpu, input, offsetX, offsetY);

    bool error = false;
    for (int y = 0; y < expected.ydim; ++y) {
        for (int x = 0; x < expected.xdim; ++x) {
            int index = (y * input.xdim) + x;
            float gpu = resGpu.data[index];
            float cpu = expected.data[index];
            float threshold = std::max(gpu, cpu) / 1000.f;
            float diff = std::abs(cpu - gpu);
            if (diff > threshold && diff > 0.001) {
                error = true;
                printf("%d gpu %.4f cpu %.4f (%f > %f)\n", index, gpu, cpu, diff, threshold);
            }
        }
    }
#ifdef DEBUG
    Image<float> img(expected.xdim, expected.ydim);
    img.data = expected;
    img.write("expected.vol");
    img.data = resGpu;
    img.write("resGpu.vol");
#endif
    printf("\n SHIFT %s\n", error ? "FAILED" : "OK");
}

---

The documentation for this class was generated from the following files:

• xmipp/libraries/reconstruction_cuda/cuda_gpu_geo_shift_transformer.h
• xmipp/libraries/reconstruction_cuda/cuda_gpu_geo_shift_transformer.cpp