Struct WmmaTile 

pub struct WmmaTile {
    pub bm: usize,
    pub bn: usize,
    pub bk: usize,
    pub warps_m: usize,
    pub warps_n: usize,
}

A pipelined-WMMA configuration: block tile BM×BN, K-stage BK, warp grid WARPS_M×WARPS_N.

Fields

bm: usize

bn: usize

bk: usize

warps_m: usize

warps_n: usize

Implementations

impl WmmaTile

pub fn warps(&self) -> usize

pub fn threads(&self) -> usize

pub fn wm(&self) -> usize

Warp-tile height/width (rows/cols of C each warp owns).

pub fn wn(&self) -> usize

pub fn smem_bytes(&self) -> usize

int8 shared-memory footprint (single-buffered): A stage + B stage.

pub fn acc_regs(&self) -> usize

Accumulator registers per thread ≈ (frags per warp) × 8 int32 regs per 16×16 acc fragment.

pub fn is_legal(&self, gpu: &GpuLimits) -> bool

Legal iff everything divides on 16-boundaries, the warp grid tiles the block, and it fits Ampere.

impl WmmaTile

pub fn cpasync_legal(&self, gpu: &GpuLimits) -> bool

cp.async needs 16-byte contiguous chunks: A rows are BK wide, B rows BN wide.

Trait Implementations

impl Clone for WmmaTile

fn clone(&self) -> WmmaTile

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for WmmaTile

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl PartialEq for WmmaTile

fn eq(&self, other: &WmmaTile) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl Copy for WmmaTile

impl Eq for WmmaTile

impl StructuralPartialEq for WmmaTile