diff --git a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_model_sensitive_pass.hpp b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_model_sensitive_pass.hpp index 810c3c5243..c5923ba10d 100644 --- a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_model_sensitive_pass.hpp +++ b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_model_sensitive_pass.hpp @@ -221,7 +221,7 @@ struct Rmsnorm2dFwdPipelineModelSensitiveT5Pass } else { - Epilogue{}(y_window_, rmsn); + Epilogue{}(y_window_, rmsn, nullptr); } } }; diff --git a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp index c77d61872e..39d7c65d3e 100644 --- a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp +++ b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_one_pass.hpp @@ -160,7 +160,7 @@ struct Rmsnorm2dFwdPipelineOnePass } else { - Epilogue{}(y_window_, rmsn); + Epilogue{}(y_window_, rmsn, nullptr); } } }; diff --git a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp index 4ca1dbc5da..d01f37879a 100644 --- a/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp +++ b/include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp @@ -195,7 +195,7 @@ struct Rmsnorm2dFwdPipelineTwoPass }); static_assert(kFusedQuant == Rmsnorm2dFusedQuantEnum::NO_SWEEP); - Epilogue{}(y_window, rmsn); + Epilogue{}(y_window, rmsn, nullptr); move_tile_window(gamma_window, {-Block_N}); move_tile_window(y_window, {0, -Block_N});