sub_view.hpp

Go to the documentation of this file.

//=======================================================================
// Copyright (c) 2014-2023 Baptiste Wicht
// Distributed under the terms of the MIT License.
// (See accompanying file LICENSE or copy at
//  http://opensource.org/licenses/MIT)
//=======================================================================

#pragma once

#ifdef ETL_CUDA
#include "etl/impl/cublas/cuda.hpp"
#endif

namespace etl {

template <matrix T, bool Aligned>
requires(!fast_sub_view_able<T>)
struct sub_view<T, Aligned> final : iterable<sub_view<T, Aligned>, false>,
                                                                              assignable<sub_view<T, Aligned>, value_t<T>>,
                                                                              value_testable<sub_view<T, Aligned>>,
                                                                              inplace_assignable<sub_view<T, Aligned>> {
    using this_type            = sub_view<T, Aligned>;                                                 
    using iterable_base_type   = iterable<this_type, false>;                                           
    using assignable_base_type = assignable<this_type, value_t<T>>;                                    
    using sub_type             = T;                                                                    
    using value_type           = value_t<sub_type>;                                                    
    using memory_type          = memory_t<sub_type>;                                                   
    using const_memory_type    = const_memory_t<sub_type>;                                             
    using return_type          = return_helper<sub_type, decltype(std::declval<sub_type>()[0])>;       
    using const_return_type    = const_return_helper<sub_type, decltype(std::declval<sub_type>()[0])>; 
    using iterator             = etl::iterator<this_type>;                                             
    using const_iterator       = etl::iterator<const this_type>;                                       

    template <typename V = default_vec>
    using vec_type = typename V::template vec_type<value_type>;

    using assignable_base_type::operator=;
    using iterable_base_type::begin;
    using iterable_base_type::end;

private:
    sub_type sub_expr;       
    const size_t i;          
    const size_t sub_offset; 

    friend struct etl_traits<sub_view>;

    static constexpr order storage_order = decay_traits<sub_type>::storage_order; 

public:
    sub_view(sub_type sub_expr, size_t i) : sub_expr(sub_expr), i(i), sub_offset(i * subsize(sub_expr)) {}

    const_return_type operator[](size_t j) const {
        return storage_order == order::RowMajor ? sub_expr[sub_offset + j] : sub_expr[i + dim<0>(sub_expr) * j];
    }

    return_type operator[](size_t j) {
        return storage_order == order::RowMajor ? sub_expr[sub_offset + j] : sub_expr[i + dim<0>(sub_expr) * j];
    }

    value_type read_flat(size_t j) const noexcept {
        return storage_order == order::RowMajor ? sub_expr.read_flat(sub_offset + j) : sub_expr.read_flat(i + dim<0>(sub_expr) * j);
    }

    template <typename... S>
    ETL_STRONG_INLINE(const_return_type)
    operator()(S... args) const requires(sizeof...(S) + 1 == decay_traits<sub_type>::dimensions()) {
        return sub_expr(i, static_cast<size_t>(args)...);
    }

    template <typename... S>
    ETL_STRONG_INLINE(return_type)
    operator()(S... args) requires(sizeof...(S) + 1 == decay_traits<sub_type>::dimensions()) {
        return sub_expr(i, static_cast<size_t>(args)...);
    }

    auto operator()(size_t x) const requires(decay_traits<sub_type>::dimensions() > 2) {
        return sub(*this, x);
    }

    template <typename V = default_vec>
    void store(vec_type<V> in, size_t x) noexcept {
        return sub_expr.template storeu<V>(in, x + sub_offset);
    }

    template <typename V = default_vec>
    void storeu(vec_type<V> in, size_t x) noexcept {
        return sub_expr.template storeu<V>(in, x + sub_offset);
    }

    template <typename V = default_vec>
    void stream(vec_type<V> in, size_t x) noexcept {
        return sub_expr.template storeu<V>(in, x + sub_offset);
    }

    template <typename V = default_vec>
    ETL_STRONG_INLINE(vec_type<V>)
    load(size_t x) const noexcept {
        return sub_expr.template loadu<V>(x + sub_offset);
    }

    template <typename V = default_vec>
    ETL_STRONG_INLINE(vec_type<V>)
    loadu(size_t x) const noexcept {
        return sub_expr.template loadu<V>(x + sub_offset);
    }

    template <typename E>
    bool alias(const E& rhs) const noexcept {
        return sub_expr.alias(rhs);
    }

    size_t& unsafe_dimension_access(size_t x) {
        return sub_expr.unsafe_dimension_access(x + 1);
    }

    // Assignment functions

    template <typename L>
    void assign_to(L&& lhs) const {
        std_assign_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_add_to(L&& lhs) const {
        std_add_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_sub_to(L&& lhs) const {
        std_sub_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_mul_to(L&& lhs) const {
        std_mul_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_div_to(L&& lhs) const {
        std_div_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_mod_to(L&& lhs) const {
        std_mod_evaluate(*this, std::forward<L>(lhs));
    }

    // Internals

    void visit(detail::evaluator_visitor& visitor) const {
        sub_expr.visit(visitor);
    }

    void ensure_cpu_up_to_date() const {
        // The sub value must be ensured
        sub_expr.ensure_cpu_up_to_date();
    }

    void ensure_gpu_up_to_date() const {
        // The sub value must be ensured
        sub_expr.ensure_gpu_up_to_date();
    }

    friend std::ostream& operator<<(std::ostream& os, const sub_view& v) {
        return os << "sub(" << v.sub_expr << ", " << v.i << ")";
    }
};

template <matrix T, bool Aligned>
requires(fast_sub_view_able<T>)
struct sub_view<T, Aligned> : iterable<sub_view<T, Aligned>, true>,
                                                                       assignable<sub_view<T, Aligned>, value_t<T>>,
                                                                       value_testable<sub_view<T, Aligned>>,
                                                                       inplace_assignable<sub_view<T, Aligned>> {
    static_assert(decay_traits<T>::storage_order == order::RowMajor, "sub_view<T, true> should only be done with RowMajor");

    using this_type            = sub_view<T, Aligned>;                                                 
    using iterable_base_type   = iterable<this_type, true>;                                            
    using assignable_base_type = assignable<this_type, value_t<T>>;                                    
    using sub_type             = T;                                                                    
    using value_type           = value_t<sub_type>;                                                    
    using memory_type          = memory_t<sub_type>;                                                   
    using const_memory_type    = const_memory_t<sub_type>;                                             
    using return_type          = return_helper<sub_type, decltype(std::declval<sub_type>()[0])>;       
    using const_return_type    = const_return_helper<sub_type, decltype(std::declval<sub_type>()[0])>; 
    using iterator             = value_type*;                                                          
    using const_iterator       = const value_type*;                                                    

    template <typename V = default_vec>
    using vec_type = typename V::template vec_type<value_type>;

    using iterable_base_type::begin;
    using iterable_base_type::end;
    using assignable_base_type::operator=;

private:
    T sub_expr;            
    const size_t i;        
    const size_t sub_size; 

    static constexpr size_t n_dimensions = decay_traits<sub_type>::dimensions() - 1; 
    static constexpr order storage_order = decay_traits<sub_type>::storage_order;    

    mutable memory_type memory; 

    mutable bool cpu_up_to_date; 
    mutable bool gpu_up_to_date; 

    friend struct etl_traits<sub_view>;

public:
    sub_view(sub_type sub_expr, size_t i) : sub_expr(sub_expr), i(i), sub_size(subsize(sub_expr)) {
        // Accessing the memory through fast sub views means evaluation
        if constexpr (decay_traits<sub_type>::is_temporary) {
            standard_evaluator::pre_assign_rhs(*this);
        }

        this->memory = this->sub_expr.memory_start() + i * sub_size;

        // A sub view inherits the CPU/GPU from parent
        this->cpu_up_to_date = this->sub_expr.is_cpu_up_to_date();
        this->gpu_up_to_date = this->sub_expr.is_gpu_up_to_date();

        cpp_assert(this->memory, "Invalid memory");
    }

    ~sub_view() {
        if (this->memory) {
            // Propagate the status on the parent
            if (!this->cpu_up_to_date) {
                if (sub_expr.is_gpu_up_to_date()) {
                    sub_expr.invalidate_cpu();
                } else {
                    // If the GPU is not up to date, cannot invalidate the CPU too
                    ensure_cpu_up_to_date();
                }
            }

            if (!this->gpu_up_to_date) {
                if (sub_expr.is_cpu_up_to_date()) {
                    sub_expr.invalidate_gpu();
                } else {
                    // If the GPU is not up to date, cannot invalidate the CPU too
                    ensure_gpu_up_to_date();
                }
            }
        }
    }

    const_return_type operator[](size_t j) const {
        ensure_cpu_up_to_date();
        return memory[j];
    }

    return_type operator[](size_t j) {
        ensure_cpu_up_to_date();
        invalidate_gpu();
        return memory[j];
    }

    value_type read_flat(size_t j) const noexcept {
        ensure_cpu_up_to_date();
        return memory[j];
    }

    template <typename... S>
    ETL_STRONG_INLINE(const_return_type)
    operator()(S... args) const requires (sizeof...(S) == n_dimensions) {
        ensure_cpu_up_to_date();
        return memory[dyn_index(*this, args...)];
    }

    template <typename... S>
    ETL_STRONG_INLINE(return_type)
    operator()(S... args) requires (sizeof...(S) == n_dimensions) {
        ensure_cpu_up_to_date();
        return memory[dyn_index(*this, args...)];
    }

    auto operator()(size_t x) const requires(n_dimensions > 1) {
        return sub(*this, x);
    }

    template <typename V = default_vec>
    void store(vec_type<V> in, size_t x) noexcept {
        return V::storeu(memory + x, in);
    }

    template <typename V = default_vec>
    void storeu(vec_type<V> in, size_t x) noexcept {
        return V::storeu(memory + x, in);
    }

    template <typename V = default_vec>
    void stream(vec_type<V> in, size_t x) noexcept {
        //TODO If the sub view is aligned (at compile-time), use stream store here
        return V::storeu(memory + x, in);
    }

    template <typename V = default_vec>
    vec_type<V> load(size_t x) const noexcept {
        return V::loadu(memory + x);
    }

    template <typename V = default_vec>
    vec_type<V> loadu(size_t x) const noexcept {
        return V::loadu(memory + x);
    }

    template <typename E>
    bool alias(const E& rhs) const noexcept {
        if constexpr (is_dma<E>) {
            return memory_alias(memory_start(), memory_end(), rhs.memory_start(), rhs.memory_end());
        } else {
            return sub_expr.alias(rhs);
        }
    }

    ETL_STRONG_INLINE(memory_type) memory_start() noexcept {
        return memory;
    }

    ETL_STRONG_INLINE(const_memory_type) memory_start() const noexcept {
        return memory;
    }

    ETL_STRONG_INLINE(memory_type) memory_end() noexcept {
        return memory + sub_size;
    }

    ETL_STRONG_INLINE(const_memory_type) memory_end() const noexcept {
        return memory + sub_size;
    }

    size_t& unsafe_dimension_access(size_t x) {
        return sub_expr.unsafe_dimension_access(x + 1);
    }

    template <size_t... I>
    std::array<size_t, decay_traits<this_type>::dimensions()> dim_array(std::index_sequence<I...>) const {
        return {{decay_traits<this_type>::dim(*this, I)...}};
    }

    // Assignment functions

    template <typename L>
    void assign_to(L&& lhs) const {
        std_assign_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_add_to(L&& lhs) const {
        std_add_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_sub_to(L&& lhs) const {
        std_sub_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_mul_to(L&& lhs) const {
        std_mul_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_div_to(L&& lhs) const {
        std_div_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename L>
    void assign_mod_to(L&& lhs) const {
        std_mod_evaluate(*this, std::forward<L>(lhs));
    }

    template <typename Y>
    auto& gpu_compute_hint([[maybe_unused]] Y& y) {
        this->ensure_gpu_up_to_date();
        return *this;
    }

    template <typename Y>
    const auto& gpu_compute_hint([[maybe_unused]] Y& y) const {
        this->ensure_gpu_up_to_date();
        return *this;
    }

    // Internals

    void visit(detail::evaluator_visitor& visitor) const {
        sub_expr.visit(visitor);
    }

    // GPU functions

    value_type* gpu_memory() const noexcept {
        return sub_expr.gpu_memory() + i * sub_size;
    }

    void gpu_evict() const noexcept {
        sub_expr.gpu_evict();
    }

    void invalidate_cpu() const noexcept {
        this->cpu_up_to_date = false;
    }

    void invalidate_gpu() const noexcept {
        this->gpu_up_to_date = false;
    }

    void validate_cpu() const noexcept {
        this->cpu_up_to_date = true;
    }

    void validate_gpu() const noexcept {
        this->gpu_up_to_date = true;
    }

    void ensure_gpu_allocated() const {
        // Allocate is done by the sub
        sub_expr.ensure_gpu_allocated();
    }

    void ensure_gpu_up_to_date() const {
        sub_expr.ensure_gpu_allocated();

#ifdef ETL_CUDA
        if (!this->gpu_up_to_date) {
            cuda_check_assert(cudaMemcpy(const_cast<std::remove_const_t<value_type>*>(gpu_memory()),
                                         const_cast<std::remove_const_t<value_type>*>(memory_start()), sub_size * sizeof(value_type), cudaMemcpyHostToDevice));

            this->gpu_up_to_date = true;

            inc_counter("gpu:sub:cpu_to_gpu");
        }
#endif
    }

    void ensure_cpu_up_to_date() const {
#ifdef ETL_CUDA
        if (!this->cpu_up_to_date) {
            cuda_check_assert(cudaMemcpy(const_cast<std::remove_const_t<value_type>*>(memory_start()),
                                         const_cast<std::remove_const_t<value_type>*>(gpu_memory()), sub_size * sizeof(value_type), cudaMemcpyDeviceToHost));

            inc_counter("gpu:sub:gpu_to_cpu");
        }
#endif

        this->cpu_up_to_date = true;
    }

    void gpu_copy_from([[maybe_unused]] const value_type* new_gpu_memory) const {
        cpp_assert(sub_expr.gpu_memory(), "GPU must be allocated before copy");

#ifdef ETL_CUDA
        cuda_check_assert(cudaMemcpy(const_cast<std::remove_const_t<value_type>*>(gpu_memory()), const_cast<std::remove_const_t<value_type>*>(new_gpu_memory),
                                     sub_size * sizeof(value_type), cudaMemcpyDeviceToDevice));
#endif

        gpu_up_to_date = true;
        cpu_up_to_date = false;
    }

    bool is_cpu_up_to_date() const noexcept {
        return cpu_up_to_date;
    }

    bool is_gpu_up_to_date() const noexcept {
        return gpu_up_to_date;
    }

    friend std::ostream& operator<<(std::ostream& os, const sub_view& v) {
        return os << "sub(" << v.sub_expr << ", " << v.i << ")";
    }
};

template <typename T, bool Aligned>
struct etl_traits<etl::sub_view<T, Aligned>> {
    using expr_t     = etl::sub_view<T, Aligned>;       
    using sub_expr_t = std::decay_t<T>;                 
    using sub_traits = etl_traits<sub_expr_t>;          
    using value_type = typename sub_traits::value_type; 

    static constexpr bool is_etl         = true;                       
    static constexpr bool is_transformer = false;                      
    static constexpr bool is_view        = true;                       
    static constexpr bool is_magic_view  = false;                      
    static constexpr bool is_fast        = sub_traits::is_fast;        
    static constexpr bool is_linear      = sub_traits::is_linear;      
    static constexpr bool is_thread_safe = sub_traits::is_thread_safe; 
    static constexpr bool is_value       = false;                      
    static constexpr bool is_direct =
        sub_traits::is_direct && sub_traits::storage_order == order::RowMajor; 
    static constexpr bool is_generator   = false;                              
    static constexpr bool is_padded      = false;                              
    static constexpr bool is_aligned     = false;                              
    static constexpr bool is_temporary   = sub_traits::is_temporary;           
    static constexpr bool gpu_computable = fast_sub_view_able<T>;              
    static constexpr order storage_order = sub_traits::storage_order;          

    template <vector_mode_t V>
    static constexpr bool vectorizable = sub_traits::template vectorizable<V>&& storage_order == order::RowMajor;

    static size_t size(const expr_t& v) noexcept {
        return sub_traits::size(v.sub_expr) / sub_traits::dim(v.sub_expr, 0);
    }

    static size_t dim(const expr_t& v, size_t d) noexcept {
        return sub_traits::dim(v.sub_expr, d + 1);
    }

    static constexpr size_t size() noexcept {
        return sub_traits::size() / sub_traits::template dim<0>();
    }

    template <size_t D>
    static constexpr size_t dim() noexcept {
        return sub_traits::template dim<D + 1>();
    }

    static constexpr size_t dimensions() noexcept {
        return sub_traits::dimensions() - 1;
    }

    static constexpr int complexity() noexcept {
        return -1;
    }
};

} //end of namespace etl