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
26720576
"tests/vscode:/vscode.git/clone" did not exist on "372b58108e19a6065931d6db11015849a0397400"
Commit
26720576
authored
Mar 02, 2023
by
rocking
Browse files
Fix bug of welford count
parent
35cffc47
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
5 additions
and
1 deletion
+5
-1
include/ck/tensor_operation/gpu/grid/normalization/gridwise_normalization_splitk_1st.hpp
.../grid/normalization/gridwise_normalization_splitk_1st.hpp
+5
-1
No files found.
include/ck/tensor_operation/gpu/grid/normalization/gridwise_normalization_splitk_1st.hpp
View file @
26720576
...
@@ -216,12 +216,16 @@ struct GridwiseNormalizationSplitK1st
...
@@ -216,12 +216,16 @@ struct GridwiseNormalizationSplitK1st
});
});
}
}
int
welford_count
=
0
;
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
static_for
<
0
,
MThreadSliceSize
,
1
>
{}([
&
](
auto
I
)
{
if
constexpr
(
I
>
0
)
if
constexpr
(
I
>
0
)
block_sync_lds
();
block_sync_lds
();
int
count
=
threadwise_welford
.
cur_count_
;
int
count
=
threadwise_welford
.
cur_count_
;
BlockwiseWelford
::
Run
(
mean_thread_buf
(
I
),
var_thread_buf
(
I
),
count
);
BlockwiseWelford
::
Run
(
mean_thread_buf
(
I
),
var_thread_buf
(
I
),
count
);
if
constexpr
(
I
==
MThreadSliceSize
-
1
)
welford_count
=
count
;
});
});
if
(
thread_k_cluster_id
==
0
)
if
(
thread_k_cluster_id
==
0
)
...
@@ -239,7 +243,7 @@ struct GridwiseNormalizationSplitK1st
...
@@ -239,7 +243,7 @@ struct GridwiseNormalizationSplitK1st
var_global_val_buf
);
var_global_val_buf
);
if
(
block_m_cluster_id
==
0
&&
thread_m_cluster_id
==
0
)
if
(
block_m_cluster_id
==
0
&&
thread_m_cluster_id
==
0
)
p_welford_count_global
[
block_k_cluster_id
]
=
threadwise_
welford
.
cur
_count
_
;
p_welford_count_global
[
block_k_cluster_id
]
=
welford_count
;
}
}
}
}
};
};
...
...
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