docc 0.8.3 (0d44999b)
C++ API Reference
Loading...
Searching...
No Matches
rocm_arch.h
Go to the documentation of this file.
1#pragma once
2
3#include <string>
4#include <vector>
5
9#include "sdfg/types/array.h"
10
11namespace sdfg::gpu::rocm {
12
13enum class RocmMmaVariant {
15 Bare = 2,
16};
17
19 const bool base_support;
20 const bool f32_support;
21 const uint16_t threads_per_mma_block;
24 const uint16_t int8_block_k;
25
30
33 return {16, 16, int8_block_k};
34 }
35
36protected:
37 static constexpr MmaBlockSize DEFAULT_BLOCK_SIZE = {16, 16, 16};
38 static constexpr MmaBlockSize CDNA_BLOCK_SIZE = {32, 32, 8};
39
40public:
41 static constexpr const char* MMA_FRAGMENT_STORAGE_TYPE = "ROCM_MMA";
42 static constexpr const char* MMA_REGS_STORAGE_TYPE = "ROCM_MMA_REGS";
43
44 bool supported_types(types::PrimitiveType input_type, types::PrimitiveType output_type) const override;
45
47 types::PrimitiveType input_type, types::PrimitiveType output_type, types::PrimitiveType desired_acc_type
48 ) const override;
49
51 const MmaBlockSize& block_size, types::PrimitiveType input_type, types::PrimitiveType acc_type
52 ) const override;
53
54 std::optional<GpuMmaTiling> try_get_mma_tiling(
55 const MmaBlockSize& block_size,
56 const BlockCounts* wave_tile_hint,
57 const symbolic::MultiExpression& res_shape,
58 types::PrimitiveType acc_type,
59 bool ceil_counts = false
60 ) const;
61
62 std::optional<GpuMmaTiling> get_mma_tiling(
63 const symbolic::MultiExpression& res_shape,
64 types::PrimitiveType input_type,
65 types::PrimitiveType acc_type,
66 const MmaBlockSize* block_size_hint,
67 const BlockCounts* wave_tile_hint
68 ) const override;
69
70 std::optional<GpuMmaTiling> get_ragged_mma_tiling(
71 const symbolic::MultiExpression& padded_shape,
72 types::PrimitiveType input_type,
73 types::PrimitiveType output_type,
74 types::PrimitiveType acc_type,
75 const MmaBlockSize* block_size_hint,
76 const BlockCounts* wave_tile_hint
77 ) const override;
78
80 types::StorageType& storage_type, const GpuMmaTiling& size, MmaFragmentType type, MmaFragmentLayout layout
81 ) const override;
82
83 static bool is_mma_type(const types::StorageType& storage);
84
86
88
90 std::ostream& os, const types::StorageType& storage_type, types::PrimitiveType element_type
91 ) const;
92
94 std::ostream& os,
95 MmaFragmentType type,
96 const MmaBlockSize& dims,
97 MmaFragmentLayout layout,
98 types::PrimitiveType scalar_type,
99 std::optional<std::pair<int, int>> coop_dims
100 );
101
103 std::ostream& os,
104 MmaFragmentType type,
105 const MmaBlockSize& dims,
106 MmaFragmentLayout layout,
107 types::PrimitiveType scalar_type,
108 std::optional<std::pair<int, int>> coop_dims
109 );
110
111 std::vector<MmaBlockSize>
113};
114
115
116class RocmArch : public GpuArch {
117 int per_cu_threads_;
118 std::string rocm_name_;
119 RocmMmaSupport mma_support_;
120
121public:
123 const std::string& name,
124 int per_cu_threads,
125 bool mma_base_support,
126 bool mma_f32_support,
127 uint16_t mma_int8_block_k = 0
128 )
129 : GpuArch(), rocm_name_(name), per_cu_threads_(per_cu_threads),
130 mma_support_(mma_base_support ? 16 : 0, mma_f32_support, per_cu_threads, mma_int8_block_k) {
131 }
132
133 std::string unique_id() const override;
134 std::string name() const override;
135
136 int per_cu_threads() const override {
137 return per_cu_threads_;
138 }
139
140 const RocmMmaSupport* mma_support() const override {
141 if (mma_support_.base_support) {
142 return &mma_support_;
143 } else {
144 return nullptr;
145 }
146 }
147
149};
150
152
155
161 std::string gfx_name;
163 std::vector<std::string> device_names;
166};
167
177const std::vector<RocmDeviceInfo>& query_rocm_devices();
178
184
191const RocmArch* rocm_arch_parse(const std::string& name);
192
193} // namespace sdfg::gpu::rocm
Array type representation for c-style arrays.
Dispatchers for dataflow blocks and library nodes.
GpuArch()
Definition gpu_arch.h:18
Definition rocm_arch.h:116
const RocmMmaSupport * mma_support() const override
Definition rocm_arch.h:140
int per_cu_threads() const override
Definition rocm_arch.h:136
structured_control_flow::ScheduleType create_schedule_type() const override
RocmArch(const std::string &name, int per_cu_threads, bool mma_base_support, bool mma_f32_support, uint16_t mma_int8_block_k=0)
Definition rocm_arch.h:122
std::string name() const override
std::string unique_id() const override
Represents a schedule type for Map nodes.
Definition structured_loop.h:31
Represents the storage location and management characteristics of data.
Definition type.h:64
std::vector< Expression > MultiExpression
Vector of symbolic expressions for multi-dimensional operations.
Definition symbolic.h:89
StringEnum ImplementationType
Implementation type identifier for library nodes.
Definition library_node.h:75
Definition rocm_arch.h:11
const RocmArch * rocm_arch_from_available_hardware()
RocmArch ROCM_ARCH_GFX90A
RocmArch ROCM_ARCH_GFX942
RocmArch ROCM_ARCH_GFX1201
const RocmArch * rocm_arch_from_schedule_type(const structured_control_flow::ScheduleType &schedule)
Wrapper, that allows us to store this inside the ScheduleType in the future.
RocmMmaVariant
Definition rocm_arch.h:13
@ Bare
raw MFMA builtin path (RocmMmaBareStrategy)
Definition rocm_arch.h:15
@ Fragment
rocwmma fragment path (RocmMmaFragmentStrategy)
Definition rocm_arch.h:14
const std::vector< RocmDeviceInfo > & query_rocm_devices()
Queries the ROCm GPUs available on this machine via rocminfo.
const RocmArch * rocm_arch_for_device(const RocmDeviceInfo &device)
Resolves a queried device to a RocmArch.
const RocmArch * rocm_arch_from_env()
Use the DOCC_ROCM_ARCH env var or if unset, query the available cards.
const RocmArch * rocm_arch_parse(const std::string &name)
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
Grouping of MMA blocks.
Definition gpu_mma_fragment.h:35
Definition gpu_mma_fragment.h:44
Definition gpu_mma_fragment.h:16
A ROCm/HIP GPU device as reported by rocminfo: its LLVM/ROCm ISA (gfx) name together with the marketi...
Definition rocm_arch.h:159
std::string gfx_name
LLVM/ROCm ISA target name, e.g. "gfx1201", "gfx942".
Definition rocm_arch.h:161
int wavefront_size
Wavefront (wave) size in threads, e.g. 32 (RDNA) or 64 (CDNA/GCN); 0 if unknown.
Definition rocm_arch.h:165
std::vector< std::string > device_names
Distinct marketing device names sharing this gfx target, useful for logging.
Definition rocm_arch.h:163
Definition rocm_arch.h:18
static constexpr const char * MMA_FRAGMENT_STORAGE_TYPE
Definition rocm_arch.h:41
void emit_block_frag_type(std::ostream &os, const types::StorageType &storage_type, types::PrimitiveType element_type) const
static void emit_block_frag_type(std::ostream &os, MmaFragmentType type, const MmaBlockSize &dims, MmaFragmentLayout layout, types::PrimitiveType scalar_type, std::optional< std::pair< int, int > > coop_dims)
bool supported_types(types::PrimitiveType input_type, types::PrimitiveType output_type) const override
bool is_valid_block_size(const MmaBlockSize &block_size, types::PrimitiveType input_type, types::PrimitiveType acc_type) const override
RocmMmaSupport(bool base_support, bool f32_support, uint16_t threads, uint16_t int8_block_k=0)
Definition rocm_arch.h:26
static constexpr MmaBlockSize CDNA_BLOCK_SIZE
Definition rocm_arch.h:38
std::optional< GpuMmaTiling > get_ragged_mma_tiling(const symbolic::MultiExpression &padded_shape, types::PrimitiveType input_type, types::PrimitiveType output_type, types::PrimitiveType acc_type, const MmaBlockSize *block_size_hint, const BlockCounts *wave_tile_hint) const override
Tiling for a ragged problem: padded_shape is a constant upper bound of each extent that need not be a...
types::PrimitiveType get_accumulator_type(types::PrimitiveType input_type, types::PrimitiveType output_type, types::PrimitiveType desired_acc_type) const override
static constexpr MmaBlockSize DEFAULT_BLOCK_SIZE
Definition rocm_arch.h:37
static void emit_bare_fragment_decl_type(std::ostream &os, MmaFragmentType type, const MmaBlockSize &dims, MmaFragmentLayout layout, types::PrimitiveType scalar_type, std::optional< std::pair< int, int > > coop_dims)
const bool base_support
Definition rocm_arch.h:19
static RocmMmaVariant select_variant(const GpuMmaTiling &tiling)
static constexpr const char * MMA_REGS_STORAGE_TYPE
Definition rocm_arch.h:42
std::optional< GpuMmaTiling > get_mma_tiling(const symbolic::MultiExpression &res_shape, types::PrimitiveType input_type, types::PrimitiveType acc_type, const MmaBlockSize *block_size_hint, const BlockCounts *wave_tile_hint) const override
const bool f32_support
Definition rocm_arch.h:20
data_flow::ImplementationType get_mma_impl_type(const GpuMmaTiling &tiling) const override
static bool is_mma_type(const types::StorageType &storage)
void set_mma_fragment_storage_type(types::StorageType &storage_type, const GpuMmaTiling &size, MmaFragmentType type, MmaFragmentLayout layout) const override
std::vector< MmaBlockSize > get_supported_block_sizes(types::PrimitiveType input_type, types::PrimitiveType acc_type) const override
std::optional< GpuMmaTiling > try_get_mma_tiling(const MmaBlockSize &block_size, const BlockCounts *wave_tile_hint, const symbolic::MultiExpression &res_shape, types::PrimitiveType acc_type, bool ceil_counts=false) const
MmaBlockSize int8_block_size() const
The block size Int8 matmuls use by default, at the full Int8 rate.
Definition rocm_arch.h:32
const uint16_t threads_per_mma_block
Definition rocm_arch.h:21
const uint16_t int8_block_k
K of the native Int8 x Int8 -> Int32 instruction of 16x16 blocks (16 on CDNA2 and RDNA,...
Definition rocm_arch.h:24