gpuSignal.h

// -*- C++ -*-
// -*- coding: utf-8 -*-
//
// Source Author: Liang Yu
// Copyright 2019

#pragma once

#include "forward.h"

#include <complex>
#include <valarray>

#include <cufft.h>
#include <thrust/complex.h>

#include <isce3/core/Common.h>

template<class T>
class isce3::cuda::signal::gpuSignal {

    public:
        // Default constructor
        gpuSignal() {};
        gpuSignal(cufftType _type);
        ~gpuSignal();

        void azimuthFFT(int ncolumns, int nrows);

        void rangeFFT(int ncolumns, int nrows);

        void FFT2D(int ncolumns, int nrows);

        void fftPlan(int rank, int* n, int howmany,
                    int* inembed, int istride, int idist,
                    int* onembed, int ostride, int odist);

        void nextPowerOfTwo(size_t N, size_t& fftLength);

        inline void _configureRangeFFT(int ncolumns, int nrows);

        inline void _configureAzimuthFFT(int ncolumns, int nrows);

        void dataToDevice(std::complex<T> *input);
        void dataToDevice(std::valarray<std::complex<T>> &input);
        void dataToHost(std::complex<T> *output);
        void dataToHost(std::valarray<std::complex<T>> &output);

        void forward();

        void forwardC2C(std::complex<T> *input, std::complex<T> *output);
        void forwardC2C(std::valarray<std::complex<T>> &input,
                        std::valarray<std::complex<T>> &output);
        void forwardZ2Z(std::complex<T> *input, std::complex<T> *output);
        void forwardZ2Z(std::valarray<std::complex<T>> &input,
                        std::valarray<std::complex<T>> &output);
        void forwardD2Z(T *input, std::complex<T> *output);

        void forward(std::complex<T> *input, std::complex<T> *output);
        void forward(std::valarray<std::complex<T>> &input,
                     std::valarray<std::complex<T>> &output);

        void forwardDevMem(T *input, T *output);
        void forwardDevMem(T *dataInPlace);

        void inverse();

        void inverseC2C(std::complex<T> *input, std::complex<T> *output);
        void inverseC2C(std::valarray<std::complex<T>> &input,
                        std::valarray<std::complex<T>> &output);
        void inverseZ2Z(std::complex<T> *input, std::complex<T> *output);
        void inverseZ2Z(std::valarray<std::complex<T>> &input,
                        std::valarray<std::complex<T>> &output);
        void inverseZ2D(std::complex<T> *input, T *output);

        void inverse(std::complex<T> *input, std::complex<T> *output);
        void inverse(std::valarray<std::complex<T>> &input,
                     std::valarray<std::complex<T>> &output);

        void inverseDevMem(T *input, T *output);
        void inverseDevMem(T *dataInPlace);

        void upsample(std::valarray<std::complex<T>> &input,
                      std::valarray<std::complex<T>> &output,
                      int row,
                      int ncols,
                      int upsampleFactor);
        void upsample(std::valarray<std::complex<T>> &input,
                      std::valarray<std::complex<T>> &output,
                      int row,
                      int ncols,
                      int upsampleFactor,
                      std::valarray<std::complex<T>> &shiftImpact);

        int getRows() {return _rows;};
        int getColumns() {return _columns;};
        int getNumElements() {return _n_elements;};

        T* getDevicePtr() {return _d_data;};

    private:
        cufftHandle _plan;
        bool _plan_set;
        cufftType _cufft_type;

        // FFT plan parameters
        int _rank;
        int _n[2];
        int _howmany;
        int _inembed[2];
        int _istride;
        int _idist;
        int _onembed[2];
        int _ostride;
        int _odist;
        int _n_elements;
        int _rows;
        int _columns;

        // device memory pointers
        T *_d_data;
        bool _d_data_set;
};

template<class T>
CUDA_GLOBAL void rangeShift_g(thrust::complex<T> *data_lo_res,
        thrust::complex<T> *data_hi_res,
        int n_rows,
        int n_cols_lo,
        int n_cols_hi);

template<class T>
CUDA_GLOBAL void rangeShiftImpactMult_g(thrust::complex<T> *data_lo_res,
        thrust::complex<T> *data_hi_res,
        thrust::complex<T> *impact_shift,
        int n_rows,
        int n_cols_lo,
        int n_cols_hi);

template<class T>
void upsample(isce3::cuda::signal::gpuSignal<T> &fwd,
        isce3::cuda::signal::gpuSignal<T> &inv,
        thrust::complex<T> *input,
        thrust::complex<T> *output);

template<class T>
void upsample(isce3::cuda::signal::gpuSignal<T> &fwd,
        isce3::cuda::signal::gpuSignal<T> &inv,
        thrust::complex<T> *input,
        thrust::complex<T> *output,
        thrust::complex<T> *shiftImpact);

template<class T>
void upsample(isce3::cuda::signal::gpuSignal<T> &fwd,
        isce3::cuda::signal::gpuSignal<T> &inv,
        std::valarray<std::complex<T>> &input,
        std::valarray<std::complex<T>> &output);

template<class T>
void upsample(isce3::cuda::signal::gpuSignal<T> &fwd,
        isce3::cuda::signal::gpuSignal<T> &inv,
        std::valarray<std::complex<T>> &input,
        std::valarray<std::complex<T>> &output,
        std::valarray<std::complex<T>> &shiftImpact);

template<class T>
CUDA_GLOBAL void normalize_g(thrust::complex<T> *data,
        T normalization,
        size_t n_elements);