Mojo struct
MatmulConfig
struct MatmulConfig[a_type: DType, b_type: DType, c_type: DType, transpose_b: Bool = True]
Static configuration of SM90 GPU matmul.
Fields
- block_tile_shape (
IndexList[3]): - mma_shape (
IndexList[3]): - cluster_shape (
IndexList[3]): - num_pipeline_stages (
Int): - num_k_partitions (
Int): - num_consumer (
Int): - partitioned_multicast (
Bool): - k_group_size (
Int):
Implemented traits
AnyType,
Copyable,
Equatable,
Hashable,
ImplicitlyCopyable,
ImplicitlyDestructible,
Movable,
RegisterPassable,
TrivialRegisterPassable,
Writable
Methods
__init__
__init__(block_tile_shape: IndexList[3], mma_shape: IndexList[3], cluster_shape: IndexList[3], num_pipeline_stages: Int, num_k_partitions: Int, num_consumer: Int, partitioned_multicast: Bool, pdl_level: PDLLevel, k_group_size: Int) -> Self
Initialize MatmulConfig with explicit values for all fields.
__init__(m: Int, n: Int, k: Int, num_k_partitions: Int = 1, partitioned_multicast: Bool = False, pdl_level: PDLLevel = PDLLevel.OFF, k_groups: Optional[Int] = None, consumer_groups: Optional[Int] = None, swapAB: Bool = False) -> Self
Initialize MatmulConfig by computing optimal values from M, N, K.
Args:
- m (
Int): The M dimension of the matmul. - n (
Int): The N dimension of the matmul. - k (
Int): The K dimension of the matmul. - num_k_partitions (
Int): Number of K partitions. - partitioned_multicast (
Bool): Whether to use partitioned multicast. - pdl_level (
PDLLevel): PDL level for grid controls. - k_groups (
Optional[Int]): How many pipeline (loads and stores) are grouped together. - consumer_groups (
Optional[Int]): The number of consumer groups. - swapAB (
Bool): Whether to swap A and B.
__eq__
adjust_kgroup_size
static adjust_kgroup_size(mma_m: Int, mma_n: Int, K: Int, BK: Int, num_pipeline_stages: Int) -> Int
Returns:
pdl_level
to_base_config
to_base_config(self) -> MatmulConfig[a_type, b_type, c_type, transpose_b]
Convert to base MatmulConfig from utils_gpu.
Returns:
write_to
write_to(self, mut writer: T)
write_repr_to
write_repr_to(self, mut writer: T)
__hash__
__hash__[H: Hasher](self, mut hasher: H)
Updates hasher with the underlying bytes.
Parameters:
- H (
Hasher): The hasher type.
Args:
- hasher (
H): The hasher instance.
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!