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
0bfd650b
"...composable_kernel.git" did not exist on "4e097ad283d5a2d977501c61d0f8c3081dfa35f6"
Commit
0bfd650b
authored
Jun 21, 2023
by
ltqin
Browse files
regular code2
parent
bab3161b
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
1 addition
and
9 deletions
+1
-9
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_bacckward_ydotygrad.hpp
...dwise_batched_multihead_attention_bacckward_ydotygrad.hpp
+1
-9
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_bacckward_ydotygrad.hpp
View file @
0bfd650b
...
@@ -34,16 +34,9 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
...
@@ -34,16 +34,9 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
static
constexpr
auto
I2
=
Number
<
2
>
{};
static
constexpr
auto
I2
=
Number
<
2
>
{};
static
constexpr
auto
I3
=
Number
<
3
>
{};
static
constexpr
auto
I3
=
Number
<
3
>
{};
static
constexpr
auto
I4
=
Number
<
4
>
{};
static
constexpr
auto
I4
=
Number
<
4
>
{};
static
constexpr
auto
I5
=
Number
<
5
>
{};
static
constexpr
auto
I6
=
Number
<
6
>
{};
static
constexpr
auto
I7
=
Number
<
7
>
{};
static
constexpr
auto
I8
=
Number
<
8
>
{};
static
constexpr
auto
I9
=
Number
<
9
>
{};
static
constexpr
auto
WaveSize
=
64
;
static
constexpr
auto
WaveSize
=
64
;
using
ThisThreadBlock
=
ThisThreadBlock
<
BlockSize
>
;
// block_id to matrix tile idx (m0, n0) mapping are controlled by {M01, N01}
// block_id to matrix tile idx (m0, n0) mapping are controlled by {M01, N01}
template
<
typename
Block2CTileMap
>
template
<
typename
Block2CTileMap
>
__host__
__device__
static
constexpr
bool
CheckValidity
(
const
CGridDesc_M_N
&
c_grid_desc_m_n
,
__host__
__device__
static
constexpr
bool
CheckValidity
(
const
CGridDesc_M_N
&
c_grid_desc_m_n
,
...
@@ -261,7 +254,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
...
@@ -261,7 +254,7 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
1
,
1
,
false
>
{
ors_grid_desc_mblock_mperblock
,
false
>
{
ors_grid_desc_mblock_mperblock
,
make_multi_index
(
block_work_idx_m
,
// mblock
make_multi_index
(
block_work_idx_m
,
// mblock
get_thread_local_1d_id
()),
// mper
xdl
get_thread_local_1d_id
()),
// mper
block
ck
::
tensor_operation
::
element_wise
::
PassThrough
{}};
ck
::
tensor_operation
::
element_wise
::
PassThrough
{}};
// copy from VGPR to Global
// copy from VGPR to Global
...
@@ -270,7 +263,6 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
...
@@ -270,7 +263,6 @@ struct GridwiseBatchedMultiheadAttentionBackward_YDotYGrad
y_dot_ygrad_thread_accum_buf
,
y_dot_ygrad_thread_accum_buf
,
ors_grid_desc_mblock_mperblock
,
ors_grid_desc_mblock_mperblock
,
ors_grid_buf
);
ors_grid_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