"git@developer.sourcefind.cn:modelzoo/resnet50_tensorflow.git" did not exist on "3a4ac1eee9039eee0f0a3aad1e55c1394446b7ba"
ConstantMatrixDescriptor.hpp 2.06 KB
Newer Older
1
2
3
#ifndef CK_CONSTANT_MATRIX_DESCRIPTOR_HPP
#define CK_CONSTANT_MATRIX_DESCRIPTOR_HPP

Chao Liu's avatar
Chao Liu committed
4
#include "common.hpp"
Chao Liu's avatar
Chao Liu committed
5

6
7
namespace ck {

Chao Liu's avatar
Chao Liu committed
8
template <index_t NRow_, index_t NCol_, index_t RowStride_>
Chao Liu's avatar
Chao Liu committed
9
10
struct ConstantMatrixDescriptor
{
Chao Liu's avatar
Chao Liu committed
11
    __host__ __device__ constexpr ConstantMatrixDescriptor()
Chao Liu's avatar
Chao Liu committed
12
    {
Chao Liu's avatar
Chao Liu committed
13
        static_assert(NCol_ <= RowStride_, "wrong! NCol > RowStride!");
Chao Liu's avatar
Chao Liu committed
14
15
    }

Chao Liu's avatar
Chao Liu committed
16
    __host__ __device__ static constexpr index_t NRow() { return NRow_; }
Chao Liu's avatar
Chao Liu committed
17

Chao Liu's avatar
Chao Liu committed
18
    __host__ __device__ static constexpr index_t NCol() { return NCol_; }
Chao Liu's avatar
Chao Liu committed
19

Chao Liu's avatar
Chao Liu committed
20
    __host__ __device__ static constexpr index_t RowStride() { return RowStride_; }
Chao Liu's avatar
Chao Liu committed
21

Chao Liu's avatar
Chao Liu committed
22
    __host__ __device__ static constexpr auto GetLengths() { return Sequence<NRow_, NCol_>{}; }
Chao Liu's avatar
Chao Liu committed
23

Chao Liu's avatar
Chao Liu committed
24
    __host__ __device__ static constexpr index_t GetElementSize() { return NRow_ * NCol_; }
Chao Liu's avatar
Chao Liu committed
25

Chao Liu's avatar
Chao Liu committed
26
    __host__ __device__ static constexpr index_t GetElementSpace() { return NRow_ * RowStride_; }
Chao Liu's avatar
Chao Liu committed
27

Chao Liu's avatar
Chao Liu committed
28
    __host__ __device__ static index_t GetOffsetFromMultiIndex(index_t irow, index_t icol)
Chao Liu's avatar
Chao Liu committed
29
    {
Chao Liu's avatar
Chao Liu committed
30
        return irow * RowStride_ + icol;
Chao Liu's avatar
Chao Liu committed
31
32
    }

Chao Liu's avatar
Chao Liu committed
33
    template <index_t SubNRow, index_t SubNCol>
Chao Liu's avatar
Chao Liu committed
34
35
    __host__ __device__ static constexpr auto MakeSubMatrixDescriptor(Number<SubNRow>,
                                                                      Number<SubNCol>)
Chao Liu's avatar
Chao Liu committed
36
    {
Chao Liu's avatar
Chao Liu committed
37
        return ConstantMatrixDescriptor<SubNRow, SubNCol, RowStride_>{};
Chao Liu's avatar
Chao Liu committed
38
39
40
    }
};

Chao Liu's avatar
Chao Liu committed
41
template <index_t NRow, index_t NCol>
Chao Liu's avatar
Chao Liu committed
42
43
44
45
46
__host__ __device__ constexpr auto make_ConstantMatrixDescriptor(Number<NRow>, Number<NCol>)
{
    return ConstantMatrixDescriptor<NRow, NCol, NCol>{};
}

Chao Liu's avatar
Chao Liu committed
47
template <index_t NRow, index_t NCol, index_t RowStride>
Chao Liu's avatar
Chao Liu committed
48
49
50
51
52
__host__ __device__ constexpr auto
    make_ConstantMatrixDescriptor(Number<NRow>, Number<NCol>, Number<RowStride>)
{
    return ConstantMatrixDescriptor<NRow, NCol, RowStride>{};
}
Chao Liu's avatar
Chao Liu committed
53
54
55
56
57
58
59
60
61
62
63

template <class TDesc>
__host__ __device__ void print_ConstantMatrixDescriptor(TDesc, const char* s)
{
    const auto desc = TDesc{};

    constexpr auto I0 = Number<0>{};
    constexpr auto I1 = Number<1>{};

    printf("%s NRow %u NCol %u RowStride %u\n", s, desc.NRow(), desc.NCol(), desc.RowStride());
}
64
65
66
67

} // namespace ck

#endif