sdfg::gpu::rocm::RocmMmaBareStrategy Class Reference

Raw v_mfma_f32_32x32x8f16 emit strategy (see RocmMfma32Support). More...

#include <sdfg/targets/rocm/rocm_mma_bare_dispatcher.h>

Inheritance diagram for sdfg::gpu::rocm::RocmMmaBareStrategy:

Public Member Functions
void	emit_fill (codegen::CodegenOutput &out, const std::string &acc_frag, types::PrimitiveType fill_type) const override
	Zero-initialize the accumulator fragment behind acc_frag.
void	emit_mma (codegen::CodegenOutput &out, const std::string &frag_d, const std::string &frag_a, const std::string &frag_b, const std::string &frag_c) const override
	frag_d = frag_a * frag_b + frag_c (single MMA step; all operands are fragments).
void	emit_fragment_load (codegen::CodegenOutput &out, const std::string &frag, const std::string &base, const GpuMmaFromMemoryLayout &layout, MmaFragmentType frag_type, types::PrimitiveType element_type, MmaFragmentLayout fragment_layout) const override
	Load a tile from memory (base + layout) into the fragment behind frag.
void	emit_fragment_store (codegen::CodegenOutput &out, const std::string &frag, const std::string &base, const GpuMmaFromMemoryLayout &layout, MmaFragmentType frag_type, types::PrimitiveType element_type) const override
	Store the fragment behind frag to memory (base + layout).
void	emit_eltwise (codegen::CodegenOutput &out, const std::string &main_frag, const std::vector< std::string > &additional_frags, const EltwiseCompute &compute) const override
	Element-wise loop over main_frag, invoking compute for each element.
std::string	fragment_element (const std::string &frag, const std::string &idx) const override
	Expression for element idx of this lane's share of frag (needed for clipped transfers).
std::string	fragment_element_count (const std::string &frag) const override
	Number of elements of this lane's share of frag (needed for clipped transfers).
std::optional< std::pair< std::string, std::string > >	fragment_element_coords (MmaFragmentType frag_type, const MmaBlockSize &block_size, const std::string &idx, std::string &condition) const override
	The raw MFMA register layouts documented above.
Public Member Functions inherited from sdfg::gpu::GpuMmaDispatcherStrategy
virtual	~GpuMmaDispatcherStrategy ()=default
virtual void	emit_needed_declarations (codegen::CodegenOutput &out) const
	Target library includes / helper declarations needed before the first emit.

Additional Inherited Members
Public Types inherited from sdfg::gpu::GpuMmaDispatcherStrategy
using	EltwiseCompute
	Per-element callback used by emit_eltwise: writes one output element from its inputs.

Detailed Description

Raw v_mfma_f32_32x32x8f16 emit strategy (see RocmMfma32Support).

Fragments are per-lane register vectors; lane l = threadIdx.x % 64 holds: A (32x8): row l%32, k 4*(l/32) .. +3 B (8x32): col l%32, k 4*(l/32) .. +3 C (32x32): element e at row 8*(e/4) + 4*(l/32) + e%4, col l%32

Member Function Documentation

emit_eltwise()

void sdfg::gpu::rocm::RocmMmaBareStrategy::emit_eltwise ( codegen::CodegenOutput & out, const std::string & main_frag, const std::vector< std::string > & additional_frags, const EltwiseCompute & compute ) const

overridevirtual

Element-wise loop over main_frag, invoking compute for each element.

Implements sdfg::gpu::GpuMmaDispatcherStrategy.

emit_fill()

void sdfg::gpu::rocm::RocmMmaBareStrategy::emit_fill ( codegen::CodegenOutput & out, const std::string & acc_frag, types::PrimitiveType fill_type ) const

overridevirtual

Zero-initialize the accumulator fragment behind acc_frag.

Implements sdfg::gpu::GpuMmaDispatcherStrategy.

emit_fragment_load()

void sdfg::gpu::rocm::RocmMmaBareStrategy::emit_fragment_load ( codegen::CodegenOutput & out, const std::string & frag, const std::string & base, const GpuMmaFromMemoryLayout & layout, MmaFragmentType frag_type, types::PrimitiveType element_type, MmaFragmentLayout fragment_layout ) const

overridevirtual

Load a tile from memory (base + layout) into the fragment behind frag.

Implements sdfg::gpu::GpuMmaDispatcherStrategy.

emit_fragment_store()

void sdfg::gpu::rocm::RocmMmaBareStrategy::emit_fragment_store ( codegen::CodegenOutput & out, const std::string & frag, const std::string & base, const GpuMmaFromMemoryLayout & layout, MmaFragmentType frag_type, types::PrimitiveType element_type ) const

overridevirtual

Store the fragment behind frag to memory (base + layout).

Implements sdfg::gpu::GpuMmaDispatcherStrategy.

emit_mma()

void sdfg::gpu::rocm::RocmMmaBareStrategy::emit_mma ( codegen::CodegenOutput & out, const std::string & frag_d, const std::string & frag_a, const std::string & frag_b, const std::string & frag_c ) const

overridevirtual

frag_d = frag_a * frag_b + frag_c (single MMA step; all operands are fragments).

Implements sdfg::gpu::GpuMmaDispatcherStrategy.

fragment_element()

std::string sdfg::gpu::rocm::RocmMmaBareStrategy::fragment_element ( const std::string & frag, const std::string & idx ) const

inlineoverridevirtual

Expression for element idx of this lane's share of frag (needed for clipped transfers).

Reimplemented from sdfg::gpu::GpuMmaDispatcherStrategy.

fragment_element_coords()

std::optional< std::pair< std::string, std::string > > sdfg::gpu::rocm::RocmMmaBareStrategy::fragment_element_coords ( MmaFragmentType frag_type, const MmaBlockSize & block_size, const std::string & idx, std::string & condition ) const

inlineoverridevirtual

The raw MFMA register layouts documented above.

Reimplemented from sdfg::gpu::GpuMmaDispatcherStrategy.

fragment_element_count()

std::string sdfg::gpu::rocm::RocmMmaBareStrategy::fragment_element_count ( const std::string & frag ) const

inlineoverridevirtual

Number of elements of this lane's share of frag (needed for clipped transfers).

Reimplemented from sdfg::gpu::GpuMmaDispatcherStrategy.

---

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

• opt/include/sdfg/targets/rocm/rocm_mma_bare_dispatcher.h