TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ > Struct Template Reference

TileFmhaBwdShape&lt; BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ > Struct Template Reference

#include <tile_fmha_shape.hpp>

Public Types

using BlockTile = remove_cvref_t<BlockTile_>
using Gemm0BlockWarps = remove_cvref_t<Gemm0BlockWarps_>
using Gemm0WarpTile = remove_cvref_t<Gemm0WarpTile_>
using Gemm1BlockWarps = remove_cvref_t<Gemm1BlockWarps_>
using Gemm1WarpTile = remove_cvref_t<Gemm1WarpTile_>
using Gemm2BlockWarps = remove_cvref_t<Gemm2BlockWarps_>
using Gemm2WarpTile = remove_cvref_t<Gemm2WarpTile_>
using Gemm3BlockWarps = remove_cvref_t<Gemm3BlockWarps_>
using Gemm3WarpTile = remove_cvref_t<Gemm3WarpTile_>
using Gemm4BlockWarps = remove_cvref_t<Gemm4BlockWarps_>
using Gemm4WarpTile = remove_cvref_t<Gemm4WarpTile_>

Static Public Attributes

static constexpr index_t NumWarps
static constexpr index_t kM0 = BlockTile::at(number<0>{})
static constexpr index_t kN0 = BlockTile::at(number<1>{})
static constexpr index_t kK0
static constexpr index_t kK1
static constexpr index_t kK2
static constexpr index_t kK3
static constexpr index_t kK4 = BlockTile::at(number<6>{})
static constexpr index_t kQKHeaddim
static constexpr index_t kVHeaddim = BlockTile::at(number<8>{})
static constexpr index_t kMaxSeqLenQ = kMaxSeqLenQ_

Member Typedef Documentation

◆ BlockTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::BlockTile = remove_cvref_t<BlockTile_>

◆ Gemm0BlockWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm0BlockWarps = remove_cvref_t<Gemm0BlockWarps_>

◆ Gemm0WarpTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm0WarpTile = remove_cvref_t<Gemm0WarpTile_>

◆ Gemm1BlockWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm1BlockWarps = remove_cvref_t<Gemm1BlockWarps_>

◆ Gemm1WarpTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm1WarpTile = remove_cvref_t<Gemm1WarpTile_>

◆ Gemm2BlockWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm2BlockWarps = remove_cvref_t<Gemm2BlockWarps_>

◆ Gemm2WarpTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm2WarpTile = remove_cvref_t<Gemm2WarpTile_>

◆ Gemm3BlockWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm3BlockWarps = remove_cvref_t<Gemm3BlockWarps_>

◆ Gemm3WarpTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm3WarpTile = remove_cvref_t<Gemm3WarpTile_>

◆ Gemm4BlockWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm4BlockWarps = remove_cvref_t<Gemm4BlockWarps_>

◆ Gemm4WarpTile

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
using ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::Gemm4WarpTile = remove_cvref_t<Gemm4WarpTile_>

Member Data Documentation

◆ kK0

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kK0
staticconstexpr
Initial value:
=
BlockTile::at(number<2>{})
ck_tile::number
constant< v > number
Definition tile/core/numeric/integral_constant.hpp:37

◆ kK1

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kK1
staticconstexpr
Initial value:
=
BlockTile::at(number<3>{})

◆ kK2

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kK2
staticconstexpr
Initial value:
=
BlockTile::at(number<4>{})

◆ kK3

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kK3
staticconstexpr
Initial value:
=
BlockTile::at(number<5>{})

◆ kK4

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kK4 = BlockTile::at(number<6>{})
staticconstexpr

◆ kM0

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kM0 = BlockTile::at(number<0>{})
staticconstexpr

◆ kMaxSeqLenQ

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kMaxSeqLenQ = kMaxSeqLenQ_
staticconstexpr

◆ kN0

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kN0 = BlockTile::at(number<1>{})
staticconstexpr

◆ kQKHeaddim

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kQKHeaddim
staticconstexpr
Initial value:
=
BlockTile::at(number<7>{})

◆ kVHeaddim

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::kVHeaddim = BlockTile::at(number<8>{})
staticconstexpr

◆ NumWarps

template<typename BlockTile_, typename Gemm0BlockWarps_, typename Gemm0WarpTile_, typename Gemm1BlockWarps_, typename Gemm1WarpTile_, typename Gemm2BlockWarps_, typename Gemm2WarpTile_, typename Gemm3BlockWarps_, typename Gemm3WarpTile_, typename Gemm4BlockWarps_, typename Gemm4WarpTile_, index_t kMaxSeqLenQ_ = 0>
index_t ck_tile::TileFmhaBwdShape< BlockTile_, Gemm0BlockWarps_, Gemm0WarpTile_, Gemm1BlockWarps_, Gemm1WarpTile_, Gemm2BlockWarps_, Gemm2WarpTile_, Gemm3BlockWarps_, Gemm3WarpTile_, Gemm4BlockWarps_, Gemm4WarpTile_, kMaxSeqLenQ_ >::NumWarps
staticconstexpr
Initial value:
=
reduce_on_sequence(Gemm0BlockWarps{}, multiplies{}, number<1>{})
ck_tile::reduce_on_sequence
CK_TILE_HOST_DEVICE constexpr index_t reduce_on_sequence(Seq, Reduce f, number< Init >)
Definition tile/core/container/sequence.hpp:982
ck_tile::TileFmhaBwdShape::Gemm0BlockWarps
remove_cvref_t< Gemm0BlockWarps_ > Gemm0BlockWarps
Definition tile_fmha_shape.hpp:84
ck_tile::multiplies
Definition tile/core/numeric/math.hpp:98
The documentation for this struct was generated from the following file: