docc 0.8.3 (0d44999b)
C++ API Reference
Loading...
Searching...
No Matches
cuda_mma_fragment_dispatcher.h
Go to the documentation of this file.
1#pragma once
2
7
8namespace sdfg::gpu::cuda {
9
11
14public:
16
17 void emit_fill(codegen::CodegenOutput& out, const std::string& acc_frag, types::PrimitiveType fill_type) const override;
18
21 const std::string& frag_d,
22 const std::string& frag_a,
23 const std::string& frag_b,
24 const std::string& frag_c
25 ) const override;
26
29 const std::string& frag,
30 const std::string& base,
31 const GpuMmaFromMemoryLayout& layout,
32 MmaFragmentType frag_type,
33 types::PrimitiveType element_type,
34 MmaFragmentLayout fragment_layout
35 ) const override;
36
39 const std::string& frag,
40 const std::string& base,
41 const GpuMmaFromMemoryLayout& layout,
42 MmaFragmentType frag_type,
43 types::PrimitiveType element_type
44 ) const override;
45
48 const std::string& main_frag,
49 const std::vector<std::string>& additional_frags,
50 const EltwiseCompute& compute
51 ) const override;
52};
53
55public:
57 static CudaWmmaLibDependency inst;
58 return &inst;
59 }
60
61 std::string_view name() const override {
62 return "nvcuda-wmma";
63 }
64 void enumerate_includes(std::vector<std::string>& out_list) const override {
65 out_list.push_back("mma.h");
66 }
67 std::vector<std::string_view>& globally_unique_ids() const override {
68 static std::vector<std::string_view> ids{"nvcuda-wmma"};
69 return ids;
70 }
71};
72
73} // namespace sdfg::gpu::cuda
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