ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-06 15:54:31 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix bug of std caculation

Browse Source
This commit is contained in:
rocking
2024-10-22 20:36:25 +00:00
parent 26f16dd20b
commit dfb4bf9488
2 changed files with 2 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp
View File

@@ -90,7 +90,7 @@ struct Layernorm2dFwdPipelineOnePass
// compute inv-std
auto inv_std = tile_elementwise_in(
[&](const auto& v_) {
return type_convert<ComputeDataType>(1.0f) / (sqrt(v_) + epsilon);
return type_convert<ComputeDataType>(1.0f) / (sqrt(v_ + epsilon));
},
var);

4
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp
View File

@@ -105,7 +105,7 @@ struct Layernorm2dFwdPipelineTwoPass
// compute inv-std
auto inv_std = tile_elementwise_in(
[&](const auto& v_) {
return type_convert<ComputeDataType>(1.0f) / (sqrt(v_) + epsilon);
return type_convert<ComputeDataType>(1.0f) / (sqrt(v_ + epsilon));
},
var);
@@ -118,8 +118,6 @@ struct Layernorm2dFwdPipelineTwoPass
ck_tile::index_t stride_to_right_most_window =
row_size % Block_N == 0 ? row_size - Block_N : row_size - row_size % Block_N;
// x_window.foo();
// gamma_window.foo();
move_tile_window(x_window, {0, -Block_N});
move_tile_window(gamma_window, {stride_to_right_most_window});
move_tile_window(beta_window, {stride_to_right_most_window});