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
59088eca
Commit
59088eca
authored
May 26, 2023
by
danyao12
Browse files
rename sgrad_slice_idx
parent
a9f57b62
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
7 additions
and
6 deletions
+7
-6
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_pt4.hpp
...batched_multihead_attention_backward_xdl_cshuffle_pt4.hpp
+7
-6
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_pt4.hpp
View file @
59088eca
...
...
@@ -2100,20 +2100,21 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V1
qgrad_thread_buf
.
Clear
();
static_for
<
0
,
num_gemm2_loop
,
1
>
{}([
&
](
auto
gemm2_loop_idx
)
{
// gemm dQ
// load VGrad Gemm A
const
auto
p
_slice_idx
=
const
auto
sgrad
_slice_idx
=
Gemm2
::
ASrcBlockSliceWindowIterator
::
GetIndexTupleOfNumber
(
gemm2_loop_idx
);
constexpr
auto
mwave_range
=
make_tuple
(
p
_slice_idx
[
I2
],
p
_slice_idx
[
I2
]
+
Gemm2Params
::
ABlockSliceLengths_M0_K0_M1_K1
::
At
(
I2
));
sgrad
_slice_idx
[
I2
],
sgrad
_slice_idx
[
I2
]
+
Gemm2Params
::
ABlockSliceLengths_M0_K0_M1_K1
::
At
(
I2
));
constexpr
auto
nwave_range
=
make_tuple
(
p
_slice_idx
[
I3
],
p
_slice_idx
[
I3
]
+
Gemm2Params
::
ABlockSliceLengths_M0_K0_M1_K1
::
At
(
I3
));
sgrad
_slice_idx
[
I3
],
sgrad
_slice_idx
[
I3
]
+
Gemm2Params
::
ABlockSliceLengths_M0_K0_M1_K1
::
At
(
I3
));
if
(
gemm2_a_copy_subgroup
.
IsBelong
(
mwave_range
,
nwave_range
))
{
qgrad_gemm_tile_sgrad_thread_copy_vgpr_to_lds
.
Run
(
Gemm2
::
a_src_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2
,
make_tuple
(
p_slice_idx
[
I0
],
p_slice_idx
[
I1
],
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
make_tuple
(
sgrad_slice_idx
[
I0
],
sgrad_slice_idx
[
I1
],
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
sgrad_thread_buf
,
Gemm2
::
a_block_desc_m0_k0_m1_k1_m2_m3_m4_k2
,
gemm2_a_block_buf
);
...
...
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