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
f0e7bc99
Commit
f0e7bc99
authored
Aug 19, 2021
by
Jing Zhang
Browse files
enable fp16 mfma
parent
cf378e46
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
8 deletions
+8
-8
composable_kernel/include/tensor_operation/xdlops_gemm.hpp
composable_kernel/include/tensor_operation/xdlops_gemm.hpp
+8
-8
No files found.
composable_kernel/include/tensor_operation/xdlops_gemm.hpp
View file @
f0e7bc99
...
@@ -506,49 +506,49 @@ struct MfmaSelector
...
@@ -506,49 +506,49 @@ struct MfmaSelector
return
MfmaInstr
::
mfma_f32_16x16x4xf32
;
return
MfmaInstr
::
mfma_f32_16x16x4xf32
;
}
}
#if 0
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
64
,
64
>
()
static
constexpr
auto
GetMfma
<
half_t
,
64
,
64
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_32x32x4f16
, 64, 64>{}
;
return
MfmaInstr
::
mfma_f32_32x32x4f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
32
,
64
>
()
static
constexpr
auto
GetMfma
<
half_t
,
32
,
64
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_32x32x4f16
, 32, 64>{}
;
return
MfmaInstr
::
mfma_f32_32x32x4f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
32
,
32
>
()
static
constexpr
auto
GetMfma
<
half_t
,
32
,
32
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_32x32x8f16
, 32, 32>{}
;
return
MfmaInstr
::
mfma_f32_32x32x8f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
16
,
16
>
()
static
constexpr
auto
GetMfma
<
half_t
,
16
,
16
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_16x16x16f16
, 16, 16>{}
;
return
MfmaInstr
::
mfma_f32_16x16x16f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
16
,
64
>
()
static
constexpr
auto
GetMfma
<
half_t
,
16
,
64
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_16x16x4f16
, 16, 64>{}
;
return
MfmaInstr
::
mfma_f32_16x16x4f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
8
,
64
>
()
static
constexpr
auto
GetMfma
<
half_t
,
8
,
64
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_4x4x4f16
, 8, 64>{}
;
return
MfmaInstr
::
mfma_f32_4x4x4f16
;
}
}
template
<
>
template
<
>
static
constexpr
auto
GetMfma
<
half_t
,
4
,
64
>
()
static
constexpr
auto
GetMfma
<
half_t
,
4
,
64
>
()
{
{
return
xdlops_info<
MfmaInstr::mfma_f32_4x4x4f16
, 4, 64>{}
;
return
MfmaInstr
::
mfma_f32_4x4x4f16
;
}
}
#if 0
template <>
template <>
static constexpr auto GetMfma<ushort, 128, 64>()
static constexpr auto GetMfma<ushort, 128, 64>()
{
{
...
...
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