functional2.hip.hpp 2.15 KB
Newer Older
1
#pragma once
Chao Liu's avatar
Chao Liu committed
2
#include "functional.hip.hpp"
3
4
#include "Sequence.hip.hpp"

Chao Liu's avatar
Chao Liu committed
5
6
7
#if 0
template <index_t Iter, index_t Remaining, index_t Increment>
struct static_for_impl
8
{
Chao Liu's avatar
Chao Liu committed
9
10
    template <class F>
    constexpr __host__ __device__ void operator()(F f) const
11
    {
Chao Liu's avatar
Chao Liu committed
12
13
        static_assert(Remaining % Increment == 0, "wrong! Remaining % Increment != 0");
        static_assert(Increment <= Remaining, "will go out-of-range");
14

Chao Liu's avatar
Chao Liu committed
15
16
        f(Number<Iter>{});
        static_for_impl<Iter + Increment, Remaining - Increment, Increment>{}(f);
17
18
19
    }
};

Chao Liu's avatar
Chao Liu committed
20
21
template <index_t Iter, index_t Increment>
struct static_for_impl<Iter, 0, Increment>
22
23
{
    template <class F>
Chao Liu's avatar
Chao Liu committed
24
    constexpr __host__ __device__ void operator()(F) const
25
    {
Chao Liu's avatar
Chao Liu committed
26
27
        // no work left, just return
        return;
28
29
30
    }
};

Chao Liu's avatar
Chao Liu committed
31
32
33
// F signature: F(Number<Iter>)
template <index_t NBegin, index_t NEnd, index_t Increment>
struct static_for
34
{
Chao Liu's avatar
Chao Liu committed
35
36
    template <class F>
    constexpr __host__ __device__ void operator()(F f) const
37
    {
Chao Liu's avatar
Chao Liu committed
38
        static_assert(NBegin <= NEnd, "wrongs! should have NBegin <= NEnd");
39

Chao Liu's avatar
Chao Liu committed
40
41
        static_assert((NEnd - NBegin) % Increment == 0,
                      "Wrong! should satisfy (NEnd - NBegin) % Increment == 0");
42

Chao Liu's avatar
Chao Liu committed
43
44
45
46
47
48
#if 0
        static_if<(NBegin < NEnd)>{}(
            [&](auto fwd) { static_for_impl<NBegin, NEnd - NBegin, fwd(Increment)>{}(f); });
#else
        static_for_impl<NBegin, NEnd - NBegin, Increment>{}(f);
#endif
49
50
    }
};
Chao Liu's avatar
Chao Liu committed
51
52
53
#else
template <class>
struct static_for_impl;
54

Chao Liu's avatar
Chao Liu committed
55
56
template <index_t... Is>
struct static_for_impl<Sequence<Is...>>
57
{
Chao Liu's avatar
Chao Liu committed
58
59
    template <class F>
    __host__ __device__ constexpr void operator()(F f) const
60
    {
Chao Liu's avatar
Chao Liu committed
61
        swallow{(f(Number<Is>{}), 0)...};
62
63
64
    }
};

Chao Liu's avatar
Chao Liu committed
65
66
67
// F signature: F(Number<Iter>)
template <index_t NBegin, index_t NEnd, index_t Increment>
struct static_for
68
69
{
    template <class F>
Chao Liu's avatar
Chao Liu committed
70
    __host__ __device__ constexpr void operator()(F f) const
71
    {
Chao Liu's avatar
Chao Liu committed
72
73
74
75
        static_assert(NBegin <= NEnd, "wrongs! should have NBegin <= NEnd");

        static_assert((NEnd - NBegin) % Increment == 0,
                      "Wrong! should satisfy (NEnd - NBegin) % Increment == 0");
76

Chao Liu's avatar
Chao Liu committed
77
        static_for_impl<typename arithmetic_sequence_gen<NBegin, NEnd, Increment>::SeqType>{}(f);
78
79
    }
};
Chao Liu's avatar
Chao Liu committed
80
#endif