Alignment::CudaShiftCorrEstimator< T > Class Template Reference

#include <cuda_shift_corr_estimator.h>

Inheritance diagram for Alignment::CudaShiftCorrEstimator< T >:

Collaboration diagram for Alignment::CudaShiftCorrEstimator< T >:

Public Member Functions
	CudaShiftCorrEstimator ()
virtual	~CudaShiftCorrEstimator ()
void	init2D (const std::vector< HW *> &hw, AlignType type, const FFTSettings< T > &dims, size_t maxShift, bool includingBatchFT, bool includingSingleFT, bool allowDataOverwrite) override
void	release ()
void	load2DReferenceOneToN (const std::complex< T > *h_ref) override
void	load2DReferenceOneToN (const T *h_ref) override
void	computeCorrelations2DOneToN (std::complex< T > *h_inOut, bool center) override
void	computeCorrelations2DOneToN (const HW &hw, std::complex< T > *inOut, const std::complex< T > *ref, const Dimensions &dims, bool center) override
void	computeShift2DOneToN (T *others) override
HW &	getHW () const override
Public Member Functions inherited from Alignment::AShiftCorrEstimator< T >
	AShiftCorrEstimator ()
	AShiftCorrEstimator (const AShiftCorrEstimator &)=delete
	AShiftCorrEstimator (const AShiftCorrEstimator &&)=delete
virtual	~AShiftCorrEstimator ()
AShiftCorrEstimator &	operator= (const AShiftCorrEstimator &)=delete
AShiftCorrEstimator &	operator= (const AShiftCorrEstimator &&)=delete
Public Member Functions inherited from Alignment::AShiftEstimator< T >
	AShiftEstimator ()
	AShiftEstimator (const AShiftEstimator &)=delete
	AShiftEstimator (const AShiftEstimator &&)=delete
virtual	~AShiftEstimator ()
AShiftEstimator &	operator= (const AShiftEstimator &)=delete
AShiftEstimator &	operator= (const AShiftEstimator &&)=delete
std::vector< Point2D< float > >	getShifts2D ()
constexpr bool	isInitialized () const
constexpr Dimensions	getDimensions () const
constexpr AlignType	getAlignType () const

Static Public Member Functions
static std::vector< Point2D< float > >	computeShifts2DOneToN (const std::vector< GPU *> &gpus, std::complex< T > *d_othersF, T *d_othersS, std::complex< T > *d_ref, const FFTSettings< T > &settings, cufftHandle plan, T *h_centers, size_t maxShift)
template<bool center>
static void	sComputeCorrelations2DOneToN (const GPU &gpu, std::complex< T > *d_inOut, const std::complex< T > *d_ref, const Dimensions &dims)

Additional Inherited Members
Protected Member Functions inherited from Alignment::AShiftCorrEstimator< T >
virtual void	init2D (AlignType type, const FFTSettings< T > &dims, size_t maxShift, bool includingBatchFT, bool includingSingleFT, bool allowDataOverwrite)
void	init2D (const std::vector< HW *> &hw, AlignType type, const Dimensions &dims, size_t batch, size_t maxShift)
Protected Member Functions inherited from Alignment::AShiftEstimator< T >
virtual void	init2D (AlignType type, const Dimensions &dims, size_t batch, size_t maxShift)
Protected Attributes inherited from Alignment::AShiftCorrEstimator< T >
FFTSettings< T > *	m_settingsInv
size_t	m_centerSize
bool	m_includingBatchFT
bool	m_includingSingleFT
bool	m_is_ref_FD_loaded
bool	m_allowDataOverwrite
Protected Attributes inherited from Alignment::AShiftEstimator< T >
AlignType	m_type
const Dimensions *	m_dims
size_t	m_batch
size_t	m_maxShift
std::vector< Point2D< float > >	m_shifts2D
bool	m_is_ref_loaded
bool	m_is_shift_computed
bool	m_isInit

Detailed Description

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

Definition at line 45 of file cuda_shift_corr_estimator.h.

Constructor & Destructor Documentation

CudaShiftCorrEstimator()

template<typename T>

Alignment::CudaShiftCorrEstimator< T >::CudaShiftCorrEstimator ( )

inline

Definition at line 47 of file cuda_shift_corr_estimator.h.

                             {
        setDefault();
    }

~CudaShiftCorrEstimator()

template<typename T>

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

inlinevirtual

Definition at line 51 of file cuda_shift_corr_estimator.h.

                                      {
        release();
    }

Member Function Documentation

computeCorrelations2DOneToN() [1/2]

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::computeCorrelations2DOneToN ( std::complex< T > *  h_inOut, bool  center  )

overridevirtual

Implements Alignment::AShiftCorrEstimator< T >.

Definition at line 184 of file cuda_shift_corr_estimator.cpp.

                                             {
    bool isReady = (this->m_isInit && (AlignType::OneToN == this->m_type) && m_is_d_single_FD_loaded);

    if ( ! isReady) {
        REPORT_ERROR(ERR_LOGIC_ERROR, "Not ready to execute. Call init() before");
    }

    m_loadStream->set();
    m_workStream->set();
    auto loadStream = *(cudaStream_t*)m_loadStream->stream();
    auto workStream = *(cudaStream_t*)m_workStream->stream();
    // process signals in batches
    for (size_t offset = 0; offset < this->m_settingsInv->fDim().n(); offset += this->m_settingsInv->batch()) {
        // how many signals to process
        size_t toProcess = std::min(this->m_settingsInv->batch(), this->m_settingsInv->fDim().n() - offset);

        // copy memory
        gpuErrchk(cudaMemcpyAsync(
                m_d_batch_FD,
                h_inOut + offset * this->m_settingsInv->fDim().xy(),
                toProcess * this->m_settingsInv->fBytesSingle(),
                cudaMemcpyHostToDevice, loadStream));
        m_loadStream->synch();

        auto dims = Dimensions(
                this->m_settingsInv->fDim().x(),
                this->m_settingsInv->fDim().y(),
                1,
                toProcess);
        if (center) {
            sComputeCorrelations2DOneToN<true>(*m_workStream, m_d_batch_FD, m_d_single_FD, dims);
        } else {
            sComputeCorrelations2DOneToN<false>(*m_workStream, m_d_batch_FD, m_d_single_FD, dims);
        }
        m_workStream->synch();

        // copy data back
        gpuErrchk(cudaMemcpyAsync(
                h_inOut + offset * this->m_settingsInv->fDim().xy(),
                m_d_batch_FD,
                toProcess * this->m_settingsInv->fBytesSingle(),
                cudaMemcpyDeviceToHost, loadStream));
    }
}

computeCorrelations2DOneToN() [2/2]

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::computeCorrelations2DOneToN ( const HW &  hw, std::complex< T > *  inOut, const std::complex< T > *  ref, const Dimensions &  dims, bool  center  )

overridevirtual

Implements Alignment::AShiftCorrEstimator< T >.

Definition at line 419 of file cuda_shift_corr_estimator.cpp.

                     {
    const GPU *gpu;
    try {
        gpu = &dynamic_cast<const GPU&>(hw);
    } catch (std::bad_cast&) {
        REPORT_ERROR(ERR_ARG_INCORRECT, "Instance of GPU expected");
    }
    if (center) {
        return sComputeCorrelations2DOneToN<true>(*gpu, inOut, ref, dims);
    } else {
        return sComputeCorrelations2DOneToN<false>(*gpu, inOut, ref, dims);
    }
}

computeShift2DOneToN()

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::computeShift2DOneToN ( T *  others )

overridevirtual

Implements Alignment::AShiftEstimator< T >.

Definition at line 281 of file cuda_shift_corr_estimator.cpp.

                   {
    bool isReady = (this->m_isInit && (AlignType::OneToN == this->m_type)
            && m_is_d_single_FD_loaded && this->m_includingBatchFT);

    if ( ! isReady) {
        REPORT_ERROR(ERR_LOGIC_ERROR, "Not ready to execute. Call init() before");
    }
    const auto isGpuPtr = m_workStream->isGpuPointer(others);
    if(!isGpuPtr) {
        if ( ! GPU::isMemoryPinned(others)) {
            REPORT_ERROR(ERR_LOGIC_ERROR, "Input memory has to be pinned (page-locked)");
        }
        if (this->m_allowDataOverwrite) {
            REPORT_ERROR(ERR_LOGIC_ERROR, "Incompatible parameters: allowDataOverwrite && 'others' data on host");
        }
    }

    m_workStream->set();
    // make sure that all work (e.g. loading and processing of the reference) is done
    m_loadStream->synch();
    m_workStream->synch();
    // start loading data at the background
    m_isDataReady = false;
    const bool createLocalCopy = ( ! this->m_allowDataOverwrite);
    std::thread loadingThread;
    if (createLocalCopy) {
        loadingThread = std::thread(&CudaShiftCorrEstimator<T>::loadThreadRoutine, this, others);
    }

    // reserve enough space for shifts
    this->m_shifts2D.reserve(this->m_settingsInv->fDim().n());
    // process signals
    for (size_t offset = 0; offset < this->m_settingsInv->fDim().n(); offset += this->m_settingsInv->batch()) {
        size_t batchOffset = createLocalCopy
                ? offset
                : std::min(offset, this->m_settingsInv->fDim().n() - this->m_settingsInv->batch());
        // how many signals to process
        size_t toProcess = std::min(
                this->m_settingsInv->batch(),
                this->m_settingsInv->fDim().n() - batchOffset);

        if (createLocalCopy) {
            waitAndConvert();
        } else {
            T *in = others + (batchOffset * this->m_settingsInv->sDim().sizeSingle());
            CudaFFT<T>::fft(*m_batchToFD, in, m_d_batch_FD);
        }

        // compute shifts
        auto shifts = computeShifts2DOneToN(
                std::vector<GPU*>{m_workStream, m_loadStream},
                m_d_batch_FD,
                m_d_batch_SD_work,
                m_d_single_FD,
                this->m_settingsInv->createSubset(toProcess),
                *m_batchToSD,
                this->m_h_centers,
                this->m_maxShift);

        size_t start = createLocalCopy
                ? 0
                : (offset - batchOffset);
        // append shifts to existing results
        this->m_shifts2D.insert(this->m_shifts2D.end(),
                shifts.begin() + start,
                shifts.end());
    }
    if (createLocalCopy) {
        loadingThread.join();
    }
    // update state
    this->m_is_shift_computed = true;
}

computeShifts2DOneToN()

template<typename T >

std::vector< Point2D< float > > Alignment::CudaShiftCorrEstimator< T >::computeShifts2DOneToN ( const std::vector< GPU *> &  gpus, std::complex< T > *  d_othersF, T *  d_othersS, std::complex< T > *  d_ref, const FFTSettings< T > &  settings, cufftHandle  plan, T *  h_centers, size_t  maxShift  )

static

Definition at line 357 of file cuda_shift_corr_estimator.cpp.

                         {
    // we need even input in order to perform the shift (in FD, while correlating) properly
    assert(0 == (settings.sDim().x() % 2));
    assert(0 == (settings.sDim().y() % 2));
    assert(1 == settings.sDim().zPadded());
    assert(2 == gpus.size());

    size_t centerSize = maxShift * 2 + 1;

    // make sure we have enough memory for centers of the correlation functions
    assert(settings.fBytesBatch() >= (settings.sDim().n() * centerSize * centerSize * sizeof(T)));

    auto workGPU = gpus.at(0);
    auto loadGPU = gpus.at(1);
    loadGPU->set();
    workGPU->set();
    auto loadStream = *(cudaStream_t*)loadGPU->stream();
    auto workStream = *(cudaStream_t*)workGPU->stream();

    // correlate signals and shift FT so that it will be centered after IFT
    sComputeCorrelations2DOneToN<true>(*workGPU,
            d_othersF, d_ref,
            settings.fDim());

    // perform IFT
    CudaFFT<T>::ifft(plan, d_othersF, d_othersS);

    // locate maxima
    auto p_pos = (float*)d_othersF;
    auto h_pos = (float*)h_centers;
    ExtremaFinder::CudaExtremaFinder<T>::sFindMax2DAroundCenter(
            *workGPU, settings.sDim(), d_othersS, p_pos, nullptr, maxShift);
    workGPU->synch();
    // copy data back
    gpuErrchk(cudaMemcpyAsync(h_pos, p_pos,
            settings.batch() * sizeof(float),
            cudaMemcpyDeviceToHost, loadStream));
    loadGPU->synch();

    // convert absolute 1D index to logical 2D index
    // FIXME DS extract this to some utils
    auto result = std::vector<Point2D<float>>();
    result.reserve(settings.sDim().n());
    for (size_t n = 0; n < settings.batch(); ++n) {
        float y = (size_t)h_pos[n] / settings.sDim().x();
        float x = (size_t)h_pos[n] % settings.sDim().x();
        result.emplace_back(
            x - settings.sDim().x() / 2,
            y - settings.sDim().y() / 2);
    }
    return result;
}

getHW()

template<typename T>

HW& Alignment::CudaShiftCorrEstimator< T >::getHW ( ) const

inlineoverridevirtual

Implements Alignment::AShiftEstimator< T >.

Definition at line 93 of file cuda_shift_corr_estimator.h.

                               {
        return *m_workStream;
    }

init2D()

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::init2D ( const std::vector< HW *> &  hw, AlignType  type, const FFTSettings< T > &  dims, size_t  maxShift, bool  includingBatchFT, bool  includingSingleFT, bool  allowDataOverwrite  )

overridevirtual

Implements Alignment::AShiftCorrEstimator< T >.

Definition at line 34 of file cuda_shift_corr_estimator.cpp.

                                {
    // FIXME DS consider tunning the size of the input (e.g. 436x436x50)
    if (2 != hw.size()) {
        REPORT_ERROR(ERR_ARG_INCORRECT, "Two GPU streams are needed");
    }
    release(); // FIXME DS implement lazy init
    m_workStream = dynamic_cast<GPU*>(hw.at(0));
    m_loadStream = dynamic_cast<GPU*>(hw.at(1));
    if ((nullptr == m_workStream)
        || (nullptr == m_loadStream)) {
        REPORT_ERROR(ERR_LOGIC_ERROR, "Instance of GPU is expected");
    }
    m_workStream->set();
    m_loadStream->set();
    AShiftCorrEstimator<T>::init2D(type, settings, maxShift,
        includingBatchFT, includingSingleFT, allowDataOvewrite);

    // synch primitives
    m_mutex = new std::mutex();
    m_cv = new std::condition_variable();

    this->m_isInit = true;
}

load2DReferenceOneToN() [1/2]

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::load2DReferenceOneToN ( const std::complex< T > *  h_ref )

overridevirtual

Implements Alignment::AShiftCorrEstimator< T >.

Definition at line 90 of file cuda_shift_corr_estimator.cpp.

                                                                                {
    auto isReady = (this->m_isInit && (AlignType::OneToN == this->m_type));
    if ( ! isReady) {
        REPORT_ERROR(ERR_LOGIC_ERROR, "Not ready to load a reference signal");
    }

    m_loadStream->set();
    auto loadStream = *(cudaStream_t*)m_loadStream->stream();
    // copy reference to GPU
    gpuErrchk(cudaMemcpyAsync(m_d_single_FD, h_ref, this->m_settingsInv->fBytesSingle(),
            cudaMemcpyHostToDevice, loadStream));

    // update state
    m_is_d_single_FD_loaded = true;
}

load2DReferenceOneToN() [2/2]

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::load2DReferenceOneToN ( const T *  h_ref )

overridevirtual

Implements Alignment::AShiftEstimator< T >.

Definition at line 107 of file cuda_shift_corr_estimator.cpp.

                                                                    {
    auto isReady = (this->m_isInit && (AlignType::OneToN == this->m_type) && this->m_includingSingleFT);
    if ( ! isReady) {
        REPORT_ERROR(ERR_LOGIC_ERROR, "Not ready to load a reference signal");
    }
    m_loadStream->set();
    m_workStream->set();
    auto loadStream = *(cudaStream_t*)m_loadStream->stream();
    auto workStream = *(cudaStream_t*)m_workStream->stream();

    // copy reference to GPU
    gpuErrchk(cudaMemcpyAsync(m_d_single_SD, h_ref, this->m_settingsInv->sBytesSingle(),
            cudaMemcpyHostToDevice, loadStream));
    m_loadStream->synch();

    // perform FT
    CudaFFT<T>::fft(*m_singleToFD, m_d_single_SD, m_d_single_FD);

    // update state
    m_is_d_single_FD_loaded = true;
}

release()

template<typename T >

void Alignment::CudaShiftCorrEstimator< T >::release ( )

virtual

Reimplemented from Alignment::AShiftCorrEstimator< T >.

Definition at line 130 of file cuda_shift_corr_estimator.cpp.

                                        {
    // device memory
    gpuErrchk(cudaFree(m_d_single_FD));
    gpuErrchk(cudaFree(m_d_batch_FD));
    gpuErrchk(cudaFree(m_d_single_SD));
    gpuErrchk(cudaFree(m_d_batch_SD_work));
    if ( ! this->m_allowDataOverwrite) {
        gpuErrchk(cudaFree(m_d_batch_SD_load));
    }

    // host memory
    delete[] m_h_centers;

    // FT plans
    CudaFFT<T>::release(m_singleToFD);
    CudaFFT<T>::release(m_batchToFD);
    CudaFFT<T>::release(m_batchToSD);

    // synch primitives
    delete m_mutex;
    delete m_cv;

    AShiftCorrEstimator<T>::release();

    CudaShiftCorrEstimator<T>::setDefault();
}

sComputeCorrelations2DOneToN()

template<typename T >

template<bool center>

void Alignment::CudaShiftCorrEstimator< T >::sComputeCorrelations2DOneToN ( const GPU &  gpu, std::complex< T > *  d_inOut, const std::complex< T > *  d_ref, const Dimensions &  dims  )

static

Definition at line 440 of file cuda_shift_corr_estimator.cpp.

                                {
    if (center) {
        // we cannot assert xDim, as we don't know if the spatial size was even
        assert(0 == (dims.y() % 2));
    }
    assert(0 < dims.x());
    assert(0 < dims.y());
    assert(1 == dims.z());
    assert(0 < dims.n());

    // create threads / blocks
    dim3 dimBlock(64, 1, 1);
    dim3 dimGrid(
            std::ceil(dims.x() / (float)dimBlock.x),
            dims.n(), 1);
    auto stream = *(cudaStream_t*)gpu.stream();
    if (std::is_same<T, float>::value) {
        computeCorrelations2DOneToNKernel<float2, center>
            <<<dimGrid, dimBlock, 0, stream>>> (
                (float2*)d_inOut, (float2*)d_ref,
                dims.x(), dims.y(), dims.n());
    } else if (std::is_same<T, double>::value) {
        computeCorrelations2DOneToNKernel<double2, center>
            <<<dimGrid, dimBlock, 0, stream>>> (
                (double2*)d_inOut, (double2*)d_ref,
                dims.x(), dims.y(), dims.n());
    } else {
        REPORT_ERROR(ERR_TYPE_INCORRECT, "Not implemented");
    }
}

---

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

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