functional.hip.hpp 1.04 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
#pragma once
#include "constant_integral.hip.hpp"

template <unsigned NLoop>
struct static_loop_n
{
    template <class F>
    __host__ __device__ void operator()(F f) const
    {
        static_assert(NLoop > 1, "out-of-range");

        f(Number<NLoop - 1>{});
        static_loop_n<NLoop - 1>{}(f);
    }
};

template <>
struct static_loop_n<1>
{
    template <class F>
    __host__ __device__ void operator()(F f) const
    {
        f(Number<0>{});
    }
};

template <unsigned NLoop>
struct static_const_reduce_n
{
    template <class F, class Reduce>
    __host__ __device__ constexpr auto operator()(F f, Reduce r) const
    {
        static_assert(NLoop > 1, "out-of-range");

        constexpr auto a = f(Number<NLoop - 1>{});
        auto b = static_const_reduce_n<NLoop - 1>{}(f, r); // cannot use constexpr here, weird
        return r(a, b);
    }
};

template <>
struct static_const_reduce_n<1>
{
    template <class F, class Reduce>
    __host__ __device__ constexpr auto operator()(F f, Reduce) const
    {
        return f(Number<0>{});
    }
};