threadwise_gemm.hip.hpp 2.39 KB
Newer Older
1
2
#pragma once

Chao Liu's avatar
Chao Liu committed
3
template <class Float, class SrcMatrix, class DstMatrix, index_t NRow, index_t NCol>
Chao Liu's avatar
Chao Liu committed
4
5
6
7
8
__device__ void threadwise_matrix_copy(SrcMatrix,
                                       const Float* __restrict__ p_src,
                                       DstMatrix,
                                       Float* __restrict__ p_dst,
                                       Sequence<NRow, NCol>)
9
{
Chao Liu's avatar
Chao Liu committed
10
11
    constexpr auto src_mtx = SrcMatrix{};
    constexpr auto dst_mtx = DstMatrix{};
12

Chao Liu's avatar
Chao Liu committed
13
    for(index_t i = 0; i < NRow; ++i)
14
    {
Chao Liu's avatar
Chao Liu committed
15
        for(index_t j = 0; j < NCol; ++j)
16
        {
Chao Liu's avatar
Chao Liu committed
17
18
            const index_t src_index = src_mtx.Get1dIndex(i, j);
            const index_t dst_index = dst_mtx.Get1dIndex(i, j);
19

Chao Liu's avatar
Chao Liu committed
20
            p_dst[dst_index] = p_src[src_index];
21
22
23
24
25
26
27
28
29
30
31
32
        }
    }
}

template <class MatrixA,
          class MatrixB,
          class MatrixC,
          bool TransA,
          bool TransB,
          bool TransC,
          class FloatA,
          class FloatB,
33
          class FloatC>
34
__device__ void threadwise_gemm(MatrixA,
Chao Liu's avatar
Chao Liu committed
35
                                integral_constant<bool, TransA>,
Chao Liu's avatar
Chao Liu committed
36
                                const FloatA* __restrict__ p_a_thread,
37
                                MatrixB,
Chao Liu's avatar
Chao Liu committed
38
                                integral_constant<bool, TransB>,
Chao Liu's avatar
Chao Liu committed
39
                                const FloatB* __restrict__ p_b_thread,
40
                                MatrixC,
Chao Liu's avatar
Chao Liu committed
41
                                integral_constant<bool, TransC>,
42
                                FloatC* __restrict__ p_c_thread)
43
44
45
{
    if(TransA && (!TransB) && (!TransC))
    {
Chao Liu's avatar
Chao Liu committed
46
47
48
        constexpr auto a_mtx = MatrixA{};
        constexpr auto b_mtx = MatrixB{};
        constexpr auto c_mtx = MatrixC{};
49

Chao Liu's avatar
Chao Liu committed
50
51
52
        constexpr index_t M = c_mtx.NRow();
        constexpr index_t N = c_mtx.NCol();
        constexpr index_t K = a_mtx.NRow(); // A is transposed
53

Chao Liu's avatar
Chao Liu committed
54
        for(index_t k = 0; k < K; ++k)
55
        {
Jing Zhang's avatar
Jing Zhang committed
56
            for(index_t i = 0; i < M; i++)
57
            {
Jing Zhang's avatar
Jing Zhang committed
58
                for(index_t j = 0; j < N; j++)
59
                {
Chao Liu's avatar
Chao Liu committed
60
61
62
                    const index_t aindex = a_mtx.Get1dIndex(k, i); // A is transposed
                    const index_t bindex = b_mtx.Get1dIndex(k, j);
                    const index_t cindex = c_mtx.Get1dIndex(i, j);
63

Jing Zhang's avatar
Jing Zhang committed
64
                    p_c_thread[cindex] += p_a_thread[aindex] * p_b_thread[bindex];
65
66
67
68
69
70
71
72
73
74
                }
            }
        }
    }
    else
    {
        // not implemented
        assert(false);
    }
}