21 const std::string& frag_d,
22 const std::string& frag_a,
23 const std::string& frag_b,
24 const std::string& frag_c
29 const std::string& frag,
30 const std::string& base,
39 const std::string& frag,
40 const std::string& base,
48 const std::string& main_frag,
49 const std::vector<std::string>& additional_frags,
61 std::string_view
name()
const override {
65 out_list.push_back(
"mma.h");
68 static std::vector<std::string_view> ids{
"nvcuda-wmma"};
Dispatchers for dataflow blocks and library nodes.
Definition code_snippet_factory.h:51
Hardware-dependent emit primitives for the composable MMA library nodes.
Definition gpu_mma_dispatcher.h:23
std::function< void( codegen::CodegenOutput &out, const std::string &main_frag_elem, const std::string &idx, const std::vector< std::string > &other_frag_elems)> EltwiseCompute
Per-element callback used by emit_eltwise: writes one output element from its inputs.
Definition gpu_mma_dispatcher.h:26
Classic rocwmma fragment-based emit strategy (load/store/mma_sync/fill_fragment).
Definition cuda_mma_fragment_dispatcher.h:13
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_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_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_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.
void emit_needed_declarations(codegen::CodegenOutput &out) const override
Target library includes / helper declarations needed before the first emit.
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.
Definition cuda_mma_fragment_dispatcher.h:54
void enumerate_includes(std::vector< std::string > &out_list) const override
Definition cuda_mma_fragment_dispatcher.h:64
static const CudaWmmaLibDependency * instance()
Definition cuda_mma_fragment_dispatcher.h:56
std::vector< std::string_view > & globally_unique_ids() const override
Definition cuda_mma_fragment_dispatcher.h:67
std::string_view name() const override
Definition cuda_mma_fragment_dispatcher.h:61
Library node definitions for calling functions beyond simple instructions.
StringEnum ImplementationType
Implementation type identifier for library nodes.
Definition library_node.h:75
Definition cuda_arch.h:11
data_flow::ImplementationType ImplementationType_CUDA_MMA_FRAGMENT("CUDA_MMA")
MmaFragmentLayout
Definition gpu_mma_fragment.h:78
MmaFragmentType
Definition gpu_mma_fragment.h:14
PrimitiveType
Enumeration of all primitive data types supported by sdfglib.
Definition type.h:35
Definition block_dispatcher.h:108
Describes how a single MMA fragment maps to memory relative to a base pointer.
Definition gpu_mma_fragment.h:163