use raw store to improve performance

example/ck_tile/02_layernorm2d/generate.py

@@ -204,7 +204,7 @@ float layernorm2d_fwd_(const S& s, A a)
     using TwoPassPipeline = ck_tile::Layernorm2dFwdPipelineTwoPass<PipelineProblem>;
     using Pipeline        = std::conditional_t<Traits_::kTwoPass, TwoPassPipeline, OnePassPipeline>;
 
-    using Default2DEpilogueProblem = ck_tile::Default2DEpilogueProblem<ComputeDataType, YDataType, false, Traits_::kPadN, false>;
+    using Default2DEpilogueProblem = ck_tile::Default2DEpilogueProblem<ComputeDataType, YDataType, false, Traits_::kPadN, true>;
     using Default2DEpilogue = ck_tile::Default2DEpilogue<Default2DEpilogueProblem>;
 
     static constexpr bool UseSmoothInputScale = Traits_::kFusedQuant == 1;
@@ -327,7 +327,7 @@ float layernorm2d_fwd(layernorm2d_fwd_traits t,
     @dataclass
     class k_pipeline_one_pass:
         F_Problem         : Any #k_problem
-    
+
     @dataclass
     class k_pipeline_two_pass:
         F_Problem         : Any #k_problem