array.hpp 3.38 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
// SPDX-License-Identifier: MIT
Illia Silin's avatar
Illia Silin committed
2
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3

Chao Liu's avatar
Chao Liu committed
4
5
6
#pragma once

#include <initializer_list>
7

Chao Liu's avatar
Chao Liu committed
8
#include "functional2.hpp"
Chao Liu's avatar
Chao Liu committed
9
#include "sequence.hpp"
10

11
12
namespace ck {

Chao Liu's avatar
Chao Liu committed
13
template <typename TData, index_t NSize>
14
15
struct Array
{
Chao Liu's avatar
Chao Liu committed
16
    using type      = Array;
Chao Liu's avatar
Chao Liu committed
17
    using data_type = TData;
18

Chao Liu's avatar
Chao Liu committed
19
    TData mData[NSize];
20

Chao Liu's avatar
Chao Liu committed
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
    __host__ __device__ constexpr Array() : mData{} {}

    __host__ __device__ constexpr Array(std::initializer_list<TData> ilist)
    {
        constexpr index_t list_size = std::initializer_list<TData>{}.size();

        static_assert(list_size <= NSize, "out of bound");

        index_t i   = 0;
        TData vlast = TData{};

        for(const TData& val : ilist)
        {
            mData[i] = val;
            vlast    = val;
            ++i;
        }

        for(; i < NSize; ++i)
        {
            mData[i] = vlast;
        }
    }

Chao Liu's avatar
Chao Liu committed
45
46
    __host__ __device__ static constexpr index_t Size() { return NSize; }

Chao Liu's avatar
Chao Liu committed
47
48
49
50
51
52
53
54
55
56
57
58
    template <index_t I>
    __host__ __device__ constexpr const TData& At() const
    {
        return mData[I];
    }

    template <index_t I>
    __host__ __device__ constexpr TData& At()
    {
        return mData[I];
    }

Chao Liu's avatar
Chao Liu committed
59
    __host__ __device__ constexpr const TData& At(index_t i) const { return mData[i]; }
60

Chao Liu's avatar
Chao Liu committed
61
62
    __host__ __device__ constexpr TData& At(index_t i) { return mData[i]; }

Chao Liu's avatar
Chao Liu committed
63
    __host__ __device__ constexpr const TData& operator[](index_t i) const { return mData[i]; }
Chao Liu's avatar
Chao Liu committed
64

Chao Liu's avatar
Chao Liu committed
65
    __host__ __device__ constexpr TData& operator()(index_t i) { return mData[i]; }
Chao Liu's avatar
Chao Liu committed
66

Chao Liu's avatar
Chao Liu committed
67
    template <typename T>
Chao Liu's avatar
Chao Liu committed
68
    __host__ __device__ constexpr auto operator=(const T& a)
Chao Liu's avatar
Chao Liu committed
69
    {
Chao Liu's avatar
Chao Liu committed
70
        static_assert(T::Size() == Size(), "wrong! size not the same");
Chao Liu's avatar
Chao Liu committed
71

Chao Liu's avatar
Chao Liu committed
72
73
74
75
        for(index_t i = 0; i < NSize; ++i)
        {
            mData[i] = a[i];
        }
Chao Liu's avatar
Chao Liu committed
76
77
78

        return *this;
    }
Chao Liu's avatar
Chao Liu committed
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97

    __host__ __device__ static constexpr bool IsStatic() { return is_static_v<TData>; };

    __host__ __device__ void Print() const
    {
        printf("Array{size: %d, data: [", NSize);

        for(index_t i = 0; i < NSize; i++)
        {
            print(mData[i]);

            if(i < NSize - 1)
            {
                printf(", ");
            }
        }

        printf("]}");
    }
Chao Liu's avatar
Chao Liu committed
98
99
};

Chao Liu's avatar
Chao Liu committed
100
101
102
// empty Array
template <typename TData>
struct Array<TData, 0>
Chao Liu's avatar
Chao Liu committed
103
{
Chao Liu's avatar
Chao Liu committed
104
105
    using type      = Array;
    using data_type = TData;
Chao Liu's avatar
Chao Liu committed
106

Chao Liu's avatar
Chao Liu committed
107
108
    __host__ __device__ constexpr Array() {}

Chao Liu's avatar
Chao Liu committed
109
    __host__ __device__ static constexpr index_t Size() { return 0; }
Chao Liu's avatar
Chao Liu committed
110
111
112
113

    __host__ __device__ static constexpr bool IsStatic() { return is_static_v<TData>; };

    __host__ __device__ void Print() const { printf("Array{size: 0, data: []}"); }
Chao Liu's avatar
Chao Liu committed
114
115
};

Chao Liu's avatar
Chao Liu committed
116
117
template <typename T, typename... Xs>
__host__ __device__ constexpr auto make_array(Xs&&... xs)
Chao Liu's avatar
Chao Liu committed
118
{
Chao Liu's avatar
Chao Liu committed
119
120
121
122
123
124
125
126
127
128
129
130
    using data_type = remove_cvref_t<T>;

    return Array<data_type, sizeof...(Xs)>{std::forward<Xs>(xs)...};
}

template <typename F, index_t N>
__host__ __device__ constexpr auto generate_array(F&& f, Number<N>)
{
    using T = remove_cvref_t<decltype(f(Number<0>{}))>;

    return unpack([&f](auto&&... is) { return Array<T, N>{f(is)...}; },
                  typename arithmetic_sequence_gen<0, N, 1>::type{});
Chao Liu's avatar
Chao Liu committed
131
}
Chao Liu's avatar
Chao Liu committed
132

Chao Liu's avatar
Chao Liu committed
133
134
template <typename T, index_t N, typename X>
__host__ __device__ constexpr auto to_array(const X& x)
Chao Liu's avatar
Chao Liu committed
135
{
Chao Liu's avatar
Chao Liu committed
136
137
138
139
140
141
142
    STATIC_ASSERT(N <= X::Size(), "");

    Array<T, N> arr;

    static_for<0, N, 1>{}([&x, &arr](auto i) { arr(i) = x[i]; });

    return arr;
143
}
144

145
} // namespace ck