gpuInterpolator.h

#pragma once
 
#include <isce3/core/forward.h>
#include <isce3/core/Common.h>
 
#include <thrust/host_vector.h>
 
using isce3::core::Matrix;

namespace isce3{ namespace cuda{ namespace core{
template <class U>
    class gpuInterpolator {
        public:
            CUDA_HOSTDEV gpuInterpolator() {}
            CUDA_HOSTDEV virtual ~gpuInterpolator() {}
            CUDA_DEV virtual U interpolate(double, double, const U*, size_t, size_t) = 0;
    };
 

template <class U>
class gpuBilinearInterpolator : public isce3::cuda::core::gpuInterpolator<U> {
    public:
        CUDA_HOSTDEV gpuBilinearInterpolator(){};
        CUDA_DEV U interpolate(double, double, const U*, size_t, size_t);
        CUDA_HOST void interpolate_h(const Matrix<double>&, Matrix<U>&, double, double, U*);
};
 

template <class U>
class gpuBicubicInterpolator : public isce3::cuda::core::gpuInterpolator<U> {
    public:
        CUDA_HOSTDEV gpuBicubicInterpolator(){};
        CUDA_DEV U interpolate(double, double, const U*, size_t, size_t);
        CUDA_HOST void interpolate_h(const Matrix<double>&, Matrix<U>&, double, double, U*);
};
 

template <class U>
class gpuSpline2dInterpolator : public isce3::cuda::core::gpuInterpolator<U> {
    protected:
        size_t _order;
    public:
        CUDA_HOSTDEV gpuSpline2dInterpolator(size_t order):_order(order){};
        CUDA_DEV U interpolate(double, double, const U*, size_t, size_t);
        CUDA_HOST void interpolate_h(const Matrix<double>&, Matrix<U>&, double, double, U*);
};
 

template <class U>
class gpuSinc2dInterpolator : public isce3::cuda::core::gpuInterpolator<U> {
    protected:
        double *_kernel;
        // Number of divisions per sample (total number samples in the lookup
        // table is kernel length * decimation factor)
        int _decimationFactor;
        // size of kernel
        int _kernelLength;
        // size of half kernel
        int _halfKernelLength;
        // True if initialized from host,
        // False if copy-constructed from gpuSinc2dInterpolator on device
        bool _owner;
    public:
        CUDA_HOSTDEV gpuSinc2dInterpolator(){};

        CUDA_HOST gpuSinc2dInterpolator(
                const int kernelLength, const int decimationFactor,
                const double beta = 1.0, const double pedestal = 0.0);

        CUDA_DEV gpuSinc2dInterpolator(double *device_filter, int kernelLength,
                int decimationFactor) :
            _kernel(device_filter),
            _decimationFactor(decimationFactor),
            _kernelLength(kernelLength),
            _halfKernelLength(kernelLength / 2),
            _owner(false) {};
 
        CUDA_HOSTDEV gpuSinc2dInterpolator(const gpuSinc2dInterpolator &i):
            _kernel(i._kernel),
            _decimationFactor(i._decimationFactor),
            _kernelLength(i._kernelLength),
            _halfKernelLength(i._halfKernelLength),
            _owner(false) {};
 
        CUDA_HOSTDEV ~gpuSinc2dInterpolator();
        CUDA_DEV U interpolate(double, double, const U*, size_t, size_t);
        CUDA_HOST void interpolate_h(const Matrix<double>&, Matrix<U>&, double, double, U*);
};

CUDA_HOST void compute_normalized_coefficients(
        const double beta, const int relfiltlen, const int decfactor,
        const double pedestal, thrust::host_vector<double>& filter);
 

template <class T>
class gpuNearestNeighborInterpolator : public isce3::cuda::core::gpuInterpolator<T> {
    public:
        CUDA_HOSTDEV gpuNearestNeighborInterpolator(){};
        CUDA_DEV T interpolate(double x, double y, const T* z, size_t nx, size_t ny = 0);
};
 
}}}