common.hip.hpp 1.9 KB
Newer Older
1
#pragma once
Chao Liu's avatar
Chao Liu committed
2
#include "vector_type.hip.hpp"
3
4
5
6
#include "constant_integral.hip.hpp"
#include "Sequence.hip.hpp"
#include "Array.hip.hpp"
#include "functional.hip.hpp"
7
#include "functional2.hip.hpp"
8

Chao Liu's avatar
Chao Liu committed
9
#if USE_AMD_INLINE_ASM
Chao Liu's avatar
Chao Liu committed
10
#include "amd_inline_asm.hip.hpp"
Chao Liu's avatar
Chao Liu committed
11
12
#endif

Chao Liu's avatar
Chao Liu committed
13
__device__ index_t get_thread_local_1d_id() { return threadIdx.x; }
Chao Liu's avatar
Chao Liu committed
14

Chao Liu's avatar
Chao Liu committed
15
__device__ index_t get_block_1d_id() { return blockIdx.x; }
16

17
18
19
20
21
22
23
24
25
26
27
template <class T1, class T2>
struct is_same
{
    static const bool value = false;
};

template <class T>
struct is_same<T, T>
{
    static const bool value = true;
};
Chao Liu's avatar
Chao Liu committed
28

29
namespace mod_conv { // namespace mod_conv
Chao Liu's avatar
Chao Liu committed
30
31
template <class T, T s>
struct scales
32
{
Chao Liu's avatar
Chao Liu committed
33
    __host__ __device__ constexpr T operator()(T a) const { return s * a; }
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
};

template <class T>
struct integer_divide_ceiler
{
    __host__ __device__ constexpr T operator()(T a, T b) const
    {
        static_assert(is_same<T, index_t>::value || is_same<T, int>::value, "wrong type");

        return (a + b - 1) / b;
    }
};

template <class T>
__host__ __device__ constexpr T integer_divide_ceil(T a, T b)
{
    static_assert(is_same<T, index_t>::value || is_same<T, int>::value, "wrong type");

Chao Liu's avatar
Chao Liu committed
52
    return (a + b - 1) / b;
Chao Liu's avatar
Chao Liu committed
53
54
}

Chao Liu's avatar
Chao Liu committed
55
56
template <class T>
__host__ __device__ constexpr T max(T x, T y)
Chao Liu's avatar
Chao Liu committed
57
{
Chao Liu's avatar
Chao Liu committed
58
    return x > y ? x : y;
Chao Liu's avatar
Chao Liu committed
59
60
}

Chao Liu's avatar
Chao Liu committed
61
62
template <class T, class... Ts>
__host__ __device__ constexpr T max(T x, Ts... xs)
Chao Liu's avatar
Chao Liu committed
63
{
Chao Liu's avatar
Chao Liu committed
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
    static_assert(sizeof...(xs) > 0, "not enough argument");

    auto y = max(xs...);

    static_assert(is_same<decltype(y), T>::value, "not the same type");

    return x > y ? x : y;
}

template <class T>
__host__ __device__ constexpr T min(T x, T y)
{
    return x < y ? x : y;
}

template <class T, class... Ts>
__host__ __device__ constexpr T min(T x, Ts... xs)
{
    static_assert(sizeof...(xs) > 0, "not enough argument");

    auto y = min(xs...);

    static_assert(is_same<decltype(y), T>::value, "not the same type");

    return x < y ? x : y;
}
Chao Liu's avatar
tidy up  
Chao Liu committed
90
} // namespace mod_conv