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
762cbddf
Commit
762cbddf
authored
May 08, 2023
by
danyao12
Browse files
fix shuffle GemmORepeat
parent
17774771
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
2 additions
and
2 deletions
+2
-2
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_pt3.hpp
...batched_multihead_attention_backward_xdl_cshuffle_pt3.hpp
+2
-2
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_pt3.hpp
View file @
762cbddf
...
@@ -2164,7 +2164,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V1
...
@@ -2164,7 +2164,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V1
// shuffle dK&dV and write
// shuffle dK&dV and write
{
{
static_assert
(
Gemm2Params_N_O_M
::
GemmNRepeat
%
CShuffleMXdlPerWavePerShuffle
==
0
&&
static_assert
(
Gemm2Params_N_O_M
::
GemmNRepeat
%
CShuffleMXdlPerWavePerShuffle
==
0
&&
Gemm2Params_N_O_M
::
Gemm
N
Repeat
%
CShuffleNXdlPerWavePerShuffle
==
0
,
Gemm2Params_N_O_M
::
Gemm
O
Repeat
%
CShuffleNXdlPerWavePerShuffle
==
0
,
"wrong!"
);
"wrong!"
);
constexpr
index_t
MWave
=
Gemm2Params_N_O_M
::
GemmNWave
;
constexpr
index_t
MWave
=
Gemm2Params_N_O_M
::
GemmNWave
;
...
@@ -2363,7 +2363,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V1
...
@@ -2363,7 +2363,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V1
// space filling curve for threadwise C in VGPR
// space filling curve for threadwise C in VGPR
constexpr
auto
sfc_c_vgpr
=
SpaceFillingCurve
<
Sequence
<
Gemm2Params_N_O_M
::
GemmNRepeat
,
constexpr
auto
sfc_c_vgpr
=
SpaceFillingCurve
<
Sequence
<
Gemm2Params_N_O_M
::
GemmNRepeat
,
Gemm2Params_N_O_M
::
Gemm
N
Repeat
,
Gemm2Params_N_O_M
::
Gemm
O
Repeat
,
1
,
1
,
1
,
1
,
1
,
1
,
...
...
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