TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ > Struct Template Reference

TopkSoftmaxWarpPerRowProblem&lt; InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ > Struct Template Reference

#include <topk_softmax_warp_per_row_problem.hpp>

Public Types

using InputType = remove_cvref_t<InputType_>
using WeightType = remove_cvref_t<WeightType_>
using IndexType = remove_cvref_t<IndexType_>

Static Public Attributes

static constexpr index_t LaunchType = LaunchType_
static constexpr index_t Experts = Experts_
static constexpr index_t BytesPerIssue = BytesPerIssue_
static constexpr index_t IssuesPerCol = IssuesPerCol_
static constexpr index_t BlockSize = BlockSize_
static constexpr index_t WarpSize = get_warp_size()
static constexpr bool ActivationIsSoftmax = ActivationIsSoftmax_
static constexpr index_t VectorSize = BytesPerIssue / sizeof(InputType)
static constexpr index_t LanesPerRow = min(Experts / VectorSize, WarpSize)
static constexpr index_t RowsPerWarpPerColIssue = WarpSize / LanesPerRow
static constexpr index_t RowsPerWarp = IssuesPerCol * RowsPerWarpPerColIssue
static constexpr index_t IssuesPerRow = Experts / (LanesPerRow * VectorSize)
static constexpr index_t WarpsPerBlock = BlockSize / WarpSize
static constexpr index_t RowsPerBlock = RowsPerWarp * WarpsPerBlock

Member Typedef Documentation

◆ IndexType

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
using ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::IndexType = remove_cvref_t<IndexType_>

◆ InputType

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
using ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::InputType = remove_cvref_t<InputType_>

◆ WeightType

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
using ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::WeightType = remove_cvref_t<WeightType_>

Member Data Documentation

◆ ActivationIsSoftmax

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
bool ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::ActivationIsSoftmax = ActivationIsSoftmax_
staticconstexpr

◆ BlockSize

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::BlockSize = BlockSize_
staticconstexpr

◆ BytesPerIssue

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::BytesPerIssue = BytesPerIssue_
staticconstexpr

◆ Experts

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::Experts = Experts_
staticconstexpr

◆ IssuesPerCol

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::IssuesPerCol = IssuesPerCol_
staticconstexpr

◆ IssuesPerRow

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::IssuesPerRow = Experts / (LanesPerRow * VectorSize)
staticconstexpr

◆ LanesPerRow

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::LanesPerRow = min(Experts / VectorSize, WarpSize)
staticconstexpr

◆ LaunchType

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::LaunchType = LaunchType_
staticconstexpr

◆ RowsPerBlock

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::RowsPerBlock = RowsPerWarp * WarpsPerBlock
staticconstexpr

◆ RowsPerWarp

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::RowsPerWarp = IssuesPerCol * RowsPerWarpPerColIssue
staticconstexpr

◆ RowsPerWarpPerColIssue

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::RowsPerWarpPerColIssue = WarpSize / LanesPerRow
staticconstexpr

◆ VectorSize

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::VectorSize = BytesPerIssue / sizeof(InputType)
staticconstexpr

◆ WarpSize

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::WarpSize = get_warp_size()
staticconstexpr

◆ WarpsPerBlock

template<typename InputType_, typename WeightType_, typename IndexType_, index_t Experts_, bool ActivationIsSoftmax_ = true, index_t IssuesPerCol_ = 2, index_t BytesPerIssue_ = sizeof(InputType_), index_t LaunchType_ = 0, index_t BlockSize_ = 256>
index_t ck_tile::TopkSoftmaxWarpPerRowProblem< InputType_, WeightType_, IndexType_, Experts_, ActivationIsSoftmax_, IssuesPerCol_, BytesPerIssue_, LaunchType_, BlockSize_ >::WarpsPerBlock = BlockSize / WarpSize
staticconstexpr
The documentation for this struct was generated from the following file: