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
44b66c41
"...sencence-bert_pytorch.git" did not exist on "24db6dab25dfbedb8bcb0f54a6bdb90507550cd0"
Commit
44b66c41
authored
Dec 23, 2022
by
rocking
Browse files
Refine naming
parent
14d29856
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
12 additions
and
12 deletions
+12
-12
include/ck/tensor_operation/gpu/grid/gemm_layernorm/gridwise_welford_second_half_layernorm2d.hpp
...mm_layernorm/gridwise_welford_second_half_layernorm2d.hpp
+12
-12
No files found.
include/ck/tensor_operation/gpu/grid/gemm_layernorm/gridwise_welford_second_half_layernorm2d.hpp
View file @
44b66c41
...
@@ -239,7 +239,7 @@ struct GridwiseWelfordSecondHalfLayernorm2d
...
@@ -239,7 +239,7 @@ struct GridwiseWelfordSecondHalfLayernorm2d
block_work_idx
[
I0
]
*
M_BlockTileSize
+
thread_m_cluster_id
*
MThreadSliceSize
,
block_work_idx
[
I0
]
*
M_BlockTileSize
+
thread_m_cluster_id
*
MThreadSliceSize
,
block_work_idx
[
I1
]
*
N_BlockTileSize
+
thread_n_cluster_id
*
NThreadSliceSize
));
block_work_idx
[
I1
]
*
N_BlockTileSize
+
thread_n_cluster_id
*
NThreadSliceSize
));
auto
threadwise_gamma_load_
m_
n
=
auto
threadwise_gamma_load_n
=
ThreadwiseTensorSliceTransfer_v2
<
GammaDataType
,
ThreadwiseTensorSliceTransfer_v2
<
GammaDataType
,
ComputeDataType
,
ComputeDataType
,
decltype
(
gamma_grid_desc_n
),
decltype
(
gamma_grid_desc_n
),
...
@@ -254,7 +254,7 @@ struct GridwiseWelfordSecondHalfLayernorm2d
...
@@ -254,7 +254,7 @@ struct GridwiseWelfordSecondHalfLayernorm2d
make_multi_index
(
block_work_idx
[
I1
]
*
N_BlockTileSize
+
make_multi_index
(
block_work_idx
[
I1
]
*
N_BlockTileSize
+
thread_n_cluster_id
*
NThreadSliceSize
));
thread_n_cluster_id
*
NThreadSliceSize
));
auto
threadwise_beta_load_
m_
n
=
auto
threadwise_beta_load_n
=
ThreadwiseTensorSliceTransfer_v2
<
BetaDataType
,
ThreadwiseTensorSliceTransfer_v2
<
BetaDataType
,
ComputeDataType
,
ComputeDataType
,
decltype
(
beta_grid_desc_n
),
decltype
(
beta_grid_desc_n
),
...
@@ -358,11 +358,11 @@ struct GridwiseWelfordSecondHalfLayernorm2d
...
@@ -358,11 +358,11 @@ struct GridwiseWelfordSecondHalfLayernorm2d
});
});
});
});
threadwise_gamma_load_
m_
n
.
Run
(
gamma_grid_desc_n
,
threadwise_gamma_load_n
.
Run
(
gamma_grid_desc_n
,
gamma_global_val_buf
,
gamma_global_val_buf
,
thread_buffer_desc_n
,
thread_buffer_desc_n
,
make_tuple
(
I0
),
make_tuple
(
I0
),
gamma_thread_buf
);
gamma_thread_buf
);
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
m
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
m
)
{
static_for
<
0
,
NThreadSliceSize
,
1
>
{}([
&
](
auto
n
)
{
static_for
<
0
,
NThreadSliceSize
,
1
>
{}([
&
](
auto
n
)
{
...
@@ -371,11 +371,11 @@ struct GridwiseWelfordSecondHalfLayernorm2d
...
@@ -371,11 +371,11 @@ struct GridwiseWelfordSecondHalfLayernorm2d
});
});
});
});
threadwise_beta_load_
m_
n
.
Run
(
beta_grid_desc_n
,
threadwise_beta_load_n
.
Run
(
beta_grid_desc_n
,
beta_global_val_buf
,
beta_global_val_buf
,
thread_buffer_desc_n
,
thread_buffer_desc_n
,
make_tuple
(
I0
),
make_tuple
(
I0
),
beta_thread_buf
);
beta_thread_buf
);
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
m
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
m
)
{
static_for
<
0
,
NThreadSliceSize
,
1
>
{}([
&
](
auto
n
)
{
static_for
<
0
,
NThreadSliceSize
,
1
>
{}([
&
](
auto
n
)
{
...
...
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