Alignment::CudaRotPolarEstimator< T > Class Template Reference

#include <cuda_rot_polar_estimator.h>

Inheritance diagram for Alignment::CudaRotPolarEstimator< T >:

Collaboration diagram for Alignment::CudaRotPolarEstimator< T >:

Public Member Functions
	CudaRotPolarEstimator ()
virtual	~CudaRotPolarEstimator ()
	CudaRotPolarEstimator (CudaRotPolarEstimator &o)=delete
CudaRotPolarEstimator &	operator= (const CudaRotPolarEstimator &other)=delete
	CudaRotPolarEstimator (CudaRotPolarEstimator &&o)
CudaRotPolarEstimator const &	operator= (CudaRotPolarEstimator &&o)=delete
template<bool FULL_CIRCLE>
void	sComputePolarTransform (const GPU &gpu, const Dimensions &dimIn, T *__restrict__ d_in, const Dimensions &dimOut, T *__restrict__ d_out, int posOfFirstRing)
template<bool FULL_CIRCLE>
void	sNormalize (const GPU &gpu, const Dimensions &dim, T *__restrict__ d_in, T *__restrict__ d_sums, T *__restrict__ d_sumsSqr, int posOfFirstRing)
Public Member Functions inherited from Alignment::ARotationEstimator< T >
	ARotationEstimator ()
void	init (const RotationEstimationSetting settings, bool reuse)
void	loadReference (const T *ref)
void	compute (T *others)
const std::vector< float > &	getRotations2D () const
virtual	~ARotationEstimator ()
HW &	getHW () const
const RotationEstimationSetting &	getSettings () const

Static Public Member Functions
static void	sComputeCorrelationsOneToN (const GPU &gpu, const std::complex< T > *d_in, std::complex< T > *d_out, const std::complex< T > *d_ref, const Dimensions &dims, int firstRingRadius)
template<bool FULL_CIRCLE>
static void	sComputePolarTransform (const GPU &gpu, const Dimensions &dimIn, T *d_in, const Dimensions &dimOut, T *d_out, int posOfFirstRing)
template<bool FULL_CIRCLE>
static void	sNormalize (const GPU &gpu, const Dimensions &dimIn, T *d_in, T *d_1, T *d_2, int posOfFirstRing)

Additional Inherited Members
Protected Member Functions inherited from Alignment::ARotationEstimator< T >
std::vector< float > &	getRotations2D ()
constexpr bool	isInitialized () const
constexpr bool	isRefLoaded () const

Detailed Description

template<typename T>
class Alignment::CudaRotPolarEstimator< T >

Definition at line 44 of file cuda_rot_polar_estimator.h.

Constructor & Destructor Documentation

CudaRotPolarEstimator() [1/3]

template<typename T>

Alignment::CudaRotPolarEstimator< T >::CudaRotPolarEstimator ( )

inline

Definition at line 46 of file cuda_rot_polar_estimator.h.

                            {
        setDefault();
    }

~CudaRotPolarEstimator()

template<typename T>

virtual Alignment::CudaRotPolarEstimator< T >::~CudaRotPolarEstimator ( )

inlinevirtual

Definition at line 50 of file cuda_rot_polar_estimator.h.

                                     {
        release();
    }

CudaRotPolarEstimator() [2/3]

template<typename T>

Alignment::CudaRotPolarEstimator< T >::CudaRotPolarEstimator ( CudaRotPolarEstimator< T > &  o )

delete

CudaRotPolarEstimator() [3/3]

template<typename T>

Alignment::CudaRotPolarEstimator< T >::CudaRotPolarEstimator ( CudaRotPolarEstimator< T > &&  o )

inline

Definition at line 56 of file cuda_rot_polar_estimator.h.

                                                     {
        m_mainStream = o.m_mainStream;
        m_backgroundStream = o.m_backgroundStream;
        m_samples = o.m_samples;

        // device memory
        m_d_ref = o.m_d_ref;
        m_d_batch = o.m_d_batch;
        m_d_batchPolarOrCorr = o.m_d_batchPolarOrCorr;
        m_d_batchPolarFD = o.m_d_batchPolarFD;
        m_d_sumsOrMaxPos = o.m_d_sumsOrMaxPos;
        m_d_sumsSqr = o.m_d_sumsSqr;
        m_d_batchCorrSumFD = o.m_d_batchCorrSumFD;

        // FT plans
        m_singleToFD = o.m_singleToFD;
        m_batchToFD = o.m_batchToFD;
        m_batchToSD = o.m_batchToSD;

        // synch primitives
        m_mutex = o.m_mutex;
        m_cv = o.m_cv;
        m_isDataReady = o.m_isDataReady;

        // host memory
        m_h_batchMaxPositions = o.m_h_batchMaxPositions;

        // remove data from other
        o.setDefault();
    }

Member Function Documentation

operator=() [1/2]

template<typename T>

CudaRotPolarEstimator& Alignment::CudaRotPolarEstimator< T >::operator= ( const CudaRotPolarEstimator< T > &  other )

delete

operator=() [2/2]

template<typename T>

CudaRotPolarEstimator const& Alignment::CudaRotPolarEstimator< T >::operator= ( CudaRotPolarEstimator< T > &&  o )

delete

sComputeCorrelationsOneToN()

template<typename T >

void Alignment::CudaRotPolarEstimator< T >::sComputeCorrelationsOneToN ( const GPU &  gpu, const std::complex< T > *  d_in, std::complex< T > *  d_out, const std::complex< T > *  d_ref, const Dimensions &  dims, int  firstRingRadius  )

static

Definition at line 216 of file cuda_rot_polar_estimator.cpp.

                             {
    auto stream = *(cudaStream_t*)gpu.stream();
    dim3 dimBlock(32);
    dim3 dimGrid(
        ceil((dims.x() * dims.n()) / (float)dimBlock.x));
    if (std::is_same<T, float>::value) {
        computePolarCorrelationsSumOneToNKernel<float2>
            <<<dimGrid, dimBlock, 0, stream>>> (
            (float2*)d_in, (float2*)d_out,
            (float2*)d_ref,
            firstRingRadius,
            dims.x(), dims.y(), dims.n());
    } else if (std::is_same<T, double>::value) {
        computePolarCorrelationsSumOneToNKernel<double2>
            <<<dimGrid, dimBlock, 0, stream>>> (
            (double2*)d_in, (double2*)d_out,
            (double2*)d_ref,
            firstRingRadius,
            dims.x(), dims.y(), dims.n());
    } else {
            REPORT_ERROR(ERR_TYPE_INCORRECT, "Not implemented");
    }
}

sComputePolarTransform() [1/2]

template<typename T>

template<bool FULL_CIRCLE>

static void Alignment::CudaRotPolarEstimator< T >::sComputePolarTransform ( const GPU &  gpu, const Dimensions &  dimIn, T *  d_in, const Dimensions &  dimOut, T *  d_out, int  posOfFirstRing  )

static

sComputePolarTransform() [2/2]

template<typename T>

template<bool FULL_CIRCLE>

void Alignment::CudaRotPolarEstimator< T >::sComputePolarTransform	(	const GPU &	gpu,
		const Dimensions &	dimIn,
		T *__restrict__	d_in,
		const Dimensions &	dimOut,
		T *__restrict__	d_out,
		int	posOfFirstRing
	)

Definition at line 281 of file cuda_rot_polar_estimator.cpp.

                            {
    assert (dimIn.x() == dimIn.y());
    assert ((dimOut.y() + 1) * 2 <= dimIn.x()); // assert that there's space around the biggest ring
    dim3 dimBlock(32);
    dim3 dimGrid(
        ceil((dimOut.x() * dimOut.n()) / (float)dimBlock.x));

    auto stream = *(cudaStream_t*)gpu.stream();

    polarFromCartesian<T, FULL_CIRCLE>
        <<<dimGrid, dimBlock, 0, stream>>> (
        d_in, dimIn.x(), dimIn.y(),
        d_out, dimOut.x(), dimOut.y(), dimOut.n(), posOfFirstRing);
}

sNormalize() [1/2]

template<typename T>

template<bool FULL_CIRCLE>

static void Alignment::CudaRotPolarEstimator< T >::sNormalize ( const GPU &  gpu, const Dimensions &  dimIn, T *  d_in, T *  d_1, T *  d_2, int  posOfFirstRing  )

static

sNormalize() [2/2]

template<typename T>

template<bool FULL_CIRCLE>

void Alignment::CudaRotPolarEstimator< T >::sNormalize	(	const GPU &	gpu,
		const Dimensions &	dim,
		T *__restrict__	d_in,
		T *__restrict__	d_sums,
		T *__restrict__	d_sumsSqr,
		int	posOfFirstRing
	)

Definition at line 304 of file cuda_rot_polar_estimator.cpp.

                            {
    dim3 dimBlock(64);
    dim3 dimGrid(
        ceil((dim.x() * dim.n()) / (float)dimBlock.x));

    auto stream = *(cudaStream_t*)gpu.stream();
    // clear the arrays
    size_t bytes = dim.n() * sizeof(T);
    gpuErrchk(cudaMemset(d_sums, 0, bytes));
    gpuErrchk(cudaMemset(d_sumsSqr, 0, bytes));
    computeSumSumSqr<T, FULL_CIRCLE>
        <<<dimGrid, dimBlock, 0, stream>>> (
        d_in, dim.x(), dim.y(), dim.n(),
        d_sums, d_sumsSqr, posOfFirstRing);

    const T piConst = FULL_CIRCLE ? (2 * M_PI) : M_PI;
    // sum of the first n terms of an arithmetic sequence
    // a1 = first radius (posOfFirstRing)
    // an = last radius (dim.y() - 1 + posOfFirstRing)
    // s = n * (a1 + an) / 2
    const T radiiSum = (dim.y() * (2 * posOfFirstRing + dim.y() - 1)) / (T)2;
    T norm = piConst * radiiSum;

    normalize<T>
        <<<dimGrid, dimBlock, 0, stream>>> (
        d_in, dim.x(), dim.y(), dim.n(),
        norm,
        d_sums, d_sumsSqr);
}

---

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

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