"demo/git@developer.sourcefind.cn:modelzoo/sam2_pytorch.git" did not exist on "3af09475e690c3ab8b5678c05e8e8939ac2c27fd"
array.hpp 11.1 KB
Newer Older
1
2
3
#ifndef CK_ARRAY_HPP
#define CK_ARRAY_HPP

Chao Liu's avatar
Chao Liu committed
4
#include "sequence.hpp"
Chao Liu's avatar
Chao Liu committed
5
#include "functional2.hpp"
6

7
8
namespace ck {

Chao Liu's avatar
Chao Liu committed
9
template <typename TData, index_t NSize>
10
11
struct Array
{
Chao Liu's avatar
Chao Liu committed
12
    using type      = Array<TData, NSize>;
Chao Liu's avatar
Chao Liu committed
13
    using data_type = TData;
14

Chao Liu's avatar
Chao Liu committed
15
    index_t mData[NSize];
16

Chao Liu's avatar
Chao Liu committed
17
    __host__ __device__ explicit constexpr Array() {}
18

Chao Liu's avatar
Chao Liu committed
19
    template <typename X, typename... Xs>
Chao Liu's avatar
Chao Liu committed
20
    __host__ __device__ constexpr Array(X x, Xs... xs)
Chao Liu's avatar
Chao Liu committed
21
        : mData{static_cast<TData>(x), static_cast<TData>(xs)...}
22
    {
Chao Liu's avatar
Chao Liu committed
23
        static_assert(sizeof...(Xs) + 1 == NSize, "wrong! size");
24
25
    }

Chao Liu's avatar
Chao Liu committed
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
#if 0
    template <typename T>
    __host__ __device__ explicit constexpr Array(const T& x)
    {
        static_assert(T::Size() == NSize, "wrong! size");

        static_for<0, NSize, 1>{}([&](auto i){
           mData[i] = x.At(i); 
        })
    }
#endif

    __host__ __device__ static constexpr index_t Size() { return NSize; }

    __host__ __device__ static constexpr index_t GetSize() { return Size(); }
41

Chao Liu's avatar
Chao Liu committed
42
    template <index_t I>
Chao Liu's avatar
Chao Liu committed
43
    __host__ __device__ constexpr const TData& At(Number<I>) const
Chao Liu's avatar
Chao Liu committed
44
    {
Chao Liu's avatar
Chao Liu committed
45
46
        static_assert(I < NSize, "wrong!");

Chao Liu's avatar
Chao Liu committed
47
48
49
50
        return mData[I];
    }

    template <index_t I>
Chao Liu's avatar
Chao Liu committed
51
    __host__ __device__ constexpr TData& At(Number<I>)
Chao Liu's avatar
Chao Liu committed
52
    {
Chao Liu's avatar
Chao Liu committed
53
54
        static_assert(I < NSize, "wrong!");

Chao Liu's avatar
Chao Liu committed
55
56
57
        return mData[I];
    }

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

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

    template <typename I>
    __host__ __device__ constexpr const TData& operator[](I i) const
Chao Liu's avatar
Chao Liu committed
64
    {
Chao Liu's avatar
Chao Liu committed
65
66
        return At(i);
    }
Chao Liu's avatar
Chao Liu committed
67

Chao Liu's avatar
Chao Liu committed
68
69
70
71
    template <typename I>
    __host__ __device__ constexpr TData& operator()(I i)
    {
        return At(i);
Chao Liu's avatar
Chao Liu committed
72
73
    }

Chao Liu's avatar
Chao Liu committed
74
75
76
77
78
79
80
    template <typename T>
    __host__ __device__ constexpr type& operator=(const T& x)
    {
        static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = x[i]; });

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

    struct lambda_PushBack // emulate constexpr lambda
    {
        const Array<TData, NSize>& old_array;
        Array<TData, NSize + 1>& new_array;

        __host__ __device__ constexpr lambda_PushBack(const Array<TData, NSize>& old_array_,
                                                      Array<TData, NSize + 1>& new_array_)
            : old_array(old_array_), new_array(new_array_)
        {
        }

        template <index_t I>
        __host__ __device__ constexpr void operator()(Number<I>) const
        {
Chao Liu's avatar
Chao Liu committed
96
            new_array(Number<I>{}) = old_array[I];
Chao Liu's avatar
Chao Liu committed
97
98
99
        }
    };

Chao Liu's avatar
Chao Liu committed
100
    __host__ __device__ constexpr auto PushBack(TData x) const
101
102
103
    {
        Array<TData, NSize + 1> new_array;

Chao Liu's avatar
Chao Liu committed
104
        static_for<0, NSize, 1>{}(lambda_PushBack(*this, new_array));
105

Chao Liu's avatar
Chao Liu committed
106
        new_array(Number<NSize>{}) = x;
107
108
109

        return new_array;
    }
110
};
111

Chao Liu's avatar
Chao Liu committed
112
// Arr: Array
Chao Liu's avatar
Chao Liu committed
113
// Picks: Sequence<...>
Chao Liu's avatar
Chao Liu committed
114
template <typename Arr, typename Picks>
Chao Liu's avatar
Chao Liu committed
115
struct ArrayElementPicker
Chao Liu's avatar
Chao Liu committed
116
{
Chao Liu's avatar
Chao Liu committed
117
    using type      = ArrayElementPicker;
Chao Liu's avatar
Chao Liu committed
118
119
    using data_type = typename Arr::data_type;

Chao Liu's avatar
Chao Liu committed
120
121
122
    __host__ __device__ constexpr ArrayElementPicker() = delete;

    __host__ __device__ explicit constexpr ArrayElementPicker(Arr& array) : mArray{array}
Chao Liu's avatar
Chao Liu committed
123
124
125
126
    {
        constexpr index_t imax =
            accumulate_on_sequence(Picks{}, math::maxer<index_t>{}, Number<0>{});

Chao Liu's avatar
Chao Liu committed
127
        static_assert(imax < Arr::Size(), "wrong! exceeding # array element");
Chao Liu's avatar
Chao Liu committed
128
129
    }

Chao Liu's avatar
Chao Liu committed
130
    __host__ __device__ static constexpr auto Size() { return Picks::Size(); }
Chao Liu's avatar
Chao Liu committed
131
132

    template <index_t I>
Chao Liu's avatar
Chao Liu committed
133
    __host__ __device__ constexpr const data_type& At(Number<I>) const
Chao Liu's avatar
Chao Liu committed
134
    {
Chao Liu's avatar
Chao Liu committed
135
136
137
138
        static_assert(I < Size(), "wrong!");

        constexpr auto IP = Picks{}[I];
        return mArray[IP];
Chao Liu's avatar
Chao Liu committed
139
140
    }

Chao Liu's avatar
Chao Liu committed
141
142
    template <index_t I>
    __host__ __device__ constexpr data_type& At(Number<I>)
Chao Liu's avatar
Chao Liu committed
143
    {
Chao Liu's avatar
Chao Liu committed
144
145
146
147
        static_assert(I < Size(), "wrong!");

        constexpr auto IP = Picks{}[I];
        return mArray(IP);
Chao Liu's avatar
Chao Liu committed
148
149
    }

Chao Liu's avatar
Chao Liu committed
150
151
    template <typename I>
    __host__ __device__ constexpr const data_type& operator[](I i) const
Chao Liu's avatar
Chao Liu committed
152
    {
Chao Liu's avatar
Chao Liu committed
153
        return At(i);
Chao Liu's avatar
Chao Liu committed
154
155
    }

Chao Liu's avatar
Chao Liu committed
156
157
    template <typename I>
    __host__ __device__ constexpr data_type& operator()(I i)
Chao Liu's avatar
Chao Liu committed
158
    {
Chao Liu's avatar
Chao Liu committed
159
        return At(i);
Chao Liu's avatar
Chao Liu committed
160
161
    }

Chao Liu's avatar
Chao Liu committed
162
163
164
165
166
167
168
169
170
    template <typename T>
    __host__ __device__ constexpr type& operator=(const T& a)
    {
        static_for<0, Size(), 1>{}([&](auto i) { operator()(i) = a[i]; });

        return *this;
    }

    Arr& mArray;
Chao Liu's avatar
Chao Liu committed
171
172
};

Chao Liu's avatar
Chao Liu committed
173
template <typename Arr, typename Picks>
Chao Liu's avatar
Chao Liu committed
174
175
176
177
178
__host__ __device__ constexpr auto pick_array_element(Arr& a, Picks)
{
    return ArrayElementPicker<Arr, Picks>(a);
}

Chao Liu's avatar
Chao Liu committed
179
180
181
182
183
184
185
186
187
188
template <typename T>
__host__ __device__ constexpr auto to_array(const T& x)
{
    Array<typename T::data_type, T::Size()> y;

    static_for<0, T::Size(), 1>{}([&](auto i) { y.At(i) = x.At(i); });

    return y;
}

Chao Liu's avatar
Chao Liu committed
189
// TODO: remove this
Chao Liu's avatar
Chao Liu committed
190
191
192
193
194
195
template <index_t... Is>
__host__ __device__ constexpr auto sequence2array(Sequence<Is...>)
{
    return Array<index_t, sizeof...(Is)>{Is...};
}

Chao Liu's avatar
Chao Liu committed
196
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
197
198
__host__ __device__ constexpr auto make_zero_array()
{
Chao Liu's avatar
Chao Liu committed
199
    constexpr auto zero_sequence = typename uniform_sequence_gen<NSize, 0>::type{};
Chao Liu's avatar
Chao Liu committed
200
201
    constexpr auto zero_array    = sequence2array(zero_sequence);
    return zero_array;
Chao Liu's avatar
Chao Liu committed
202
203
}

Chao Liu's avatar
Chao Liu committed
204
template <typename TData, index_t NSize, index_t... IRs>
Chao Liu's avatar
Chao Liu committed
205
__host__ __device__ constexpr auto reorder_array_given_new2old(const Array<TData, NSize>& old_array,
Chao Liu's avatar
Chao Liu committed
206
                                                               Sequence<IRs...> /*new2old*/)
207
208
209
{
    static_assert(NSize == sizeof...(IRs), "NSize not consistent");

Chao Liu's avatar
Chao Liu committed
210
    static_assert(is_valid_sequence_map<Sequence<IRs...>>{}, "wrong! invalid reorder map");
211

212
    return Array<TData, NSize>{old_array[IRs]...};
213
214
}

Chao Liu's avatar
Chao Liu committed
215
template <typename TData, index_t NSize, typename MapOld2New>
Chao Liu's avatar
Chao Liu committed
216
struct lambda_reorder_array_given_old2new
Chao Liu's avatar
Chao Liu committed
217
{
Chao Liu's avatar
Chao Liu committed
218
219
    const Array<TData, NSize>& old_array;
    Array<TData, NSize>& new_array;
Chao Liu's avatar
Chao Liu committed
220

Chao Liu's avatar
Chao Liu committed
221
222
223
    __host__ __device__ constexpr lambda_reorder_array_given_old2new(
        const Array<TData, NSize>& old_array_, Array<TData, NSize>& new_array_)
        : old_array(old_array_), new_array(new_array_)
Chao Liu's avatar
Chao Liu committed
224
225
226
227
228
229
    {
    }

    template <index_t IOldDim>
    __host__ __device__ constexpr void operator()(Number<IOldDim>) const
    {
Chao Liu's avatar
Chao Liu committed
230
        TData old_data = old_array[IOldDim];
Chao Liu's avatar
Chao Liu committed
231

Chao Liu's avatar
Chao Liu committed
232
        constexpr index_t INewDim = MapOld2New::At(Number<IOldDim>{});
Chao Liu's avatar
Chao Liu committed
233

Chao Liu's avatar
Chao Liu committed
234
        new_array(Number<INewDim>{}) = old_data;
Chao Liu's avatar
Chao Liu committed
235
236
237
    }
};

Chao Liu's avatar
Chao Liu committed
238
template <typename TData, index_t NSize, index_t... IRs>
Chao Liu's avatar
Chao Liu committed
239
__host__ __device__ constexpr auto reorder_array_given_old2new(const Array<TData, NSize>& old_array,
Chao Liu's avatar
Chao Liu committed
240
                                                               Sequence<IRs...> /*old2new*/)
Chao Liu's avatar
Chao Liu committed
241
242
243
244
245
{
    Array<TData, NSize> new_array;

    static_assert(NSize == sizeof...(IRs), "NSize not consistent");

Chao Liu's avatar
Chao Liu committed
246
247
    static_assert(is_valid_sequence_map<Sequence<IRs...>>::value, "wrong! invalid reorder map");

Chao Liu's avatar
Chao Liu committed
248
    static_for<0, NSize, 1>{}(
Chao Liu's avatar
Chao Liu committed
249
        lambda_reorder_array_given_old2new<TData, NSize, Sequence<IRs...>>(old_array, new_array));
Chao Liu's avatar
Chao Liu committed
250
251
252

    return new_array;
}
Chao Liu's avatar
Chao Liu committed
253

Chao Liu's avatar
Chao Liu committed
254
template <typename TData, index_t NSize, typename ExtractSeq>
Chao Liu's avatar
Chao Liu committed
255
__host__ __device__ constexpr auto extract_array(const Array<TData, NSize>& old_array, ExtractSeq)
256
257
258
259
260
261
262
{
    Array<TData, ExtractSeq::GetSize()> new_array;

    constexpr index_t new_size = ExtractSeq::GetSize();

    static_assert(new_size <= NSize, "wrong! too many extract");

Chao Liu's avatar
Chao Liu committed
263
    static_for<0, new_size, 1>{}([&](auto I) { new_array(I) = old_array[ExtractSeq::At(I)]; });
264
265
266
267

    return new_array;
}

Chao Liu's avatar
Chao Liu committed
268
269
template <typename F, typename X, typename Y, typename Z> // emulate constepxr lambda for array
// math
Chao Liu's avatar
Chao Liu committed
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
struct lambda_array_math
{
    const F& f;
    const X& x;
    const Y& y;
    Z& z;

    __host__ __device__ constexpr lambda_array_math(const F& f_, const X& x_, const Y& y_, Z& z_)
        : f(f_), x(x_), y(y_), z(z_)
    {
    }

    template <index_t IDim_>
    __host__ __device__ constexpr void operator()(Number<IDim_>) const
    {
        constexpr auto IDim = Number<IDim_>{};
Chao Liu's avatar
Chao Liu committed
286
        z(IDim)             = f(x[IDim], y[IDim]);
Chao Liu's avatar
Chao Liu committed
287
288
289
    }
};

290
// Array = Array + Array
Chao Liu's avatar
Chao Liu committed
291
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
292
__host__ __device__ constexpr auto operator+(Array<TData, NSize> a, Array<TData, NSize> b)
Chao Liu's avatar
Chao Liu committed
293
294
295
{
    Array<TData, NSize> result;

296
    auto f = math::plus<index_t>{};
Chao Liu's avatar
Chao Liu committed
297

Chao Liu's avatar
Chao Liu committed
298
299
300
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
Chao Liu's avatar
Chao Liu committed
301
302
303

    return result;
}
Chao Liu's avatar
Chao Liu committed
304

305
// Array = Array - Array
Chao Liu's avatar
Chao Liu committed
306
template <typename TData, index_t NSize>
307
308
309
310
__host__ __device__ constexpr auto operator-(Array<TData, NSize> a, Array<TData, NSize> b)
{
    Array<TData, NSize> result;

311
    auto f = math::minus<index_t>{};
312

Chao Liu's avatar
Chao Liu committed
313
314
315
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
316
317
318
319

    return result;
}

Chao Liu's avatar
Chao Liu committed
320
// Array += Array
Chao Liu's avatar
Chao Liu committed
321
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
322
323
324
325
326
327
328
__host__ __device__ constexpr auto operator+=(Array<TData, NSize>& a, Array<TData, NSize> b)
{
    a = a + b;
    return a;
}

// Array -= Array
Chao Liu's avatar
Chao Liu committed
329
template <typename TData, index_t NSize>
Chao Liu's avatar
Chao Liu committed
330
331
332
333
334
__host__ __device__ constexpr auto operator-=(Array<TData, NSize>& a, Array<TData, NSize> b)
{
    a = a - b;
    return a;
}
335
// Array = Array + Sequence
Chao Liu's avatar
Chao Liu committed
336
template <typename TData, index_t NSize, index_t... Is>
337
338
339
340
341
342
__host__ __device__ constexpr auto operator+(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

343
    auto f = math::plus<index_t>{};
344

Chao Liu's avatar
Chao Liu committed
345
346
347
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
348
349
350
351
352

    return result;
}

// Array = Array - Sequence
Chao Liu's avatar
Chao Liu committed
353
template <typename TData, index_t NSize, index_t... Is>
354
355
356
357
358
359
__host__ __device__ constexpr auto operator-(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

360
    auto f = math::minus<index_t>{};
361

Chao Liu's avatar
Chao Liu committed
362
363
364
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
365
366
367
368

    return result;
}

Chao Liu's avatar
Chao Liu committed
369
// Array = Array * Sequence
Chao Liu's avatar
Chao Liu committed
370
template <typename TData, index_t NSize, index_t... Is>
Chao Liu's avatar
Chao Liu committed
371
372
373
374
375
376
__host__ __device__ constexpr auto operator*(Array<TData, NSize> a, Sequence<Is...> b)
{
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;

377
    auto f = math::multiplies<index_t>{};
Chao Liu's avatar
Chao Liu committed
378

Chao Liu's avatar
Chao Liu committed
379
380
381
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
Chao Liu's avatar
Chao Liu committed
382
383
384

    return result;
}
385

386
// Array = Sequence - Array
Chao Liu's avatar
Chao Liu committed
387
template <typename TData, index_t NSize, index_t... Is>
388
__host__ __device__ constexpr auto operator-(Sequence<Is...> a, Array<TData, NSize> b)
389
{
390
391
392
    static_assert(sizeof...(Is) == NSize, "wrong! size not the same");

    Array<TData, NSize> result;
393

394
    auto f = math::minus<index_t>{};
395

Chao Liu's avatar
Chao Liu committed
396
397
398
    static_for<0, NSize, 1>{}(
        lambda_array_math<decltype(f), decltype(a), decltype(b), decltype(result)>(
            f, a, b, result));
399
400
401
402

    return result;
}

Chao Liu's avatar
Chao Liu committed
403
template <typename TData, index_t NSize, typename Reduce>
404
405
406
407
408
409
410
__host__ __device__ constexpr TData
accumulate_on_array(const Array<TData, NSize>& a, Reduce f, TData init)
{
    TData result = init;

    static_assert(NSize > 0, "wrong");

Chao Liu's avatar
Chao Liu committed
411
    static_for<0, NSize, 1>{}([&](auto I) { result = f(result, a[I]); });
412
413
414

    return result;
}
415

416
417
} // namespace ck
#endif