GpuView.hpp

/*
  Copyright 2024 SINTEF AS
 
  This file is part of the Open Porous Media project (OPM).
 
  OPM is free software: you can redistribute it and/or modify
  it under the terms of the GNU General Public License as published by
  the Free Software Foundation, either version 3 of the License, or
  (at your option) any later version.
 
  OPM is distributed in the hope that it will be useful,
  but WITHOUT ANY WARRANTY; without even the implied warranty of
  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  GNU General Public License for more details.
 
  You should have received a copy of the GNU General Public License
  along with OPM.  If not, see <http://www.gnu.org/licenses/>.
*/
#ifndef OPM_GPUVIEW_HEADER_HPP
#define OPM_GPUVIEW_HEADER_HPP
 
#include <dune/common/fvector.hh>
 
#include <opm/common/ErrorMacros.hpp>
 
#include <opm/common/utility/gpuDecorators.hpp>
#include <opm/simulators/linalg/gpuistl/detail/safe_conversion.hpp>
 
#include <stdexcept>
#include <vector>
 
#include <fmt/core.h>
 
namespace Opm::gpuistl
{

template <typename T>
class GpuView
{
public:
    using value_type = T;
    explicit GpuView() = default;
 
    //TODO: we probably dont need anything like this or is it useful to have views also be able to handle things on CPU?
    GpuView(std::vector<T>& data);

    __host__ __device__ T& operator[](size_t idx){
#ifndef NDEBUG
        assertInRange(idx);
#endif
    return m_dataPtr[idx];
    }

    __host__ __device__ T operator[](size_t idx) const {
#ifndef NDEBUG
        assertInRange(idx);
#endif
    return m_dataPtr[idx];
    }
 

    GpuView(T* dataOnHost, size_t numberOfElements)
        : m_dataPtr(dataOnHost), m_numberOfElements(numberOfElements)
    {
    }

    ~GpuView() = default;

    __host__ __device__ T* data(){
        return m_dataPtr;
    }

    __host__ __device__ const T* data() const{
        return m_dataPtr;
    }

    __host__ __device__ bool empty() const {
        return m_numberOfElements == 0;
    }

    __host__ __device__ T& front()
    {
#ifndef NDEBUG
        assertHasElements();
#endif
        return m_dataPtr[0];
    }

    __host__ __device__ T& back()
    {
#ifndef NDEBUG
        assertHasElements();
#endif
        return m_dataPtr[m_numberOfElements-1];
    }

    __host__ __device__ T front() const
    {
#ifndef NDEBUG
        assertHasElements();
#endif
        return m_dataPtr[0];
    }

    __host__ __device__ T back() const
    {
#ifndef NDEBUG
        assertHasElements();
#endif
        return m_dataPtr[m_numberOfElements-1];
    }

    void copyFromHost(const T* dataPointer, size_t numberOfElements);

    void copyToHost(T* dataPointer, size_t numberOfElements) const;

    void copyFromHost(const std::vector<T>& data);

    void copyToHost(std::vector<T>& data) const;

    __host__ __device__ size_t size() const{
        return m_numberOfElements;
    }

    std::vector<T> asStdVector() const;
    class iterator {
    public:
        // Iterator typedefs
        using iterator_category = std::forward_iterator_tag;
        using difference_type = std::ptrdiff_t;
        using value_type = T;
        using pointer = T*;
        using reference = T&;

        __host__ __device__ iterator(T* ptr) : m_ptr(ptr) {}

        __host__ __device__ reference operator*() const {
            return *m_ptr;
        }

        __host__ __device__ iterator& operator++() {
            ++m_ptr;
            return *this;
        }

        __host__ __device__ iterator operator++(int) {
            iterator tmp = *this;
            ++m_ptr;
            return tmp;
        }

        __host__ __device__ iterator& operator--() {
            --m_ptr;
            return *this;
        }

        __host__ __device__ iterator operator--(int) {
            iterator tmp = *this;
            --m_ptr;
            return tmp;
        }

        __host__ __device__ bool operator!=(const iterator& other) const {
            return !(m_ptr == other.m_ptr);
        }

        __host__ __device__ bool operator==(const iterator& other) const {
            return m_ptr == other.m_ptr;
        }

        __host__ __device__ difference_type operator-(const iterator& other) const {
            return std::distance(other.m_ptr, m_ptr);
        }

        __host__ __device__ iterator operator-(difference_type n) const {
            return iterator(m_ptr-n);
        }

        __host__ __device__ iterator operator+(difference_type n) const {
            return iterator(m_ptr + n);
        }

        __host__ __device__ bool operator<(const iterator& other) const {
            return m_ptr < other.m_ptr;
        }

        __host__ __device__ bool operator>(const iterator& other) const {
            return m_ptr > other.m_ptr;
        }
 
    private:
        // Pointer to the current element
        T* m_ptr;
    };

    __host__ __device__ iterator begin(){
        return iterator(m_dataPtr);
    }

    __host__ __device__ iterator begin() const {
        return iterator(m_dataPtr);
    }

    __host__ __device__ iterator end(){
        return iterator(m_dataPtr + m_numberOfElements);
    }

    __host__ __device__ iterator end() const {
        return iterator(m_dataPtr + m_numberOfElements);
    }
 
private:
    T* m_dataPtr;
    size_t m_numberOfElements;

    __host__ __device__ void assertSameSize(const GpuView<T>& other) const
    {
        assertSameSize(other.m_numberOfElements);
    }
    __host__ __device__ void assertSameSize(size_t size) const
    {
#if OPM_IS_INSIDE_DEVICE_FUNCTION
        // TODO: find a better way to handle exceptions in kernels, this will possibly be printed many times
        assert(size == m_numberOfElements && "Views did not have the same size");
#else
        if (size != m_numberOfElements) {
            OPM_THROW(std::invalid_argument,
                    fmt::format(fmt::runtime("Given view has {}, while this View has {}."), size, m_numberOfElements));
        }
#endif
    }

    __host__ __device__ void assertHasElements() const
    {
#if OPM_IS_INSIDE_DEVICE_FUNCTION
        // TODO: find a better way to handle exceptions in kernels, this will possibly be printed many times
        assert(m_numberOfElements > 0 && "View has 0 elements");
#else
        if (m_numberOfElements <= 0) {
            OPM_THROW(std::invalid_argument, "View has 0 elements");
        }
#endif
    }

    __host__ __device__ void assertInRange(size_t idx) const
    {
#if OPM_IS_INSIDE_DEVICE_FUNCTION
        // TODO: find a better way to handle exceptions in kernels, this will possibly be printed many times
        assert(idx < m_numberOfElements && "The index provided was not in the range [0, buffersize-1]");
#else
        if (idx >= m_numberOfElements) {
            OPM_THROW(std::invalid_argument,
                    fmt::format(fmt::runtime("The index provided was not in the range [0, buffersize-1]")));
        }
#endif
    }
};
 
} // namespace Opm::gpuistl
 
#endif