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
c33d0d64
Commit
c33d0d64
authored
Nov 02, 2022
by
Qianfeng Zhang
Browse files
Add comments to explain inv-variance in batchnorm forward and backward
parent
d2000114
Changes
4
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
4 additions
and
1 deletion
+4
-1
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_welford_second_half_batchnorm_forward_final.hpp
...ultiblock_welford_second_half_batchnorm_forward_final.hpp
+1
-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
+1
-1
include/ck/tensor_operation/gpu/grid/gridwise_batchnorm_backward_blockwise_welford.hpp
...pu/grid/gridwise_batchnorm_backward_blockwise_welford.hpp
+1
-0
include/ck/tensor_operation/gpu/grid/gridwise_batchnorm_forward_blockwise_welford.hpp
...gpu/grid/gridwise_batchnorm_forward_blockwise_welford.hpp
+1
-0
No files found.
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_welford_second_half_batchnorm_forward_final.hpp
View file @
c33d0d64
...
@@ -529,6 +529,7 @@ struct GridwiseWelfordSecondHalfBatchNormForwardFinal
...
@@ -529,6 +529,7 @@ struct GridwiseWelfordSecondHalfBatchNormForwardFinal
auto
result_inv_var_global_buf
=
make_dynamic_buffer
<
AddressSpaceEnum
::
Global
>
(
auto
result_inv_var_global_buf
=
make_dynamic_buffer
<
AddressSpaceEnum
::
Global
>
(
resultSaveInvVariance
,
mean_var_grid_desc_m
.
GetElementSpaceSize
());
resultSaveInvVariance
,
mean_var_grid_desc_m
.
GetElementSpaceSize
());
// calculate inv-variance as 1/sqrt(epsilon+variance)
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
welford_var_thread_buf
(
I
)
=
welford_var_thread_buf
(
I
)
=
type_convert
<
AccDataType
>
(
1.0
f
)
/
sqrt
(
epsilon
+
welford_var_thread_buf
[
I
]);
type_convert
<
AccDataType
>
(
1.0
f
)
/
sqrt
(
epsilon
+
welford_var_thread_buf
[
I
]);
...
...
include/ck/tensor_operation/gpu/grid/batchnorm_multiblock/gridwise_multiblock_welford_second_half_multiblock_reduce_first_half.hpp
View file @
c33d0d64
...
@@ -359,7 +359,7 @@ struct GridwiseWelfordSecondHalfReduceFirstHalf
...
@@ -359,7 +359,7 @@ struct GridwiseWelfordSecondHalfReduceFirstHalf
welford_count_thread_buf
(
I
));
welford_count_thread_buf
(
I
));
});
});
// calculate inv-variance
from
variance, stored in place
// calculate inv-variance
as 1/sqrt(epsilon+
variance
)
, stored in place
of variance
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
welford_var_thread_buf
(
I
)
=
welford_var_thread_buf
(
I
)
=
type_convert
<
AccDataType
>
(
1.0
)
/
sqrt
(
welford_var_thread_buf
[
I
]
+
epsilon
);
type_convert
<
AccDataType
>
(
1.0
)
/
sqrt
(
welford_var_thread_buf
[
I
]
+
epsilon
);
...
...
include/ck/tensor_operation/gpu/grid/gridwise_batchnorm_backward_blockwise_welford.hpp
View file @
c33d0d64
...
@@ -410,6 +410,7 @@ struct GridwiseBatchNormBackwardWithBlockwiseWelford
...
@@ -410,6 +410,7 @@ struct GridwiseBatchNormBackwardWithBlockwiseWelford
BlockwiseWelford
::
Run
(
mean_thread_buf
(
I
),
var_thread_buf
(
I
),
count
);
BlockwiseWelford
::
Run
(
mean_thread_buf
(
I
),
var_thread_buf
(
I
),
count
);
});
});
// calculate inv-variance as 1/sqrt(epsilon+variance)
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
inv_var_thread_buf
(
I
)
=
inv_var_thread_buf
(
I
)
=
type_convert
<
AccDataType
>
(
1.0
)
/
sqrt
(
var_thread_buf
[
I
]
+
epsilon
);
type_convert
<
AccDataType
>
(
1.0
)
/
sqrt
(
var_thread_buf
[
I
]
+
epsilon
);
...
...
include/ck/tensor_operation/gpu/grid/gridwise_batchnorm_forward_blockwise_welford.hpp
View file @
c33d0d64
...
@@ -441,6 +441,7 @@ struct GridwiseBatchNormForwardWithBlockwiseWelford
...
@@ -441,6 +441,7 @@ struct GridwiseBatchNormForwardWithBlockwiseWelford
auto
result_inv_var_global_buf
=
make_dynamic_buffer
<
AddressSpaceEnum
::
Global
>
(
auto
result_inv_var_global_buf
=
make_dynamic_buffer
<
AddressSpaceEnum
::
Global
>
(
resultSaveInvVariance
,
mean_var_grid_desc_m
.
GetElementSpaceSize
());
resultSaveInvVariance
,
mean_var_grid_desc_m
.
GetElementSpaceSize
());
// calculate inv-variance as 1/sqrt(epsilon+variance), stored in place of variance
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
var_thread_buf
(
I
)
=
var_thread_buf
(
I
)
=
type_convert
<
AccDataType
>
(
1.0
f
)
/
sqrt
(
epsilon
+
var_thread_buf
[
I
]);
type_convert
<
AccDataType
>
(
1.0
f
)
/
sqrt
(
epsilon
+
var_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