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
90546fbe
Commit
90546fbe
authored
Dec 21, 2022
by
rocking
Browse files
We don't need to pad in N dimension in gemm for mean/var/count. Set NPerTile 1
parent
119cb7b1
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
3 additions
and
2 deletions
+3
-2
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
...ce/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
+3
-2
No files found.
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
View file @
90546fbe
...
@@ -523,12 +523,13 @@ struct DeviceGemmMultipleDLayernorm_Xdl_CShuffle
...
@@ -523,12 +523,13 @@ struct DeviceGemmMultipleDLayernorm_Xdl_CShuffle
gemm_nblock_
{
math
::
integer_divide_ceil
(
NRaw
,
NPerBlock
)},
gemm_nblock_
{
math
::
integer_divide_ceil
(
NRaw
,
NPerBlock
)},
epsilon_
{
static_cast
<
AccDataType
>
(
epsilon
)}
epsilon_
{
static_cast
<
AccDataType
>
(
epsilon
)}
{
{
// We don't need to pad in N dimension in gemm for mean/var/count. Set NPerTile 1.
gemm_mean_var_grid_desc_m_nblock_
=
gemm_mean_var_grid_desc_m_nblock_
=
DeviceOp
::
MakeMeanVarDescriptor_M_N
<
Sequence
<
true
,
false
>
,
MPerBlock
,
NPerBlock
>
(
DeviceOp
::
MakeMeanVarDescriptor_M_N
<
Sequence
<
true
,
false
>
,
MPerBlock
,
1
>
(
MRaw
,
gemm_nblock_
);
MRaw
,
gemm_nblock_
);
gemm_count_grid_desc_m_nblock_
=
gemm_count_grid_desc_m_nblock_
=
DeviceOp
::
MakeCountDescriptor_M_N
<
Sequence
<
true
,
false
>
,
MPerBlock
,
NPerBlock
>
(
DeviceOp
::
MakeCountDescriptor_M_N
<
Sequence
<
true
,
false
>
,
MPerBlock
,
1
>
(
MRaw
,
gemm_nblock_
);
MRaw
,
gemm_nblock_
);
layernorm_mean_var_grid_desc_m_nblock_
=
layernorm_mean_var_grid_desc_m_nblock_
=
...
...
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