59 bool ceil_counts =
false
99 std::optional<std::pair<int, int>> coop_dims
108 std::optional<std::pair<int, int>> coop_dims
111 std::vector<MmaBlockSize>
118 std::string rocm_name_;
123 const std::string&
name,
125 bool mma_base_support,
126 bool mma_f32_support,
127 uint16_t mma_int8_block_k = 0
130 mma_support_(mma_base_support ? 16 : 0, mma_f32_support,
per_cu_threads, mma_int8_block_k) {
134 std::string
name()
const override;
137 return per_cu_threads_;
141 if (mma_support_.base_support) {
142 return &mma_support_;
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