Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
60b12489
Commit
60b12489
authored
Oct 17, 2022
by
wangshaojie6
Browse files
add dynamic kernel example
parent
16467e0e
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
60 additions
and
1 deletion
+60
-1
example/01_gemm/gemm_xdl_fp16_splitk.cpp
example/01_gemm/gemm_xdl_fp16_splitk.cpp
+59
-0
include/ck/utility/common_header.hpp
include/ck/utility/common_header.hpp
+1
-1
No files found.
example/01_gemm/gemm_xdl_fp16_splitk.cpp
View file @
60b12489
...
...
@@ -207,6 +207,7 @@ int main(int argc, char* argv[])
std
::
cout
<<
"b device buf: "
<<
b_k_n_device_buf
.
GetDeviceBuffer
()
<<
std
::
endl
;
std
::
cout
<<
"c device buf: "
<<
c_m_n_device_buf
.
GetDeviceBuffer
()
<<
std
::
endl
;
#if USEING_STATIC_KERNEL
// do GEMM
if
(
M
==
16
&&
N
==
1152
&&
K
==
5120
&&
splitk
==
8
)
{
...
...
@@ -421,5 +422,63 @@ int main(int argc, char* argv[])
}
}
#else
// dynamic kernel
{
auto
gemm
=
DeviceGemmInstance
{};
auto
invoker
=
gemm
.
MakeInvoker
();
auto
argument
=
gemm
.
MakeArgument
(
static_cast
<
ADataType
*>
(
a_m_k_device_buf
.
GetDeviceBuffer
()),
static_cast
<
BDataType
*>
(
b_k_n_device_buf
.
GetDeviceBuffer
()),
static_cast
<
CDataType
*>
(
c_m_n_device_buf
.
GetDeviceBuffer
()),
M
,
N
,
K
,
StrideA
,
StrideB
,
StrideC
,
a_element_op
,
b_element_op
,
c_element_op
,
splitk
);
if
(
!
gemm
.
IsSupportedArgument
(
argument
))
{
std
::
cout
<<
gemm
.
GetTypeString
()
<<
" does not support this problem"
<<
std
::
endl
;
return
0
;
}
float
ave_time
=
invoker
.
Run
(
argument
,
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
M
*
N
*
K
;
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
N
+
sizeof
(
CDataType
)
*
M
*
N
;
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
ave_time
;
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms, "
<<
tflops
<<
" TFlops, "
<<
gb_per_sec
<<
" GB/s, "
<<
gemm
.
GetTypeString
()
<<
std
::
endl
;
c_m_n_device_buf
.
FromDevice
(
c_m_n_device_result
.
mData
.
data
());
if
(
do_verification
)
{
auto
ref_gemm
=
ReferenceGemmInstance
{};
auto
ref_invoker
=
ref_gemm
.
MakeInvoker
();
auto
ref_argument
=
ref_gemm
.
MakeArgument
(
a_m_k
,
b_k_n
,
c_m_n_host_result
,
a_element_op
,
b_element_op
,
c_element_op
);
ref_invoker
.
Run
(
ref_argument
);
return
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
)
?
0
:
1
;
}
}
#endif
return
0
;
}
include/ck/utility/common_header.hpp
View file @
60b12489
...
...
@@ -48,7 +48,7 @@
#include "amd_xdlops.hpp"
#endif
#define USEING_STATIC_KERNEL
1
#define USEING_STATIC_KERNEL
0
#define MNKB_0_8 0
#define MNKB_1_4 0
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment