GeoTransformer< T > Class Template Reference

#include <cuda_gpu_geo_transformer.h>

Inheritance diagram for GeoTransformer< T >:

Public Member Functions
	GeoTransformer ()
	~GeoTransformer ()
void	initForMatrix (size_t x, size_t y, size_t z)
void	initLazyForMatrix (size_t x, size_t y=1, size_t z=1)
void	initForBSpline (size_t inX, size_t inY, size_t inN, size_t splineX, size_t splineY, size_t splineN, const GPU &gpu)
void	initLazyForBSpline (size_t inX, size_t inY, size_t inN, size_t splineX, size_t splineY, size_t splineN, const GPU &gpu)
void	release ()
template<typename T_MAT >
void	applyGeometry (int splineDegree, MultidimArray< T > &output, const MultidimArray< T > &input, const Matrix2D< T_MAT > &transform, bool isInv, bool wrap, T outside=0, const MultidimArray< T > *bCoeffsPtr=NULL)
void	applyBSplineTransform (int splineDegree, MultidimArray< T > &output, const MultidimArray< T > &input, const std::pair< Matrix1D< T >, Matrix1D< T >> &coeffs, size_t imageIdx, T outside=0)
void	test ()

Protected Member Functions
void	applyBSplineTransformRef (int splineDegree, MultidimArray< T > &output, const MultidimArray< T > &input, const std::pair< Matrix1D< T >, Matrix1D< T >> &coeffs, size_t imageIdx, T outside=0)
void	produceAndLoadCoeffs (const MultidimArray< T > &input)
std::unique_ptr< T[]>	copy_out_d_in (size_t size) const

Detailed Description

template<typename T>
class GeoTransformer< T >

Definition at line 55 of file cuda_gpu_geo_transformer.h.

Constructor & Destructor Documentation

GeoTransformer()

template<typename T>

GeoTransformer< T >::GeoTransformer ( )

inline

Constructor

Definition at line 59 of file cuda_gpu_geo_transformer.h.

{ setDefaultValues(); };

~GeoTransformer()

template<typename T>

GeoTransformer< T >::~GeoTransformer ( )

inline

Definition at line 61 of file cuda_gpu_geo_transformer.h.

                      {
        release();
    }

Member Function Documentation

applyBSplineTransform()

template<typename T >

void GeoTransformer< T >::applyBSplineTransform	(	int	splineDegree,
		MultidimArray< T > &	output,
		const MultidimArray< T > &	input,
		const std::pair< Matrix1D< T >, Matrix1D< T >> &	coeffs,
		size_t	imageIdx,
		T	outside = 0
	)

Apply local transformation defined by a BSpline

Parameters

splineDegree	used for interpolation
output	where resulting image will be stored
input	to process
coeffs	for the X and Y dimension of the input
imageIdx	index of the current image. This function assumes that multiple calls will be done and that interpolation is done also over time
outside	value of the output, where the interpolation does not store anything

Definition at line 187 of file cuda_gpu_geo_transformer.cpp.

                                                                                     {
    checkRestrictions(3, output, input, coeffs, imageIdx);
    auto stream = *(cudaStream_t*)gpu->stream();

    setOutputSize(output);
    if ( splineDegree > 1 ) {
        produceAndLoadCoeffs(input);
    }
    else
    {
        gpuErrchk(cudaMemcpyAsync(d_in, input.data, input.yxdim * sizeof(T), cudaMemcpyHostToDevice, stream));
    }

    loadCoefficients(coeffs.first, coeffs.second);

    dim3 dimBlock(16, 16);
    dim3 dimGrid(ceil(inX / (T) dimBlock.x), ceil((inY / (T) dimBlock.y) / (T) pixelsPerThread)); //more pixels

    // take into account end points
    T hX = (splineX == 3) ? inX : (inX / (T) ((splineX - 3)));
    T hY = (splineY == 3) ? inY : (inY / (T) ((splineY - 3)));
    T hT = (splineN == 3) ? inN : (inN / (T) ((splineN - 3)));
    T tPos = imageIdx / hT;

    switch (splineDegree) {
    case 1:
        applyLocalShiftGeometryKernelMorePixels<T, 1, pixelsPerThread><<<dimGrid, dimBlock, 0, stream>>>(d_coeffsX, d_coeffsY,
                                                                                              d_out, (int)inX, (int)inY, (int)inN,
                                                                                              d_in, imageIdx, (int)splineX, (int)splineY, (int)splineN,
                                                                                              hX, hY, tPos);
        gpuErrchk(cudaPeekAtLastError());
        break;
    case 3:
        applyLocalShiftGeometryKernelMorePixels<T, 3, pixelsPerThread><<<dimGrid, dimBlock, 0, stream>>>(d_coeffsX, d_coeffsY,
                                                                                              d_out, (int)inX, (int)inY, (int)inN,
                                                                                              d_in, imageIdx, (int)splineX, (int)splineY, (int)splineN,
                                                                                              hX, hY, tPos);
        gpuErrchk(cudaPeekAtLastError());
        break;
    default:
        REPORT_ERROR(ERR_NOT_IMPLEMENTED, formatString("applyBSplineTransform not implemented for spline degree %d.", splineDegree));
    }

    gpuErrchk(
        cudaMemcpyAsync(output.data, d_out, output.zyxdim * sizeof(T),
                   cudaMemcpyDeviceToHost, stream));
}

applyBSplineTransformRef()

template<typename T >

void GeoTransformer< T >::applyBSplineTransformRef ( int  splineDegree, MultidimArray< T > &  output, const MultidimArray< T > &  input, const std::pair< Matrix1D< T >, Matrix1D< T >> &  coeffs, size_t  imageIdx, T  outside = 0  )

protected

Definition at line 155 of file cuda_gpu_geo_transformer.cpp.

                                                                                     {
    checkRestrictions(3, output, input, coeffs, imageIdx);

    loadOutput(output, outside);
    produceAndLoadCoeffs(input);

    loadCoefficients(coeffs.first, coeffs.second);

    dim3 dimBlock(BLOCK_DIM_X, BLOCK_DIM_X);
    dim3 dimGrid(ceil(inX / (T) dimBlock.x), ceil(inY / (T) dimBlock.y));

    switch (splineDegree) {
    case 3:
        applyLocalShiftGeometryKernel<T, 3><<<dimGrid, dimBlock>>>(d_coeffsX, d_coeffsY,
                d_out, (int)inX, (int)inY, (int)inN,
                d_in, imageIdx, (int)splineX, (int)splineY, (int)splineN);
            gpuErrchk(cudaPeekAtLastError());
        break;
    default:
        throw std::logic_error("not implemented");
    }

    gpuErrchk(
            cudaMemcpy(output.data, d_out, output.zyxdim * sizeof(T),
                    cudaMemcpyDeviceToHost));

}

applyGeometry()

template<typename T >

template<typename T_MAT >

void GeoTransformer< T >::applyGeometry	(	int	splineDegree,
		MultidimArray< T > &	output,
		const MultidimArray< T > &	input,
		const Matrix2D< T_MAT > &	transform,
		bool	isInv,
		bool	wrap,
		T	outside = 0,
		const MultidimArray< T > *	bCoeffsPtr = NULL
	)

Apply geometry transformation Currently supported transformations: FIXME

• 2D shift using cubic interpolation and mirrorOffBound

  Parameters

  splineDegree	used for interpolation
  output	where resulting image will be stored
  input	to process
  transform	to apply
  isInv	if true, 'transform' is expected to be from the resulting image to source image
  wrap	true to wrap after boundaris
  outside	value to be used when reading outside of the image and wrap == false
  bCoeffsPtr	spline coefficients to use

Definition at line 240 of file cuda_gpu_geo_transformer.cpp.

                                            {
    checkRestrictions(splineDegree, output, input, transform);
    if (transform.isIdentity()) {
        typeCast(input, output);
        return;
    }

    loadTransform(transform, isInv);
    loadOutput(output, outside);

    if (splineDegree > 1) {
        if (NULL != bCoeffsPtr) {
            loadInput(*bCoeffsPtr);
        } else {
            produceAndLoadCoeffs(input);
        }
    } else {
        loadInput(input);
    }

    if (input.getDim() == 2) {
        if (wrap) {
            applyGeometry_2D_wrap(splineDegree);
        } else {
            throw std::logic_error("Not implemented yet");
        }
    } else {
        throw std::logic_error("Not implemented yet");
    }

    gpuErrchk(
            cudaMemcpy(output.data, d_out, output.zyxdim * sizeof(T),
                    cudaMemcpyDeviceToHost));
}

copy_out_d_in()

template<typename T >

std::unique_ptr< T[]> GeoTransformer< T >::copy_out_d_in ( size_t  size ) const

protected

Definition at line 429 of file cuda_gpu_geo_transformer.cpp.

                                                                     {
        auto copy_in = std::unique_ptr<T[]>(new T[size]);

        cudaMemcpy(copy_in.get(), d_in , sizeof(T) * size, cudaMemcpyDeviceToHost);

        return copy_in;
    }

initForBSpline()

template<typename T >

void GeoTransformer< T >::initForBSpline	(	size_t	inX,
		size_t	inY,
		size_t	inN,
		size_t	splineX,
		size_t	splineY,
		size_t	splineN,
		const GPU &	gpu
	)

Release previously obtained resources and initialize the transformer for processing images using BSpline coefficients. It also allocates all resources on GPU.

Parameters

sizes	of the input images and number of images to be processed
number	of BSpline control points, including end points

Definition at line 71 of file cuda_gpu_geo_transformer.cpp.

{
    release();

    this->inX = inX;
    this->inY = inY;
    this->inZ = 1;
    this->inN = inN;
    this->splineX = splineX;
    this->splineY = splineY;
    this->splineN = splineN;
    // take into account end control points

    // padding for produceAndLoadCoeffs; Y dimension has to be a multiple of BLOCK_SIZE
    const int BLOCK_SIZE = iirConvolve2D_Cardinal_BSpline_3_MirrorOffBoundKernels::BLOCK_SIZE;
    const int Y_padded = (inY / BLOCK_SIZE) * BLOCK_SIZE + BLOCK_SIZE * (inY % BLOCK_SIZE != 0);

    size_t inOutSize = inX * inY;
    size_t inOutSize_padded = inX* Y_padded;
    size_t coeffsSize = splineX * splineY * splineN;
    gpuErrchk(cudaMalloc((void** ) &d_coeffsX, coeffsSize * sizeof(T)));
    gpuErrchk(cudaMalloc((void** ) &d_coeffsY, coeffsSize * sizeof(T)));
    gpuErrchk(cudaMalloc((void** ) &d_in, inOutSize_padded * sizeof(T)));
    gpuErrchk(cudaMalloc((void** ) &d_out, inOutSize * sizeof(T)));

    this->gpu = &gpu;

    isReadyForBspline = true;
}

initForMatrix()

template<typename T >

void GeoTransformer< T >::initForMatrix	(	size_t	x,
		size_t	y,
		size_t	z
	)

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

Parameters

x	dim (inner-most) of the resulting image
y	dim of the resulting image
z	dim (outer-most) of the resulting image

Definition at line 49 of file cuda_gpu_geo_transformer.cpp.

                                                                  {
    release();

    inX = x;
    inY = y;
    inZ = z;
    size_t matSize = (0 == z) ? 9 : 16;
    gpuErrchk(cudaMalloc((void** ) &d_trInv, matSize * sizeof(T)));
    gpuErrchk(cudaMalloc((void** ) &d_in, x * y * z * sizeof(T)));
    gpuErrchk(cudaMalloc((void** ) &d_out, x * y * z * sizeof(T)));

    isReadyForMatrix = true;
}

initLazyForBSpline()

template<typename T >

void GeoTransformer< T >::initLazyForBSpline	(	size_t	inX,
		size_t	inY,
		size_t	inN,
		size_t	splineX,
		size_t	splineY,
		size_t	splineN,
		const GPU &	gpu
	)

Similar as the other init() function, 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

Definition at line 103 of file cuda_gpu_geo_transformer.cpp.

{
    if (!isReadyForBspline)
    {
        initForBSpline(inX, inY, inZ, splineX, splineY, splineN, gpu);
    }
}

initLazyForMatrix()

template<typename T >

void GeoTransformer< T >::initLazyForMatrix	(	size_t	x,
		size_t	y = 1,
		size_t	z = 1
	)

Similar as other init() function, 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

Definition at line 64 of file cuda_gpu_geo_transformer.cpp.

                                                                      {
    if (!isReadyForMatrix) {
        initForMatrix(x, y, z);
    }
}

produceAndLoadCoeffs()

template<typename T >

void GeoTransformer< T >::produceAndLoadCoeffs ( const MultidimArray< T > &  input )

protected

Computes spline coefficients of the image and load them to GPU

Parameters

splineDegree	to be used
input	image used to generate the coefficients

Definition at line 303 of file cuda_gpu_geo_transformer.cpp.

{
    gpuErrchk(
        cudaMemcpyAsync(d_in, input.data, input.yxdim * sizeof(T), cudaMemcpyHostToDevice, *(cudaStream_t*)gpu->stream()));

    iirConvolve2D_Cardinal_Bspline_3_MirrorOffBoundInplace(d_in, input.xdim, input.ydim, *(cudaStream_t*)gpu->stream());
}

release()

template<typename T >

void GeoTransformer< T >::release

(

)

Release all resources hold by this instance

Definition at line 32 of file cuda_gpu_geo_transformer.cpp.

                                {
    cudaFree(d_in);
    cudaFree(d_out);
    cudaFree(d_trInv);
    cudaFree(d_coeffsX);
    cudaFree(d_coeffsY);
    setDefaultValues();
}

test()

template<typename T >

void GeoTransformer< T >::test

(

)

Definition at line 113 of file cuda_gpu_geo_transformer.cpp.

                             {
    Matrix1D<T> shift(2);
    shift.vdata[0] = 0.45;
    shift.vdata[1] = 0.62;
    Matrix2D<T> transform;
    translation2DMatrix(shift, transform, true);
    test(transform);
}

---

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

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