cubecl_ir/

features.rs

use crate::{AddressType, ElemType, OpaqueType, SemanticType, Type};
use alloc::collections::{BTreeMap, BTreeSet};

use crate::EnumSetType;

pub use crate::EnumSet;

/// Features supported by a runtime
#[derive(Debug, Clone, PartialEq, Eq, Default, Hash)]
pub struct Features {
    /// Plane features supported by this runtime.
    pub plane: EnumSet<Plane>,
    /// Clustered launches and intra-cluster operations like cluster shared memory
    pub cube_cluster: bool,
    /// Enables changing the type of containers during kernel execution.
    pub memory_reinterpret: bool,
    /// Enables explicit alignment. If false, alignment still compiles, but isn't actually applied.
    pub alignment: bool,

    /// Type support
    pub types: Types,
    /// Matrix multiplication features
    pub matmul: MatmulFeatures,

    /// Whether `copy_async` is supported
    pub copy_async: bool,
    /// Whether a [`SyncScope::Device`](crate::dialect::synchronization::SyncScope::Device)
    /// synchronization is a release and an acquire at device scope, so that one cube's writes to
    /// storage are visible to another that synchronizes after it. Without it the same
    /// synchronization is a cube barrier and nothing more, which is all WebGPU's memory model
    /// promises, so a kernel whose cubes hand each other data must ask before it runs.
    pub device_memory_scope: bool,
    /// Tensor Memory Accelerator supported features
    pub tma: EnumSet<Tma>,
    /// Whether vectors can be read from / stored to addresses not aligned
    /// with the `vector_size`
    pub unaligned_io: bool,
}

/// Type support for a device
#[derive(Debug, Clone, PartialEq, Eq, Default, Hash)]
pub struct Types {
    /// Valid address types
    pub address: BTreeSet<AddressType>,
    /// Types supported by this runtime, and which usages they support.
    pub elem: BTreeMap<ElemType, EnumSet<TypeUsage>>,
    /// Complex-specific capability families supported by this runtime.
    pub complex: BTreeMap<ElemType, EnumSet<ComplexUsage>>,
    /// Semantic constructs supported by this runtime.
    pub semantic: BTreeSet<SemanticType>,
    /// Opaque types supported by this runtime.
    pub opaque: BTreeSet<OpaqueType>,
    /// Supported vector types for atomic ops, only specific vectorizations for specific types are
    /// supported here. Not all vector types are supported as scalars, i.e. Vulkan on Nvidia only
    /// supports vectorized `f16`, not scalar. Only use the exact vectorizations registered here.
    /// These may not be supported everywhere - in practice, f32 vectors are only supported in global
    /// memory.
    pub atomic: BTreeMap<Type, EnumSet<AtomicUsage>>,
}

/// Matrix multiplication-related features
#[derive(Debug, Clone, PartialEq, Eq, Default, Hash)]
pub struct MatmulFeatures {
    /// The cmma feature enables cooperative matrix-multiply and accumulate operations.
    pub cmma: BTreeSet<MmaConfig>,
    /// Cube MMA is like cmma but at the cube level, rather than the plane level.
    /// Loading may be staged in shared memory by the driver on Vulkan - check
    /// [`cube_mma_reserved_shared_memory`](crate::HardwareProperties::cube_mma_reserved_shared_memory)
    /// to take this into account when generating a matmul config.
    pub cube_mma: BTreeSet<CubeMmaConfig>,
    /// The manual MMA feature enables cooperative matrix-multiply with manually managed data
    /// movement
    pub mma: BTreeSet<MmaConfig>,
    /// Scaled MMA allows combining matrix multiplication with unscaling quantized values into a single
    /// instruction. Scales must fit a specific layout and block size.
    pub scaled_mma: BTreeSet<ScaledMmaConfig>,
    /// Types supported for ldmatrix, if any
    pub ldmatrix: BTreeSet<ElemType>,
    /// Types supported by stmatrix, if any
    pub stmatrix: BTreeSet<ElemType>,
    /// Whether tensor addressing is supported for CMMA load/store
    pub cmma_tensor_addressing: bool,
}

/// Operations allowed for this type. CMMA is defined separately.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord, EnumSetType)]
pub enum TypeUsage {
    /// Conversion to/from the type. All types should support this.
    Conversion,
    /// All math/logic instructions except dot product
    Arithmetic,
    /// Dot product, mainly for BF16 on Intel
    DotProduct,
    /// Whether this type can be stored in a buffer
    Buffer,
}

/// Complex capability families allowed for a complex element type.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord, EnumSetType)]
pub enum ComplexUsage {
    /// Arithmetic, negation, conjugation, and real/imaginary extraction.
    Core,
    /// Equality and inequality comparisons.
    Compare,
    /// Higher-level complex math functions.
    Math,
}

impl TypeUsage {
    pub fn all() -> EnumSet<Self> {
        EnumSet::all()
    }

    pub fn no_store() -> EnumSet<Self> {
        TypeUsage::Conversion | TypeUsage::Arithmetic
    }

    pub fn maybe_store(storable: bool) -> EnumSet<Self> {
        if storable {
            EnumSet::all()
        } else {
            Self::no_store()
        }
    }
}

/// Atomic operations allowed for this type.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord, EnumSetType)]
pub enum AtomicUsage {
    /// Atomic loads and stores
    LoadStore,
    /// Atomic exchange
    Exchange,
    /// Atomic add/sub
    Add,
    /// Atomic min/max
    MinMax,
    /// Atomic bitwise and/or/xor
    Bitwise,
    /// Atomic compare-and-exchange
    CompareExchange,
}

impl AtomicUsage {
    pub fn all() -> EnumSet<Self> {
        EnumSet::all()
    }
}

/// Supported plane features
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord, EnumSetType)]
pub enum Plane {
    /// Basic plane-wide operations
    Ops,
    /// Plane-wide sync
    Sync,
    /// Allows using plane operations with divergent control flow.
    NonUniformControlFlow,
}

/// Shape and element types of a valid MMA configuration
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct MmaConfig {
    /// Element of the A matrix
    pub a_type: ElemType,
    /// Element of the B matrix
    pub b_type: ElemType,
    /// Element of the C/D matrices
    pub cd_type: ElemType,
    /// The size of the matrix on the `m` dimension
    pub m: u32,
    /// The size of the matrix on the `n` dimension
    pub n: u32,
    /// The size of the matrix on the `k` dimension
    pub k: u32,
}

/// Shape and element types of a valid flexible MMA configuration
/// Only Vulkan for now, but this should also be usable for wgmma/xmma on datacenter CUDA.
/// Actual matrix size must be multiple of `granularity` and `<= max`.
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct CubeMmaConfig {
    /// Element of the A matrix
    pub a_type: ElemType,
    /// Element of the B matrix
    pub b_type: ElemType,
    /// Element of the C/D matrices
    pub cd_type: ElemType,
    /// The granularity of the matrix on the `m` dimension
    pub m_granularity: u32,
    /// The maximum value for `m`
    pub m_max: u32,
    /// The size of the matrix on the `n` dimension
    pub n_granularity: u32,
    /// The maximum value for `n`
    pub n_max: u32,
    /// The size of the matrix on the `k` dimension
    pub k_granularity: u32,
    /// The maximum value for `k`
    pub k_max: u32,
    /// The number of units that must be in the cube for this configuration to be valid.
    /// `None` means it's always valid (but might still have an optimal value).
    pub units_per_block: Option<u32>,
}

/// Shape and element types of a valid block-scaled MMA configuration
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct ScaledMmaConfig {
    /// Element of the A matrix
    pub a_type: ElemType,
    /// Element of the B matrix
    pub b_type: ElemType,
    /// Element of the C/D matrices
    pub cd_type: ElemType,
    /// Element of the blocks scales
    pub scales_type: ElemType,
    /// The size of the matrix on the `m` dimension
    pub m: u32,
    /// The size of the matrix on the `n` dimension
    pub n: u32,
    /// The size of the matrix on the `k` dimension
    pub k: u32,
    /// Number of scales per tile row/col.
    /// A scale factor of 2 means `m x 2` scales for A and `2 x n` for B (in CUDA)
    /// Scales blocks must be organized along the natural `vector_layout` of the operation
    pub scales_factor: u32,
}

/// Atomic features that may be supported by a ``Runtime``.
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, EnumSetType)]
pub enum Tma {
    /// Base feature set for tensor memory accelerator features. Includes tiling and im2col
    Base,
    /// im2colWide encoding for tensor map.
    Im2colWide,
    /// Different atomicities for 128-byte swizzle, i.e. 128-byte with 32-byte atomicity.
    SwizzleAtomicity,
}

impl Features {
    /// Get the usages for a type
    pub fn type_usage(&self, ty: ElemType) -> EnumSet<TypeUsage> {
        self.types
            .elem
            .get(&ty)
            .cloned()
            .unwrap_or_else(EnumSet::empty)
    }

    /// Get the complex capability families for a type.
    pub fn complex_usage(&self, ty: ElemType) -> EnumSet<ComplexUsage> {
        self.types
            .complex
            .get(&ty)
            .cloned()
            .unwrap_or_else(EnumSet::empty)
    }

    /// Whether a complex type supports the requested capability family.
    pub fn supports_complex_usage(&self, ty: ElemType, usage: ComplexUsage) -> bool {
        self.complex_usage(ty).contains(usage)
    }

    /// Get the usages for an atomic type
    pub fn atomic_type_usage(&self, ty: Type) -> EnumSet<AtomicUsage> {
        self.types
            .atomic
            .get(&ty)
            .cloned()
            .unwrap_or_else(EnumSet::empty)
    }

    /// Whether the type is supported in any way
    pub fn supports_type(&self, ty: impl Into<Type>) -> bool {
        match ty.into() {
            Type::Semantic(semantic_type) => self.types.semantic.contains(&semantic_type),
            Type::Opaque(opaque_type) => self.types.opaque.contains(&opaque_type),
            ty => self.types.elem.contains_key(&ty.elem_type()),
        }
    }

    /// Whether the address type is supported in any way
    pub fn supports_address(&self, ty: impl Into<AddressType>) -> bool {
        self.types.address.contains(&ty.into())
    }
}