functional2.hpp 1.66 KB
Newer Older
1
2
3
#ifndef CK_FUNCTIONAL2_HPP
#define CK_FUNCTIONAL2_HPP

Chao Liu's avatar
Chao Liu committed
4
5
#include "composable_kernel/utility/functional.hpp"
#include "composable_kernel/utility/Sequence.hpp"
6

7
8
namespace ck {

Chao Liu's avatar
Chao Liu committed
9
10
template <class>
struct static_for_impl;
11

Chao Liu's avatar
Chao Liu committed
12
13
template <index_t... Is>
struct static_for_impl<Sequence<Is...>>
14
15
{
    template <class F>
Chao Liu's avatar
Chao Liu committed
16
    __host__ __device__ constexpr void operator()(F f) const
17
    {
Chao Liu's avatar
Chao Liu committed
18
        swallow{(f(Number<Is>{}), 0)...};
19
20
21
    }
};

Chao Liu's avatar
Chao Liu committed
22
23
24
// F signature: F(Number<Iter>)
template <index_t NBegin, index_t NEnd, index_t Increment>
struct static_for
25
{
Chao Liu's avatar
Chao Liu committed
26
    template <class F>
Chao Liu's avatar
Chao Liu committed
27
    __host__ __device__ constexpr void operator()(F f) const
28
    {
Chao Liu's avatar
Chao Liu committed
29
        static_assert(NBegin <= NEnd, "wrongs! should have NBegin <= NEnd");
30

Chao Liu's avatar
Chao Liu committed
31
32
        static_assert((NEnd - NBegin) % Increment == 0,
                      "Wrong! should satisfy (NEnd - NBegin) % Increment == 0");
33

Chao Liu's avatar
Chao Liu committed
34
        static_for_impl<typename arithmetic_sequence_gen<NBegin, NEnd, Increment>::SeqType>{}(f);
35
36
37
    }
};

Chao Liu's avatar
Chao Liu committed
38
39
template <class Seq, class Reduce>
struct lambda_accumulate_on_sequence
40
{
Chao Liu's avatar
Chao Liu committed
41
42
43
44
45
    const Reduce& f;
    index_t& result;

    __host__ __device__ constexpr lambda_accumulate_on_sequence(const Reduce& f_, index_t& result_)
        : f(f_), result(result_)
46
    {
Chao Liu's avatar
Chao Liu committed
47
48
49
50
51
52
    }

    template <class IDim>
    __host__ __device__ constexpr index_t operator()(IDim) const
    {
        return result = f(result, Seq::Get(IDim{}));
53
54
55
    }
};

Chao Liu's avatar
Chao Liu committed
56
57
58
template <class Seq, class Reduce, index_t Init>
__host__ __device__ constexpr index_t
accumulate_on_sequence(Seq, Reduce f, Number<Init> /*initial_value*/)
59
{
Chao Liu's avatar
Chao Liu committed
60
    index_t result = Init;
Chao Liu's avatar
Chao Liu committed
61

Chao Liu's avatar
Chao Liu committed
62
    static_for<0, Seq::mSize, 1>{}(lambda_accumulate_on_sequence<Seq, Reduce>(f, result));
63

Chao Liu's avatar
Chao Liu committed
64
65
    return result;
}
66
67
68

} // namespace ck
#endif