page_addresser< T, Layout > Struct Template Reference

page_addresser&lt; T, Layout &gt; Struct Template Reference#

Composable Kernel: ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout > Struct Template Reference
Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout > Struct Template Reference

#include <naive_attention.hpp>

Public Member Functions

__device__ page_addresser (int s_, int h_, int d_, void *base_ptr_, void *pptr_)
__device__ int64_t get_phy_page_idx (int i_s)
__device__ int get_phy_page_offset (int i_s)
__device__ int64_t get_offset (int i_s, int i_d)
__device__ void init (int, int i_h_)
__device__ T load (int i_s, int i_d)
__device__ void store (T, int, int)

Public Attributes

int s
int h
int d
T * base_ptr
int * page_table_ptr
int i_h

Static Public Attributes

static constexpr int x = 16 / sizeof(T)

Constructor & Destructor Documentation

◆ page_addresser()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::page_addresser ( int s_,
int h_,
int d_,
void * base_ptr_,
void * pptr_ )
inline

Member Function Documentation

◆ get_offset()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ int64_t ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::get_offset ( int i_s,
int i_d )
inline

◆ get_phy_page_idx()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ int64_t ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::get_phy_page_idx ( int i_s)
inline

◆ get_phy_page_offset()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::get_phy_page_offset ( int i_s)
inline

◆ init()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ void ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::init ( int ,
int i_h_ )
inline

◆ load()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ T ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::load ( int i_s,
int i_d )
inline

◆ store()

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
__device__ void ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::store ( T ,
int ,
int  )
inline

Member Data Documentation

◆ base_ptr

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
T* ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::base_ptr

◆ d

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::d

◆ h

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::h

◆ i_h

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::i_h

◆ page_table_ptr

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int* ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::page_table_ptr

◆ s

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::s

◆ x

template<typename QType, typename KType, typename VType, typename OType, typename AccType, typename KVScaleType, naive_attention_layout_enum QLayout, naive_attention_layout_enum KLayout, naive_attention_layout_enum VLayout, naive_attention_layout_enum OLayout, naive_attention_layout_enum KScaleLayout, naive_attention_layout_enum VScaleLayout, typename Traits>
template<typename T, naive_attention_layout_enum Layout>
int ck_tile::naive_attention_fwd_kernel< QType, KType, VType, OType, AccType, KVScaleType, QLayout, KLayout, VLayout, OLayout, KScaleLayout, VScaleLayout, Traits >::page_addresser< T, Layout >::x = 16 / sizeof(T)
staticconstexpr
The documentation for this struct was generated from the following file: