From dfb4bf9488879ff34d860df625e4035a2abf1cf4 Mon Sep 17 00:00:00 2001 From: rocking Date: Tue, 22 Oct 2024 20:36:25 +0000 Subject: [PATCH] Fix bug of std caculation --- .../pipeline/layernorm2d_fwd_pipeline_one_pass.hpp | 2 +- .../pipeline/layernorm2d_fwd_pipeline_two_pass.hpp | 4 +--- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp index d73bcb29e4..bf002141b8 100644 --- a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp +++ b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp @@ -90,7 +90,7 @@ struct Layernorm2dFwdPipelineOnePass // compute inv-std auto inv_std = tile_elementwise_in( [&](const auto& v_) { - return type_convert(1.0f) / (sqrt(v_) + epsilon); + return type_convert(1.0f) / (sqrt(v_ + epsilon)); }, var); diff --git a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp index dcbfc87dab..3347c2cb90 100644 --- a/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp +++ b/include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp @@ -105,7 +105,7 @@ struct Layernorm2dFwdPipelineTwoPass // compute inv-std auto inv_std = tile_elementwise_in( [&](const auto& v_) { - return type_convert(1.0f) / (sqrt(v_) + epsilon); + return type_convert(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});