WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ > Struct Template Reference

WarpGemmAttributeSmfmac&lt; WarpGemmAttributeSmfmacImpl_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ > Struct Template Reference

Class describing structured sparsity mfma instructions. More...

#include <warp_gemm_attribute_smfmac.hpp>

Public Types

using Impl = remove_cvref_t<WarpGemmAttributeSmfmacImpl_>
using ADataType = typename Impl::ADataType
using BDataType = typename Impl::BDataType
using IdxDataType = typename Impl::IdxDataType
using CDataType = typename Impl::CDataType
using AVecType = typename Impl::AVecType
using BVecType = typename Impl::BVecType
using CVecType = typename Impl::CVecType
using AWarpDstrEncoding
using BWarpDstrEncoding
using CWarpDstrEncoding

Public Member Functions

template<bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, const int32_t &idx, bool_constant< post_nop_ >={}) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto get_num_of_access ()

Static Public Attributes

static constexpr index_t kM = Impl::kM
static constexpr index_t kN = Impl::kN
static constexpr index_t kK = Impl::kK
static constexpr index_t kKPerThread = Impl::kABKPerLane
static constexpr index_t kCompressionRatio = Impl::CompressionRatio

Detailed Description

template<typename WarpGemmAttributeSmfmacImpl_>
struct ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >

Class describing structured sparsity mfma instructions.

Overview

Currently only 2:4 structured sparsity is supported, which is based on requirement that in every groups of four continuous elements there are at most two non-zero, which results in processing only half of elements in smfmac instruction. Because of structured sparsity A vector in smfmac instruction will be smaller than B vector by the factor of CompressionRatio. The indexes of non-zero elements are stored in index which is an additional parameter to assembly instruction. Every pair of two bit indexes are containing information about which two elements in current group of 4 values are non-zero and should be used inside smfmac instruction. Structured sparsity format is supported only for A matrix for now.

Member Typedef Documentation

◆ ADataType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::ADataType = typename Impl::ADataType

◆ AVecType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::AVecType = typename Impl::AVecType

◆ AWarpDstrEncoding

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::AWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kAMLane>, sequence<Impl::kABKLane, Impl::kABKPerLane>>,
tuple<sequence<2, 1>>,
tuple<sequence<0, 0>>,
sequence<2>,
sequence<1>>
ck_tile::sequence<>
ck_tile::tile_distribution_encoding
Definition tile_distribution_encoding.hpp:26
ck_tile::tuple
Definition tile/core/container/tuple.hpp:192

◆ BDataType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::BDataType = typename Impl::BDataType

◆ BVecType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::BVecType = typename Impl::BVecType

◆ BWarpDstrEncoding

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::BWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kBNLane>, sequence<Impl::kABKLane, Impl::kABKPerLane>>,
tuple<sequence<2, 1>>,
tuple<sequence<0, 0>>,
sequence<2>,
sequence<1>>

◆ CDataType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::CDataType = typename Impl::CDataType

◆ CVecType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::CVecType = typename Impl::CVecType

◆ CWarpDstrEncoding

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::CWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kCM0PerLane, Impl::kCMLane, Impl::kCM1PerLane>,
sequence<Impl::kCNLane>>,
tuple<sequence<1, 2>>,
tuple<sequence<1, 0>>,
sequence<1, 1>,
sequence<0, 2>>

◆ IdxDataType

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::IdxDataType = typename Impl::IdxDataType

◆ Impl

template<typename WarpGemmAttributeSmfmacImpl_>
using ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::Impl = remove_cvref_t<WarpGemmAttributeSmfmacImpl_>

Member Function Documentation

◆ get_num_of_access()

template<typename WarpGemmAttributeSmfmacImpl_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::get_num_of_access ( )
inlinestaticconstexpr

◆ operator()()

template<typename WarpGemmAttributeSmfmacImpl_>
template<bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::operator() ( CVecType & c_vec,
const AVecType & a_vec,
const BVecType & b_vec,
const int32_t & idx,
bool_constant< post_nop_ > = {} ) const
inline

Member Data Documentation

◆ kCompressionRatio

template<typename WarpGemmAttributeSmfmacImpl_>
index_t ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::kCompressionRatio = Impl::CompressionRatio
staticconstexpr

◆ kK

template<typename WarpGemmAttributeSmfmacImpl_>
index_t ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::kK = Impl::kK
staticconstexpr

◆ kKPerThread

template<typename WarpGemmAttributeSmfmacImpl_>
index_t ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::kKPerThread = Impl::kABKPerLane
staticconstexpr

◆ kM

template<typename WarpGemmAttributeSmfmacImpl_>
index_t ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::kM = Impl::kM
staticconstexpr

◆ kN

template<typename WarpGemmAttributeSmfmacImpl_>
index_t ck_tile::WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImpl_ >::kN = Impl::kN
staticconstexpr
The documentation for this struct was generated from the following file: