/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

/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#

Composable Kernel: /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

#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 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