From 6cd6bf04fb3c172cf46daf907eb6ae11628f69f4 Mon Sep 17 00:00:00 2001 From: dummycoderfe Date: Fri, 25 Oct 2024 11:17:45 +0800 Subject: [PATCH] hot_fix epsilon pos (#1597) Co-authored-by: dummycoderfe [ROCm/composable_kernel commit: 9183ce69cac01374d0eafbdb4258cf1744b5a548] --- .../layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp | 2 +- .../layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp | 2 +- 2 files changed, 2 insertions(+), 2 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..db094ac2a8 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);