functional2.hpp 1.27 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.

Chao Liu's avatar
Chao Liu committed
4
#pragma once
5

Chao Liu's avatar
Chao Liu committed
6
7
#include "ck/utility/functional.hpp"
#include "ck/utility/sequence.hpp"
8

9
10
namespace ck {

Chao Liu's avatar
Chao Liu committed
11
12
namespace detail {

Chao Liu's avatar
Chao Liu committed
13
14
template <class>
struct static_for_impl;
15

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

Chao Liu's avatar
Chao Liu committed
26
27
} // namespace detail

Chao Liu's avatar
Chao Liu committed
28
29
30
// F signature: F(Number<Iter>)
template <index_t NBegin, index_t NEnd, index_t Increment>
struct static_for
31
{
32
    __host__ __device__ constexpr static_for()
33
    {
34
        static_assert(Increment != 0 && (NEnd - NBegin) % Increment == 0,
Chao Liu's avatar
Chao Liu committed
35
                      "Wrong! should satisfy (NEnd - NBegin) % Increment == 0");
36
        static_assert((Increment > 0 && NBegin <= NEnd) || (Increment < 0 && NBegin >= NEnd),
Chao Liu's avatar
Chao Liu committed
37
38
                      "wrongs! should (Increment > 0 && NBegin <= NEnd) || (Increment < 0 && "
                      "NBegin >= NEnd)");
39
    }
40

41
42
43
    template <class F>
    __host__ __device__ constexpr void operator()(F f) const
    {
Chao Liu's avatar
Chao Liu committed
44
45
        detail::static_for_impl<typename arithmetic_sequence_gen<NBegin, NEnd, Increment>::type>{}(
            f);
46
47
48
    }
};

49
} // namespace ck