/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/add_rmsnorm2d_rdquant/pipeline/add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp Source File

• docs-6.4.3
• include
• ck_tile
• ops
• add_rmsnorm2d_rdquant
• pipeline

add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

2 // Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.

3 

4 #pragma once

5 

6 #include "ck_tile/core/utility/type_traits.hpp"

7 

8 namespace ck_tile {

9 

10 // X = A + B, Y = Rmsnorm2d(X), QY = RowwiseDynamicQuant(Y) = SaturateCast(Y / YScale)

11 template <typename ADataType_,

12  typename BDataType_,

13  typename GammaDataType_,

14  typename ComputeDataType_,

15  typename XDataType_,

16  typename YScaleDataType_,

17  typename QYDataType_,

18  typename BlockShape_,

19  bool kPadN_,

20  bool kSaveX_,

21  bool kThreePass_>

22 struct AddRmsnorm2dRdquantFwdPipelineProblem

23 {

24  using ADataType = remove_cvref_t<ADataType_>;

25  using BDataType = remove_cvref_t<BDataType_>;

26  using GammaDataType = remove_cvref_t<GammaDataType_>;

27  using ComputeDataType = remove_cvref_t<ComputeDataType_>;

28  using XDataType = remove_cvref_t<XDataType_>;

29  using YScaleDataType = remove_cvref_t<YScaleDataType_>;

30  using QYDataType = remove_cvref_t<QYDataType_>;

31  using BlockShape = remove_cvref_t<BlockShape_>;

32 

33  static constexpr bool kNeedCrossLaneSync = BlockShape::ThreadPerWarp_N > 1;

34  static constexpr bool kNeedCrossWarpSync = BlockShape::WarpPerBlock_N > 1;

35 

36  static constexpr bool kPadN = kPadN_;

37  static constexpr bool kSaveX = kSaveX_;

38  static constexpr bool kThreePass = kThreePass_;

39 };

40 

41 } // namespace ck_tile

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::remove_cvref_t

remove_cv_t< std::remove_reference_t< T > > remove_cvref_t

Definition: type_traits.hpp:20

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:23

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::kPadN

static constexpr bool kPadN

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:36

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::ADataType

remove_cvref_t< ADataType_ > ADataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:24

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::kNeedCrossLaneSync

static constexpr bool kNeedCrossLaneSync

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:33

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::kNeedCrossWarpSync

static constexpr bool kNeedCrossWarpSync

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:34

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::kThreePass

static constexpr bool kThreePass

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:38

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::QYDataType

remove_cvref_t< QYDataType_ > QYDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:30

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::ComputeDataType

remove_cvref_t< ComputeDataType_ > ComputeDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:27

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::kSaveX

static constexpr bool kSaveX

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:37

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::GammaDataType

remove_cvref_t< GammaDataType_ > GammaDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:26

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::BlockShape

remove_cvref_t< BlockShape_ > BlockShape

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:31

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::BDataType

remove_cvref_t< BDataType_ > BDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:25

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::XDataType

remove_cvref_t< XDataType_ > XDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:28

ck_tile::AddRmsnorm2dRdquantFwdPipelineProblem::YScaleDataType

remove_cvref_t< YScaleDataType_ > YScaleDataType

Definition: add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp:29

type_traits.hpp