gemm.cpp 2.64 KB
Newer Older
wsttiger's avatar
wsttiger committed
1
2
3
4
5
6
7
#include <migraph/gpu/gemm.hpp>
#include <migraph/operators.hpp>
#include <migraph/manage_ptr.hpp>
#include <migraph/gpu/miopen.hpp>
#include <utility>

namespace migraph {
8
inline namespace MIGRAPH_INLINE_NS {
wsttiger's avatar
wsttiger committed
9
10
namespace gpu {

Paul's avatar
Paul committed
11
template <class... Ts>
Paul's avatar
Paul committed
12
13
void generic_rocblas_gemm(shape::as<float>, Ts&&... xs)
{
Paul's avatar
Paul committed
14
    rocblas_sgemm(std::forward<Ts>(xs)...);
Paul's avatar
Paul committed
15
16
}

Paul's avatar
Paul committed
17
template <class... Ts>
Paul's avatar
Paul committed
18
19
void generic_rocblas_gemm(shape::as<double>, Ts&&... xs)
{
Paul's avatar
Paul committed
20
    rocblas_dgemm(std::forward<Ts>(xs)...);
Paul's avatar
Paul committed
21
22
}

Paul's avatar
Paul committed
23
template <class... Ts>
Paul's avatar
Paul committed
24
25
void generic_rocblas_gemm(shape::as<half>, Ts&&... xs)
{
Paul's avatar
Paul committed
26
    rocblas_hgemm(std::forward<Ts>(xs)...);
Paul's avatar
Paul committed
27
28
}

Paul's avatar
Paul committed
29
template <class T, class... Ts>
Paul's avatar
Paul committed
30
31
void generic_rocblas_gemm(shape::as<T>, Ts&&...)
{
Paul's avatar
Paul committed
32
    MIGRAPH_THROW("Type unsupported by rocblas");
Paul's avatar
Paul committed
33
34
}

Paul's avatar
Paul committed
35
template <class T>
Paul's avatar
Paul committed
36
37
T to_rocblas_type(T x)
{
Paul's avatar
Paul committed
38
    return x;
Paul's avatar
Paul committed
39
40
}

Paul's avatar
Paul committed
41
rocblas_half to_rocblas_type(half x) { return reinterpret_cast<const rocblas_half&>(x); }
Paul's avatar
Paul committed
42

wsttiger's avatar
wsttiger committed
43
44
45
46
47
shape miopen_gemm::compute_shape(const std::vector<shape>& inputs) const
{
    check_shapes{inputs, *this}.has(3);
    return op.compute_shape({inputs.at(0), inputs.at(1)});
}
wsttiger's avatar
wsttiger committed
48
49
50
argument miopen_gemm::compute(context& ctx,
                              const shape& output_shape,
                              const std::vector<argument>& args) const
wsttiger's avatar
wsttiger committed
51
52
53
54
55
56
57
58
59
60
61
{
    float alpha     = 1.0f;
    float beta      = 0.0f;
    bool transa     = args[0].get_shape().transposed();
    bool transb     = args[1].get_shape().transposed();
    rocblas_int lda = args[0].get_shape().strides()[transa ? 1 : 0];
    rocblas_int ldb = args[1].get_shape().strides()[transb ? 1 : 0];
    rocblas_int ldc = args[2].get_shape().strides()[0];
    rocblas_int m   = output_shape.lens()[0];
    rocblas_int n   = output_shape.lens()[1];
    rocblas_int k   = args[0].get_shape().lens()[1];
Paul's avatar
Paul committed
62
    output_shape.visit_type([&](auto as) {
Paul's avatar
Paul committed
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
        auto alpha_r = to_rocblas_type(as(alpha));
        auto beta_r  = to_rocblas_type(as(beta));
        generic_rocblas_gemm(as,
                             ctx.get_stream().get_rocblas(),
                             transb ? rocblas_operation_transpose : rocblas_operation_none,
                             transa ? rocblas_operation_transpose : rocblas_operation_none,
                             n,
                             m,
                             k,
                             &alpha_r,
                             args[1].implicit(),
                             ldb,
                             args[0].implicit(),
                             lda,
                             &beta_r,
                             args[2].implicit(),
                             ldc);

Paul's avatar
Paul committed
81
    });
wsttiger's avatar
wsttiger committed
82
83
84
85
    return args[2];
}

} // namespace gpu
86
} // namespace MIGRAPH_INLINE_NS
wsttiger's avatar
wsttiger committed
87
} // namespace migraph