"docs/vscode:/vscode.git/clone" did not exist on "12dc3e144f8cf306ba5338d3d939d86bc98f99d6"
Sequence.hip.hpp 2.54 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
#pragma once
#include "constant_integral.hip.hpp"
#include "functional.hip.hpp"

template <unsigned... Is>
struct Sequence
{
    using Type = Sequence<Is...>;

    static constexpr unsigned nDim = sizeof...(Is);

    const unsigned mData[nDim] = {Is...};

    template <unsigned I>
    __host__ __device__ constexpr unsigned Get(Number<I>) const
    {
        return mData[I];
    }

    // this is ugly, only for nDIm = 4
    template <unsigned I0, unsigned I1, unsigned I2, unsigned I3>
    __host__ __device__ constexpr auto ReorderByGetNewFromOld(Sequence<I0, I1, I2, I3>) const
    {
        static_assert(nDim == 4, "nDim != 4");

        constexpr auto old_sequence = Type{};

        constexpr unsigned NR0 = old_sequence.mData[I0];
        constexpr unsigned NR1 = old_sequence.mData[I1];
        constexpr unsigned NR2 = old_sequence.mData[I2];
        constexpr unsigned NR3 = old_sequence.mData[I3];

        return Sequence<NR0, NR1, NR2, NR3>{};
    }

    template <unsigned I0, unsigned I1, unsigned I2, unsigned I3>
    __host__ __device__ constexpr auto ReorderByPutOldToNew(Sequence<I0, I1, I2, I3>) const
    {
        // don't know how to implement this
        printf("Sequence::ReorderByPutOldToNew not implemented");
        assert(false);
    }

    template <unsigned I>
    __host__ __device__ constexpr auto PushBack(Number<I>) const
    {
        return Sequence<Is..., I>{};
    }

    __host__ __device__ constexpr auto PopBack() const;

    template <class F>
    __host__ __device__ constexpr auto Transform(F f) const
    {
        return Sequence<f(Is)...>{};
    }
};

template <unsigned... Is, unsigned I>
__host__ __device__ constexpr auto sequence_pop_back(Sequence<Is..., I>)
{
    static_assert(sizeof...(Is) >= 1, "empty Sequence!");
    return Sequence<Is...>{};
}

template <class F, unsigned... Xs, unsigned... Ys>
__host__ __device__ constexpr auto sequence_sequence_op(Sequence<Xs...>, Sequence<Ys...>, F f)
{
    static_assert(Sequence<Xs...>::nDim == Sequence<Ys...>::nDim, "Dim not the same");

    return Sequence<f(Xs, Ys)...>{};
}

template <unsigned... Xs, unsigned... Ys>
__host__ __device__ constexpr auto sequence_sequence_add(Sequence<Xs...>, Sequence<Ys...>)
{
    struct add
    {
        __host__ __device__ constexpr unsigned operator()(unsigned x, unsigned y) const
        {
            return x + y;
        }
    };

    return sequence_sequence_op(Sequence<Xs...>{}, Sequence<Ys...>{}, add{});
}

template <unsigned... Is>
__host__ __device__ constexpr auto Sequence<Is...>::PopBack() const
{
    return sequence_pop_back(Type{});
}