"vscode:/vscode.git/clone" did not exist on "8cf196ccb54a13d8cf5cada640599fcdc827e535"
math.hpp 4.51 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
Chao Liu's avatar
Chao Liu committed
5

Chao Liu's avatar
Chao Liu committed
6
#include "ck/ck.hpp"
Chao Liu's avatar
Chao Liu committed
7
#include "integral_constant.hpp"
Chao Liu's avatar
Chao Liu committed
8
#include "number.hpp"
Chao Liu's avatar
Chao Liu committed
9
#include "type.hpp"
Chao Liu's avatar
Chao Liu committed
10
#include "enable_if.hpp"
Chao Liu's avatar
Chao Liu committed
11
12
13
14

namespace ck {
namespace math {

zjing14's avatar
zjing14 committed
15
template <typename T, T s>
Chao Liu's avatar
Chao Liu committed
16
17
18
19
20
struct scales
{
    __host__ __device__ constexpr T operator()(T a) const { return s * a; }
};

zjing14's avatar
zjing14 committed
21
template <typename T>
Chao Liu's avatar
Chao Liu committed
22
23
24
25
26
struct plus
{
    __host__ __device__ constexpr T operator()(T a, T b) const { return a + b; }
};

zjing14's avatar
zjing14 committed
27
template <typename T>
Chao Liu's avatar
Chao Liu committed
28
29
30
31
32
33
struct minus
{
    __host__ __device__ constexpr T operator()(T a, T b) const { return a - b; }
};

struct multiplies
Chao Liu's avatar
Chao Liu committed
34
35
36
37
38
39
40
41
{
    template <typename A, typename B>
    __host__ __device__ constexpr auto operator()(const A& a, const B& b) const
    {
        return a * b;
    }
};

zjing14's avatar
zjing14 committed
42
template <typename T>
Chao Liu's avatar
Chao Liu committed
43
struct maximize
Chao Liu's avatar
Chao Liu committed
44
45
46
47
{
    __host__ __device__ constexpr T operator()(T a, T b) const { return a >= b ? a : b; }
};

zjing14's avatar
zjing14 committed
48
template <typename T>
Chao Liu's avatar
Chao Liu committed
49
50
51
52
53
struct minimize
{
    __host__ __device__ constexpr T operator()(T a, T b) const { return a <= b ? a : b; }
};

zjing14's avatar
zjing14 committed
54
template <typename T>
Chao Liu's avatar
Chao Liu committed
55
56
57
58
59
60
struct integer_divide_ceiler
{
    __host__ __device__ constexpr T operator()(T a, T b) const
    {
        static_assert(is_same<T, index_t>{} || is_same<T, int>{}, "wrong type");

zjing14's avatar
zjing14 committed
61
        return (a + b - Number<1>{}) / b;
Chao Liu's avatar
Chao Liu committed
62
63
64
    }
};

zjing14's avatar
zjing14 committed
65
template <typename X, typename Y>
66
67
68
69
70
__host__ __device__ constexpr auto integer_divide_floor(X x, Y y)
{
    return x / y;
}

zjing14's avatar
zjing14 committed
71
template <typename X, typename Y>
Chao Liu's avatar
Chao Liu committed
72
__host__ __device__ constexpr auto integer_divide_ceil(X x, Y y)
Chao Liu's avatar
Chao Liu committed
73
{
74
    return (x + y - Number<1>{}) / y;
Chao Liu's avatar
Chao Liu committed
75
76
}

zjing14's avatar
zjing14 committed
77
template <typename X, typename Y>
Chao Liu's avatar
Chao Liu committed
78
__host__ __device__ constexpr auto integer_least_multiple(X x, Y y)
Chao Liu's avatar
Chao Liu committed
79
{
Chao Liu's avatar
Chao Liu committed
80
    return y * integer_divide_ceil(x, y);
Chao Liu's avatar
Chao Liu committed
81
82
}

zjing14's avatar
zjing14 committed
83
template <typename T>
Chao Liu's avatar
Chao Liu committed
84
85
86
87
88
__host__ __device__ constexpr T max(T x)
{
    return x;
}

zjing14's avatar
zjing14 committed
89
90
template <typename T>
__host__ __device__ constexpr T max(T x, T y)
Chao Liu's avatar
Chao Liu committed
91
{
zjing14's avatar
zjing14 committed
92
93
    return x > y ? x : y;
}
Chao Liu's avatar
Chao Liu committed
94

zjing14's avatar
zjing14 committed
95
96
97
98
99
template <index_t X>
__host__ __device__ constexpr index_t max(Number<X>, index_t y)
{
    return X > y ? X : y;
}
Chao Liu's avatar
Chao Liu committed
100

zjing14's avatar
zjing14 committed
101
102
103
104
105
template <index_t Y>
__host__ __device__ constexpr index_t max(index_t x, Number<Y>)
{
    return x > Y ? x : Y;
}
Chao Liu's avatar
Chao Liu committed
106

zjing14's avatar
zjing14 committed
107
108
109
110
111
112
template <typename X, typename... Ys>
__host__ __device__ constexpr auto max(X x, Ys... ys)
{
    static_assert(sizeof...(Ys) > 0, "not enough argument");

    return max(x, max(ys...));
Chao Liu's avatar
Chao Liu committed
113
114
}

zjing14's avatar
zjing14 committed
115
template <typename T>
Chao Liu's avatar
Chao Liu committed
116
117
118
119
120
__host__ __device__ constexpr T min(T x)
{
    return x;
}

zjing14's avatar
zjing14 committed
121
122
123
124
125
126
127
128
template <typename T>
__host__ __device__ constexpr T min(T x, T y)
{
    return x < y ? x : y;
}

template <index_t X>
__host__ __device__ constexpr index_t min(Number<X>, index_t y)
Chao Liu's avatar
Chao Liu committed
129
{
zjing14's avatar
zjing14 committed
130
131
    return X < y ? X : y;
}
Chao Liu's avatar
Chao Liu committed
132

zjing14's avatar
zjing14 committed
133
134
135
136
137
template <index_t Y>
__host__ __device__ constexpr index_t min(index_t x, Number<Y>)
{
    return x < Y ? x : Y;
}
Chao Liu's avatar
Chao Liu committed
138

zjing14's avatar
zjing14 committed
139
140
141
142
template <typename X, typename... Ys>
__host__ __device__ constexpr auto min(X x, Ys... ys)
{
    static_assert(sizeof...(Ys) > 0, "not enough argument");
Chao Liu's avatar
Chao Liu committed
143

zjing14's avatar
zjing14 committed
144
    return min(x, min(ys...));
Chao Liu's avatar
Chao Liu committed
145
146
}

147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
// disallow implicit type casting
template <typename T>
__device__ T exp(T x);

template <>
__device__ float exp<float>(float x)
{
    return __expf(x);
}

template <>
__device__ double exp<double>(double x)
{
    return exp(x);
}

163
// greatest common divisor, aka highest common factor
Chao Liu's avatar
Chao Liu committed
164
__host__ __device__ constexpr index_t gcd(index_t x, index_t y)
Chao Liu's avatar
Chao Liu committed
165
{
166
167
168
169
170
171
172
173
174
    if(x < 0)
    {
        return gcd(-x, y);
    }
    else if(y < 0)
    {
        return gcd(x, -y);
    }
    else if(x == y || x == 0)
Chao Liu's avatar
Chao Liu committed
175
176
177
    {
        return y;
    }
Chao Liu's avatar
Chao Liu committed
178
    else if(y == 0)
Chao Liu's avatar
Chao Liu committed
179
180
181
    {
        return x;
    }
Chao Liu's avatar
Chao Liu committed
182
    else if(x > y)
Chao Liu's avatar
Chao Liu committed
183
    {
184
        return gcd(x % y, y);
Chao Liu's avatar
Chao Liu committed
185
    }
Chao Liu's avatar
Chao Liu committed
186
    else
Chao Liu's avatar
Chao Liu committed
187
    {
188
        return gcd(x, y % x);
Chao Liu's avatar
Chao Liu committed
189
190
191
192
    }
}

template <index_t X, index_t Y>
193
__host__ __device__ constexpr auto gcd(Number<X>, Number<Y>)
Chao Liu's avatar
Chao Liu committed
194
{
Chao Liu's avatar
Chao Liu committed
195
196
197
    constexpr auto r = gcd(X, Y);

    return Number<r>{};
Chao Liu's avatar
Chao Liu committed
198
199
}

Chao Liu's avatar
Chao Liu committed
200
template <typename X, typename... Ys, typename enable_if<sizeof...(Ys) >= 2, bool>::type = false>
201
__host__ __device__ constexpr auto gcd(X x, Ys... ys)
Chao Liu's avatar
Chao Liu committed
202
{
203
    return gcd(x, gcd(ys...));
Chao Liu's avatar
Chao Liu committed
204
205
206
}

// least common multiple
Chao Liu's avatar
Chao Liu committed
207
208
template <typename X, typename Y>
__host__ __device__ constexpr auto lcm(X x, Y y)
Chao Liu's avatar
Chao Liu committed
209
{
210
    return (x * y) / gcd(x, y);
Chao Liu's avatar
Chao Liu committed
211
212
}

Chao Liu's avatar
Chao Liu committed
213
template <typename X, typename... Ys, typename enable_if<sizeof...(Ys) >= 2, bool>::type = false>
Chao Liu's avatar
Chao Liu committed
214
__host__ __device__ constexpr auto lcm(X x, Ys... ys)
Chao Liu's avatar
Chao Liu committed
215
{
Chao Liu's avatar
Chao Liu committed
216
    return lcm(x, lcm(ys...));
Chao Liu's avatar
Chao Liu committed
217
218
}

zjing14's avatar
zjing14 committed
219
template <typename T>
Chao Liu's avatar
Chao Liu committed
220
221
222
223
224
struct equal
{
    __host__ __device__ constexpr bool operator()(T x, T y) const { return x == y; }
};

zjing14's avatar
zjing14 committed
225
template <typename T>
Chao Liu's avatar
Chao Liu committed
226
227
228
229
230
struct less
{
    __host__ __device__ constexpr bool operator()(T x, T y) const { return x < y; }
};

Chao Liu's avatar
Chao Liu committed
231
} // namespace math
Chao Liu's avatar
Chao Liu committed
232
} // namespace ck