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
7a190876
"docs/source/en/vscode:/vscode.git/clone" did not exist on "dbcb15c25fee7122c5184eb0e0fcef8e29495227"
Commit
7a190876
authored
Nov 03, 2022
by
Qianfeng Zhang
Browse files
Add block_sync_lds between two consecutive blockwise reductions
parent
bb220a7a
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
6 additions
and
0 deletions
+6
-0
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_reduce_second_half_batchnorm_backward_final.hpp
...ultiblock_reduce_second_half_batchnorm_backward_final.hpp
+3
-0
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_welford_second_half_multiblock_reduce_first_half.hpp
...lock_welford_second_half_multiblock_reduce_first_half.hpp
+3
-0
No files found.
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_reduce_second_half_batchnorm_backward_final.hpp
View file @
7a190876
...
...
@@ -330,6 +330,9 @@ struct GridwiseReduceSecondHalfBatchNormBackwardFinal
}
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
if
constexpr
(
I
>
0
)
block_sync_lds
();
BlockwiseReduce
::
Reduce
(
reduce_work_buf
,
dscale_thread_buf
(
I
));
block_sync_lds
();
BlockwiseReduce
::
Reduce
(
reduce_work_buf
,
dbias_thread_buf
(
I
));
...
...
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_welford_second_half_multiblock_reduce_first_half.hpp
View file @
7a190876
...
...
@@ -495,6 +495,9 @@ struct GridwiseWelfordSecondHalfReduceFirstHalf
};
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
if
constexpr
(
I
>
0
)
block_sync_lds
();
BlockwiseReduce
::
Reduce
(
reduce_work_buf
,
reduce_dscale_thread_buf
(
I
));
block_sync_lds
();
BlockwiseReduce
::
Reduce
(
reduce_work_buf
,
reduce_dbias_thread_buf
(
I
));
...
...
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