tex_linear_interpolation_1D.hh

Go to the documentation of this file.

#pragma once
 
// DiFfRG
#include <DiFfRG/common/cuda_prefix.hh>
#include <DiFfRG/common/math.hh>
#include <DiFfRG/physics/interpolation/common.hh>
 
// standard library
#include <memory>
 
// external libraries
#include <autodiff/forward/real.hpp>
 
namespace DiFfRG
{
  template <typename NT, typename Coordinates> class TexLinearInterpolator1D
  {
    static_assert(Coordinates::dim == 1, "TexLinearInterpolator1D requires 1D coordinates");
 
    static constexpr int max_device_count = 4;
 
  public:
    TexLinearInterpolator1D(const std::vector<NT> &data, const Coordinates &coordinates);
    TexLinearInterpolator1D(const NT *data, const Coordinates &coordinates);
    TexLinearInterpolator1D(const Coordinates &coordinates);
    TexLinearInterpolator1D(const TexLinearInterpolator1D &other);
 
    ~TexLinearInterpolator1D();
 
    template <typename NT2> void update(const NT2 *data)
    {
      if (!owner) throw std::runtime_error("Cannot update data of non-owner interpolator");
 
      if constexpr (std::is_same_v<NT2, autodiff::real>)
        for (uint i = 0; i < size; ++i) {
          m_data[i] = static_cast<float>(val(data[i]));
          m_data_AD[i] = static_cast<float>(derivative(data[i]));
        }
      else
        for (uint i = 0; i < size; ++i)
          m_data[i] = static_cast<float>(data[i]);
 
      update();
    }
    void update();
 
    float *data() const;
    float *data_AD() const;
 
    using ReturnType = typename internal::__TLITypes<NT>::ReturnType;
 
    __device__ __host__ ReturnType operator()(const float x) const
    {
#ifdef __CUDA_ARCH__
      if constexpr (std::is_same_v<ReturnType, autodiff::real>)
        return std::array<double, 2>{tex1D<float>(texture[0], coordinates.backward(x) + 0.5f),
                                     tex1D<float>(texture_AD[0], coordinates.backward(x) + 0.5f)};
      else if constexpr (std::is_same_v<ReturnType, float>)
        return tex1D<float>(texture[0], coordinates.backward(x) + 0.5f);
#else
      float idx = coordinates.backward(x);
      idx = std::max(0.f, std::min(idx, static_cast<float>(size - 1)));
      if constexpr (std::is_same_v<ReturnType, autodiff::real>)
        return std::array<double, 2>{{m_data[uint(std::floor(idx))] * (1.f - idx + std::floor(idx)) +
                                          m_data[uint(std::ceil(idx))] * (idx - std::floor(idx)),
                                      m_data_AD[uint(std::floor(idx))] * (1.f - idx + std::floor(idx)) +
                                          m_data_AD[uint(std::ceil(idx))] * (idx - std::floor(idx))}};
      else if constexpr (std::is_same_v<ReturnType, float>)
        return m_data[uint(std::floor(idx))] * (1.f - idx + std::floor(idx)) +
               m_data[uint(std::ceil(idx))] * (idx - std::floor(idx));
#endif
    }
 
    ReturnType &operator[](const uint i);
    const ReturnType &operator[](const uint i) const;
 
    const Coordinates &get_coordinates() const { return coordinates; }
 
  private:
    const uint size;
    const Coordinates coordinates;
 
    std::shared_ptr<float[]> m_data;
    std::vector<cudaArray_t> device_array;
    cudaTextureObject_t texture[max_device_count];
 
    std::shared_ptr<float[]> m_data_AD;
    std::vector<cudaArray_t> device_array_AD;
    cudaTextureObject_t texture_AD[max_device_count];
 
    const bool owner;
    int n_devices;
  };
} // namespace DiFfRG