Kernels.h

#pragma once

#include "forward.h"

#include <cmath>
#include <thrust/device_vector.h>

#include <isce3/core/Common.h>
#include <isce3/core/Kernels.h>

namespace isce3 { namespace cuda { namespace core {

template<typename T, class Derived>
class Kernel {
public:
    using value_type = T;

    explicit constexpr Kernel(double width) noexcept
        : _halfwidth(std::abs(width) / 2.)
    {}

    constexpr double halfwidth() const noexcept { return _halfwidth; }

    constexpr double width() const noexcept { return 2. * _halfwidth; }

    CUDA_HOSTDEV T operator()(double t) const;

private:
    double _halfwidth;
};

template<typename T>
class BartlettKernel : public Kernel<T, BartlettKernel<T>> {
    using Base = Kernel<T, BartlettKernel<T>>;
    friend Base;

public:
    using view_type = BartlettKernel<T>;

    explicit constexpr BartlettKernel(double width) : Base(width) {}

    BartlettKernel(const isce3::core::BartlettKernel<T>& other)
        : BartlettKernel(other.width())
    {}

protected:
    constexpr T eval(double t) const;
};

template<typename T>
class LinearKernel : public BartlettKernel<T> {
    using Base = BartlettKernel<T>;

public:
    using view_type = LinearKernel<T>;

    constexpr LinearKernel() : Base(2.) {}

    LinearKernel(const isce3::core::LinearKernel<T>&) : LinearKernel() {}
};

template<typename T>
class KnabKernel : public Kernel<T, KnabKernel<T>> {
    using Base = Kernel<T, KnabKernel<T>>;
    friend Base;

public:
    using view_type = KnabKernel<T>;

    constexpr KnabKernel(double width, double bandwidth);

    KnabKernel(const isce3::core::KnabKernel<T>& other)
        : Base(other.width()), _bandwidth(other.bandwidth())
    {}

    constexpr double bandwidth() const noexcept { return _bandwidth; }

protected:
    CUDA_HOSTDEV T eval(double t) const;

private:
    double _bandwidth;
};

template<typename T>
class TabulatedKernelView : public Kernel<T, TabulatedKernelView<T>> {
    using Base = Kernel<T, TabulatedKernelView<T>>;
    friend Base;

    friend class TabulatedKernel<T>;

public:
    using view_type = TabulatedKernelView<T>;

    TabulatedKernelView(const TabulatedKernel<T>& kernel);

    CUDA_DEV T operator()(double t) const { return eval(t); }

protected:
    CUDA_DEV T eval(double t) const;

private:
    const T* _table;
    int _imax;
    T _rdx;
};

template<typename T>
class TabulatedKernel : public Kernel<T, TabulatedKernel<T>> {
    using Base = Kernel<T, TabulatedKernel<T>>;
    friend Base;

    friend class TabulatedKernelView<T>;

public:
    using view_type = TabulatedKernelView<T>;

    template<class OtherKernel>
    TabulatedKernel(const OtherKernel& kernel, int n);

    TabulatedKernel(const isce3::core::TabulatedKernel<T>&);

    CUDA_DEV T operator()(double t) const { return eval(t); }

protected:
    CUDA_DEV T eval(double t) const;

private:
    thrust::device_vector<T> _table;
    int _imax;
    T _rdx;
};

template<typename T>
class ChebyKernelView : public Kernel<T, ChebyKernelView<T>> {
    using Base = Kernel<T, ChebyKernelView<T>>;
    friend Base;

    friend class ChebyKernel<T>;

public:
    using view_type = ChebyKernelView<T>;

    ChebyKernelView(const ChebyKernel<T>& kernel);

    CUDA_DEV T operator()(double t) const { return eval(t); }

protected:
    CUDA_DEV T eval(double t) const;

private:
    const T* _coeffs;
    int _n;
    T _scale;
};

template<typename T>
class ChebyKernel : public Kernel<T, ChebyKernel<T>> {
    using Base = Kernel<T, ChebyKernel<T>>;
    friend Base;

    friend class ChebyKernelView<T>;

public:
    using view_type = ChebyKernelView<T>;

    template<class OtherKernel>
    ChebyKernel(const OtherKernel& kernel, int n);

    ChebyKernel(const isce3::core::ChebyKernel<T>& other)
        : Base(other.width()), _scale(4. / other.width()),
          _coeffs(other.coeffs())
    {}

    CUDA_DEV T operator()(double t) const { return eval(t); }

protected:
    CUDA_DEV T eval(double t) const;

private:
    T _scale;
    thrust::device_vector<T> _coeffs;
};

}}} // namespace isce3::cuda::core

#include "Kernels.icc"