BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ > Struct Template Reference

BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ > Struct Template Reference#

Composable Kernel: ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ > Struct Template Reference

#include <block_dropout.hpp>

Public Member Functions
CK_TILE_HOST_DEVICE	BlockDropoutBwd (index_t i_batch, index_t i_head, index_t nheads, unsigned long long seed, unsigned long long offset, float rp_undrop_, uint8_t p_undrop_in_uint8_t_)
template<typename BlockGemm, typename RandValOutputDataType, typename PComputeWindow, typename RandValDramWindow>
CK_TILE_HOST_DEVICE void	Run (const index_t start_m0_idx, const index_t start_n0_idx, PComputeWindow &p_compute, RandValDramWindow &randval_dram_window) const

Static Public Member Functions
template<typename BlockGemm, bool IsFwd = false, typename RandValDramBlockWindowTmp>
static CK_TILE_HOST_DEVICE constexpr auto	MakeRandvalDramWindow (RandValDramBlockWindowTmp &randval_dram_block_window_tmp, index_t seqlen_qk_start)
template<typename BlockGemm>
static CK_TILE_HOST_DEVICE constexpr auto	MakeRandValTileDistribution ()

Public Attributes
const unsigned long long	ph_seed
const unsigned long long	ph_head_offset
const float	rp_undrop
const uint8_t	p_undrop_in_uint8_t

Static Public Attributes
static constexpr bool	IsDropout = true
static constexpr bool	IsStoreRandval = IsStoreRandval_

Constructor & Destructor Documentation

◆ BlockDropoutBwd()

template<bool IsWG32_, bool IsStoreRandval_>

CK_TILE_HOST_DEVICE ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::BlockDropoutBwd	(	index_t	i_batch,
		index_t	i_head,
		index_t	nheads,
		unsigned long long	seed,
		unsigned long long	offset,
		float	rp_undrop_,
		uint8_t	p_undrop_in_uint8_t_ )

inline

Member Function Documentation

◆ MakeRandvalDramWindow()

template<bool IsWG32_, bool IsStoreRandval_>

template<typename BlockGemm, bool IsFwd = false, typename RandValDramBlockWindowTmp>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::MakeRandvalDramWindow	(	RandValDramBlockWindowTmp &	randval_dram_block_window_tmp,
		index_t	seqlen_qk_start )

inlinestaticconstexpr

◆ MakeRandValTileDistribution()

template<bool IsWG32_, bool IsStoreRandval_>

template<typename BlockGemm>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::MakeRandValTileDistribution ( )

inlinestaticconstexpr

◆ Run()

template<bool IsWG32_, bool IsStoreRandval_>

template<typename BlockGemm, typename RandValOutputDataType, typename PComputeWindow, typename RandValDramWindow>

CK_TILE_HOST_DEVICE void ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::Run	(	const index_t	start_m0_idx,
		const index_t	start_n0_idx,
		PComputeWindow &	p_compute,
		RandValDramWindow &	randval_dram_window ) const

inline

Member Data Documentation

◆ IsDropout

template<bool IsWG32_, bool IsStoreRandval_>

bool ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::IsDropout = true

staticconstexpr

◆ IsStoreRandval

template<bool IsWG32_, bool IsStoreRandval_>

bool ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::IsStoreRandval = IsStoreRandval_

staticconstexpr

◆ p_undrop_in_uint8_t

template<bool IsWG32_, bool IsStoreRandval_>

const uint8_t ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::p_undrop_in_uint8_t

◆ ph_head_offset

template<bool IsWG32_, bool IsStoreRandval_>

const unsigned long long ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::ph_head_offset

◆ ph_seed

template<bool IsWG32_, bool IsStoreRandval_>

const unsigned long long ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::ph_seed

◆ rp_undrop

template<bool IsWG32_, bool IsStoreRandval_>

const float ck_tile::BlockDropoutBwd< true, IsWG32_, IsStoreRandval_ >::rp_undrop

The documentation for this struct was generated from the following file:

block_dropout.hpp