/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fmha/block/block_masking.hpp Source File#
block_masking.hpp
Go to the documentation of this file.
200 // Attention! assume the idex passed in this function is with in range of GetTileRangeAlongX/Y()
204 IsEdgeTile(index_t i_tile_top, index_t i_tile_left, number<TileHeight>, number<TileWidth>) const
369 // Attention! assume the idex passed in this function is with in range of GetTileRangeAlongX/Y()
407 template<> struct SimplifiedRatioMaskName<false> { static constexpr const char * name = "nomask"; };
408 template<> struct SimplifiedRatioMaskName<true> { static constexpr const char * name = "mask"; };
412 // this version is used for cases that the step length of y-direction changes greater than one. It
556 // Attention! assume the idex passed in this function is with in range of GetTileRangeAlongX/Y()
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_HOST_DEVICE auto integer_divide_ceil(X x, Y y)
Definition: math.hpp:149
constexpr CK_TILE_HOST_DEVICE auto make_generic_attention_mask_from_lr_window(index_t left_size, index_t right_size, index_t y_total, index_t x_total, bool is_top_left=true)
Definition: block_masking.hpp:623
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
@ MASK_GENERIC
@ MASK_FROM_TOP_LEFT
@ MASK_FROM_BOTTOM_RIGHT
constexpr CK_TILE_HOST_DEVICE auto make_generic_attention_mask_coordinates_from_lr_window(index_t left_size, index_t right_size, index_t y_total, index_t x_total, bool is_top_left=true)
Definition: block_masking.hpp:599
Definition: block_masking.hpp:81
CK_TILE_HOST_DEVICE GenericAttentionMask(index_t y_, index_t x_, index_t y_total_, index_t x_total_)
Definition: block_masking.hpp:94
constexpr CK_TILE_HOST_DEVICE auto IsEdgeTile(index_t i_tile_top, index_t i_tile_left, number< TileHeight >, number< TileWidth >) const
Definition: block_masking.hpp:204
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongX(index_t i_y, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:112
constexpr CK_TILE_HOST_DEVICE auto IsOutOfBound(index_t i_y, index_t i_x) const
Definition: block_masking.hpp:175
CK_TILE_HOST_DEVICE GenericAttentionMask(index_t y_total_, index_t x_total_)
Definition: block_masking.hpp:88
static constexpr const char * name
Definition: block_masking.hpp:86
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongY(index_t i_x, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:149
CK_TILE_HOST_DEVICE GenericAttentionMask(const MaskCoordinates &mask_coord)
Definition: block_masking.hpp:99
static constexpr bool IsMasking
Definition: block_masking.hpp:82
Definition: block_masking.hpp:248
constexpr CK_TILE_HOST_DEVICE auto IsEdgeTile(index_t i_y, index_t i_x, number< TileHeight >, number< TileWidth >) const
Definition: block_masking.hpp:373
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongX(index_t i_y, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:277
constexpr CK_TILE_HOST_DEVICE auto IsOutOfBound(index_t i_y, index_t i_x) const
Definition: block_masking.hpp:350
static constexpr const char * name
Definition: block_masking.hpp:251
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongY(index_t i_x, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:324
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongX(index_t i_y, number< TileHeight > height, number< TileWidth > width, index_t num_splits, index_t i_split) const
Definition: block_masking.hpp:303
CK_TILE_HOST_DEVICE SimplifiedGenericAttentionMask(const MaskCoordinates &mask_coord)
Definition: block_masking.hpp:264
static constexpr bool IsMasking
Definition: block_masking.hpp:249
CK_TILE_HOST_DEVICE SimplifiedGenericAttentionMask(index_t y_total_, index_t x_total_)
Definition: block_masking.hpp:253
CK_TILE_HOST_DEVICE SimplifiedGenericAttentionMask(index_t y_, index_t x_, index_t y_total_, index_t x_total_)
Definition: block_masking.hpp:259
Definition: block_masking.hpp:432
CK_TILE_HOST_DEVICE SimplifiedRatioAttentionMask(index_t y_, index_t x_, index_t y_total_, index_t x_total_, index_t y_real_, index_t y_ratio_, mdiv y_ratio_mdiv_)
Definition: block_masking.hpp:456
static constexpr const char * name
Definition: block_masking.hpp:435
CK_TILE_HOST_DEVICE SimplifiedRatioAttentionMask(index_t y_real_, index_t x_, index_t y_total_, index_t x_total_, mdiv y_ratio_mdiv_)
Definition: block_masking.hpp:443
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongX(index_t i_y, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:478
CK_TILE_HOST_DEVICE SimplifiedRatioAttentionMask(index_t y_total_, index_t x_total_)
Definition: block_masking.hpp:437
constexpr CK_TILE_HOST_DEVICE auto GetTileRangeAlongY(index_t i_x, number< YTile >, number< XTile >) const
Definition: block_masking.hpp:512
constexpr CK_TILE_HOST_DEVICE auto IsEdgeTile(index_t i_y, index_t i_x, number< TileHeight >, number< TileWidth >) const
Definition: block_masking.hpp:560
static constexpr bool IsMasking
Definition: block_masking.hpp:433
constexpr CK_TILE_HOST_DEVICE auto IsOutOfBound(index_t i_y, index_t i_x) const
Definition: block_masking.hpp:538
Definition: integral_constant.hpp:13
Definition: block_masking.hpp:71
Definition: block_masking.hpp:237
Definition: block_masking.hpp:406
Definition: magic_div.hpp:186
CK_TILE_HOST_DEVICE uint32_t div(uint32_t dividend_) const
Definition: magic_div.hpp:212