/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_v3_pipeline.hpp File Reference

• develop
• include
• ck_tile
• ops
• fmha
• pipeline

Classes | Namespaces | Macros | Functions

block_fmha_fwd_v3_pipeline.hpp File Reference

#include "ck_tile/core.hpp"
#include "ck_tile/ops/fmha/pipeline/block_fmha_fwd_v3_pipeline_default_policy.hpp"
#include "ck_tile/ops/reduce/block/block_reduce.hpp"

Go to the source code of this file.

Classes
struct	ck_tile::CoreLoopScheduler< PipelineProblem, true >
struct	ck_tile::CoreLoopScheduler< PipelineProblem, false >
struct	ck_tile::BlockFmhaFwdV3Pipeline< Problem_, Policy_ >

Namespaces
	ck_tile
	ck_tile::detail

Macros
#define	ENABLE_ASM_MARKER   1
#define	ASM_MARKER(marker)
#define	ADD_SBARRIER_FOR_PHASE0   1
#define	CK_TILE_DISABLE_PACKED_FP32   0
#define	WARP_ID   0
#define	LANE_ID   0
#define	ENABLE_DEBUG_STMTS   1
#define	DEBUG_STMTS    if(get_block_1d_id() == 0 && get_warp_id() == WARP_ID && get_lane_id() == LANE_ID)

Functions
CK_TILE_DEVICE float	ck_tile::detail::fma_impl_vsv (float a, float b, float c)
CK_TILE_DEVICE float	ck_tile::detail::add_impl_vv (float lhs, float rhs)
CK_TILE_DEVICE float	ck_tile::detail::mul_impl_vv (float lhs, float rhs)
CK_TILE_DEVICE fp16x2_t	ck_tile::detail::cvt_pk_fp16_f32 (float a, float b)
CK_TILE_DEVICE bf16x2_t	ck_tile::detail::cvt_pk_bf16_f32 (float a, float b)
CK_TILE_DEVICE fp32x2_t	ck_tile::detail::pk_mul_f32 (fp32x2_t lhs, fp32x2_t rhs)

Macro Definition Documentation

◆ ADD_SBARRIER_FOR_PHASE0

#define ADD_SBARRIER_FOR_PHASE0   1

◆ ASM_MARKER

#define ASM_MARKER

(

marker

)

Value:

__builtin_amdgcn_sched_barrier(0); \

asm volatile("; [POYENC] " #marker); \

__builtin_amdgcn_sched_barrier(0);

◆ CK_TILE_DISABLE_PACKED_FP32

#define CK_TILE_DISABLE_PACKED_FP32   0

◆ DEBUG_STMTS

#define DEBUG_STMTS    if(get_block_1d_id() == 0 && get_warp_id() == WARP_ID && get_lane_id() == LANE_ID)

◆ ENABLE_ASM_MARKER

#define ENABLE_ASM_MARKER   1

◆ ENABLE_DEBUG_STMTS

#define ENABLE_DEBUG_STMTS   1

◆ LANE_ID

#define LANE_ID   0

◆ WARP_ID

#define WARP_ID   0