slice_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

namespace etl {

template <typename T>
requires (!fast_slice_view_able<T>)
struct slice_view<T>
        : assignable<slice_view<T>, value_t<T>>, value_testable<slice_view<T>>, iterable<slice_view<T>, fast_slice_view_able<T>> {
    using this_type            = slice_view<T>;                                                        
    using iterable_base_type   = iterable<this_type, fast_slice_view_able<T>>;                         
    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>;                                       

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

private:
    T sub;              
    const size_t first; 
    const size_t last;  

    friend struct etl_traits<slice_view>;

public:
    slice_view(sub_type sub, size_t first, size_t last) : sub(sub), first(first), last(last) {}

    const_return_type operator[](size_t j) const {
        if constexpr (decay_traits<sub_type>::storage_order == order::RowMajor) {
            return sub[first * (etl::size(sub) / dim<0>(sub)) + j];
        } else {
            const auto sa = dim<0>(sub);
            const auto ss = (etl::size(sub) / sa);
            return sub[(j % ss) * sa + j / ss + first];
        }
    }

    return_type operator[](size_t j) {
        if constexpr (decay_traits<sub_type>::storage_order == order::RowMajor) {
            return sub[first * (etl::size(sub) / dim<0>(sub)) + j];
        } else {
            const auto sa = dim<0>(sub);
            const auto ss = (etl::size(sub) / sa);
            return sub[(j % ss) * sa + j / ss + first];
        }
    }

    value_type read_flat(size_t j) const noexcept {
        if constexpr (decay_traits<sub_type>::storage_order == order::RowMajor) {
            return sub.read_flat(first * (etl::size(sub) / dim<0>(sub)) + j);
        } else {
            const auto sa = dim<0>(sub);
            const auto ss = (etl::size(sub) / sa);
            return sub.read_flat((j % ss) * sa + j / ss + first);
        }
    }

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

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

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

    template <typename V = default_vec>
    auto load(size_t x) const noexcept {
        return sub.template loadu<V>(x + first * (etl::size(sub) / etl::dim<0>(sub)));
    }

    template <typename V = default_vec>
    auto loadu(size_t x) const noexcept {
        return sub.template loadu<V>(x + first * (etl::size(sub) / etl::dim<0>(sub)));
    }

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

    // 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.visit(visitor);
    }

    void ensure_cpu_up_to_date() const {
        // Need to ensure sub value
        sub.ensure_cpu_up_to_date();
    }

    void ensure_gpu_up_to_date() const {
        // Need to ensure both LHS and RHS
        sub.ensure_gpu_up_to_date();
    }
};

template <typename T>
requires (fast_slice_view_able<T>)
struct slice_view<T>
        : assignable<slice_view<T>, value_t<T>>, value_testable<slice_view<T>>, iterable<slice_view<T>, true> {
    using this_type            = slice_view<T>;                                                        
    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             = memory_type;                                                          
    using const_iterator       = const_memory_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:
    T sub;                 
    const size_t first;    
    const size_t last;     
    const size_t sub_size; 

    mutable memory_type memory; 

    mutable bool cpu_up_to_date; 
    mutable bool gpu_up_to_date; 

    friend struct etl_traits<slice_view>;

public:
    slice_view(sub_type sub, size_t first, size_t last) : sub(sub), first(first), last(last), sub_size((etl::size(sub) / etl::dim<0>(sub)) * (last - first)) {
        // 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.memory_start() + first * (etl::size(this->sub) / etl::dim<0>(this->sub));

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

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

    ~slice_view() {
        if (this->memory) {
            // Propagate the status on the parent
            if (!this->cpu_up_to_date) {
                if (sub.is_gpu_up_to_date()) {
                    sub.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.is_cpu_up_to_date()) {
                    sub.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>
    const_return_type operator()(size_t i, S... args) const requires(sizeof...(S) + 1 == decay_traits<sub_type>::dimensions()) {
        ensure_cpu_up_to_date();
        return memory[dyn_index(*this, i, args...)];
    }

    template <typename... S>
    return_type operator()(size_t i, S... args) requires(sizeof...(S) + 1 == decay_traits<sub_type>::dimensions()) {
        ensure_cpu_up_to_date();
        invalidate_gpu();
        return memory[dyn_index(*this, i, args...)];
    }

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

    template <typename V = default_vec>
    auto load(size_t x) const noexcept {
        return V::loadu(memory + 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 slice view is aligned (at compile-time), use stream store here
        return V::storeu(memory + x, in);
    }

    template <typename V = default_vec>
    auto 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.alias(rhs);
        }
    }

    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;
    }

    memory_type memory_start() noexcept {
        return memory;
    }

    const_memory_type memory_start() const noexcept {
        return memory;
    }

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

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

    // 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.visit(visitor);
    }

    // GPU functions

    value_type* gpu_memory() const noexcept {
        return sub.gpu_memory() + first * (etl::size(sub) / etl::dim<0>(sub));
    }

    void gpu_evict() const noexcept {
        sub.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.ensure_gpu_allocated();
    }

    void ensure_gpu_up_to_date() const {
        sub.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:slice: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:slice: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.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;
    }
};

template <typename T>
struct etl_traits<etl::slice_view<T>> {
    using expr_t     = etl::slice_view<T>;                          
    using sub_expr_t = std::decay_t<T>;                             
    using sub_traits = etl_traits<sub_expr_t>;                      
    using value_type = typename etl_traits<sub_expr_t>::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        = false;                      
    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      = fast_slice_view_able<T>;    
    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 = is_direct;                  
    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) {
        return (sub_traits::size(v.sub) / sub_traits::dim(v.sub, 0)) * (v.last - v.first);
    }

    static size_t dim(const expr_t& v, size_t d) {
        if (d == 0) {
            return v.last - v.first;
        } else {
            return sub_traits::dim(v.sub, d);
        }
    }

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

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

} //end of namespace etl