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
4347bab1
Commit
4347bab1
authored
Jan 18, 2023
by
ltqin
Browse files
fix P computer
parent
bb06d009
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
4 additions
and
2 deletions
+4
-2
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_v2.hpp
..._batched_multihead_attention_backward_xdl_cshuffle_v2.hpp
+4
-2
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_v2.hpp
View file @
4347bab1
...
@@ -1609,6 +1609,8 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V2
...
@@ -1609,6 +1609,8 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V2
const
index_t
num_gemm1_k_block_outer_loop
=
k_grid_desc_k0_n_k1
.
GetLength
(
I1
)
/
NPerBlock
;
const
index_t
num_gemm1_k_block_outer_loop
=
k_grid_desc_k0_n_k1
.
GetLength
(
I1
)
/
NPerBlock
;
constexpr
index_t
num_gemm1_k_block_inner_loop
=
NPerBlock
/
Gemm1KPerBlock
;
constexpr
index_t
num_gemm1_k_block_inner_loop
=
NPerBlock
/
Gemm1KPerBlock
;
const
index_t
K
=
k_grid_desc_k0_n_k1
.
GetLength
(
I0
)
*
k_grid_desc_k0_n_k1
.
GetLength
(
I2
);
const
float
scale
=
1.0
f
/
std
::
sqrt
(
K
);
// Initialize dQ
// Initialize dQ
qgrad_thread_buf
.
Clear
();
qgrad_thread_buf
.
Clear
();
...
@@ -1689,14 +1691,14 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V2
...
@@ -1689,14 +1691,14 @@ struct GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V2
}
}
else
else
{
{
s_element_op
(
s_slash_p_thread_buf
(
i
)
,
s_slash_p_thread_buf
[
i
]
)
;
s_slash_p_thread_buf
(
i
)
=
scale
*
s_slash_p_thread_buf
[
i
];
}
}
});
});
}
}
else
else
{
{
static_for
<
0
,
s_slash_p_thread_buf
.
Size
(),
1
>
{}(
static_for
<
0
,
s_slash_p_thread_buf
.
Size
(),
1
>
{}(
[
&
](
auto
i
)
{
s_
element_op
(
acc
_thread_buf
(
i
)
,
s_slash_p_thread_buf
[
i
]
)
;
});
[
&
](
auto
i
)
{
s_
slash_p
_thread_buf
(
i
)
=
scale
*
s_slash_p_thread_buf
[
i
];
});
}
}
block_sync_lds
();
// wait for lds read in gemm0 blockwise gemm
block_sync_lds
();
// wait for lds read in gemm0 blockwise gemm
...
...
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