include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_problem.hpp Source File

include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_problem.hpp Source File#

Composable Kernel: include/ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_problem.hpp Source File

 // SPDX-License-Identifier: MIT
 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
  
 #pragma once
  
 #include "ck_tile/core/utility/type_traits.hpp"
  
 namespace ck_tile {
  
 template <typename XDataType_,
           typename GammaDataType_,
           typename ComputeDataType_,
           typename YDataType_,
           typename InvRmsDataType_,
           typename SmoothScaleDataType_,
           typename YScaleDataType_,
           typename BlockShape_,
           typename Traits_>
 struct Rmsnorm2dFwdPipelineProblem
 {
     using XDataType           = remove_cvref_t<XDataType_>;
     using GammaDataType       = remove_cvref_t<GammaDataType_>;
     using ComputeDataType     = remove_cvref_t<ComputeDataType_>;
     using YDataType           = remove_cvref_t<YDataType_>;
     using InvRmsDataType      = remove_cvref_t<InvRmsDataType_>;
     using SmoothScaleDataType = remove_cvref_t<SmoothScaleDataType_>;
     using YScaleDataType      = remove_cvref_t<YScaleDataType_>;
     using BlockShape          = remove_cvref_t<BlockShape_>;
  
     static constexpr bool kNeedCrossLaneSync = BlockShape::ThreadPerWarp_N > 1;
     static constexpr bool kNeedCrossWarpSync = BlockShape::WarpPerBlock_N > 1;
  
     using Traits = remove_cvref_t<Traits_>;
 };
  
 } // namespace ck_tile