"test/vscode:/vscode.git/clone" did not exist on "542a660d052a6f639cdeabf928832f31a1db10a5"
mul.cpp 1.58 KB
Newer Older
Paul's avatar
Paul committed
1
2
#include <migraphx/gpu/device/mul.hpp>
#include <migraphx/gpu/device/nary.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
3
4
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
Khalique's avatar
Khalique committed
5

Paul's avatar
Paul committed
6
namespace migraphx {
Paul's avatar
Paul committed
7
inline namespace MIGRAPHX_INLINE_NS {
Khalique's avatar
Khalique committed
8
9
10
namespace gpu {
namespace device {

Shucai Xiao's avatar
Shucai Xiao committed
11
12
13
__global__ void mul_kernel(__half* a, __half* b, __half* r, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
Shucai Xiao's avatar
Shucai Xiao committed
14
    if(tid < n)
Shucai Xiao's avatar
Shucai Xiao committed
15
    {
Shucai Xiao's avatar
Shucai Xiao committed
16
        r[tid] = a[tid] * b[tid % 768];
Shucai Xiao's avatar
Shucai Xiao committed
17
18
19
    }
}

Khalique's avatar
Khalique committed
20
void mul(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
Khalique's avatar
Khalique committed
21
{
Shucai Xiao's avatar
Shucai Xiao committed
22
    auto s2 = arg2.get_shape();
Shucai Xiao's avatar
Shucai Xiao committed
23
    if(s2.element_space() == 768 and s2.type() == shape::half_type)
Shucai Xiao's avatar
Shucai Xiao committed
24
    {
Shucai Xiao's avatar
Shucai Xiao committed
25
        auto elem_num  = s2.elements();
Shucai Xiao's avatar
Shucai Xiao committed
26
        int block_size = 1024;
Shucai Xiao's avatar
Shucai Xiao committed
27
        int block_num  = (elem_num + block_size - 1) / block_size;
Shucai Xiao's avatar
Shucai Xiao committed
28
        mul_kernel<<<block_num, block_size>>>(reinterpret_cast<__half*>(arg1.data()),
Shucai Xiao's avatar
Shucai Xiao committed
29
30
31
                                              reinterpret_cast<__half*>(arg2.data()),
                                              reinterpret_cast<__half*>(result.data()),
                                              elem_num);
Shucai Xiao's avatar
Shucai Xiao committed
32
33
34
    }
    else
    {
Shucai Xiao's avatar
Shucai Xiao committed
35
        nary(stream, result, arg1, arg2)([](auto x, auto y) __device__ { return x * y; });
Shucai Xiao's avatar
Shucai Xiao committed
36
    }
Khalique's avatar
Khalique committed
37
38
39
40
41
42
43
44
}

void mul(hipStream_t stream,
         const argument& result,
         const argument& arg1,
         const argument& arg2,
         const argument& arg3)
{
45
46
    nary(stream, result, arg1, arg2, arg3)([](auto x, auto y, auto z)
                                               __device__ { return x * y * z; });
Khalique's avatar
Khalique committed
47
48
49
50
}

} // namespace device
} // namespace gpu
Paul's avatar
Paul committed
51
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
52
} // namespace migraphx