For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /get-started.md).

Mojo struct

TuningConfigSM100

struct TuningConfigSM100

Holds SM100 matmul kernel launch parameters for a range of MxNxK shapes.

Stores the MMA shape, block tile shape, cluster shape, swizzle and rasterization settings, pipeline stage counts, and split-K factor that select an optimized kernel configuration for matmuls whose M dimension falls in the half-open interval [M, M_end).

Fields

• ​M (Int):
• ​M_end (Int):
• ​N (Int):
• ​K (Int):
• ​mma_shape (IndexList[Int(3)]):
• ​block_tile_shape (IndexList[Int(3)]):
• ​cluster_shape (IndexList[Int(3)]):
• ​block_swizzle_size (Int):
• ​rasterize_order (RasterOrder):
• ​cta_group (Int):
• ​swapAB (Bool):
• ​k_group_size (Int):
• ​num_accum_pipeline_stages (Int):
• ​num_clc_pipeline_stages (Int):
• ​num_split_k (Int):
• ​num_pipeline_stages (Int):
• ​is_small_bn (Bool):
• ​batch_size (Int):

Implemented traits

AnyType, Copyable, Deinitable, ImplicitlyCopyable, Movable, RegisterPassable, TrivialRegisterPassable, TuningConfig, Writable

Methods

__init__

def __init__(M: Int, N: Int, K: Int, mma_shape: IndexList[Int(3)], block_tile_shape: IndexList[Int(3)], cluster_shape: IndexList[Int(3)], block_swizzle_size: Int, rasterize_order: RasterOrder, cta_group: Int = Int(2), swapAB: Bool = False, k_group_size: Int = Int(1), num_accum_pipeline_stages: Int = Int(2), num_clc_pipeline_stages: Int = Int(2), num_split_k: Int = Int(1), num_pipeline_stages: Int = Int(0), is_small_bn: Bool = False, batch_size: Int = Int(1)) -> Self

def __init__(M: Int, M_end: Int, N: Int, K: Int, mma_shape: IndexList[Int(3)], cta_group: Int, cluster_shape: IndexList[Int(3)], block_swizzle_size: Int, rasterize_order: RasterOrder, swapAB: Bool = False, k_group_size: Int = Int(1), num_accum_pipeline_stages: Int = Int(2), num_clc_pipeline_stages: Int = Int(2), num_split_k: Int = Int(1), num_pipeline_stages: Int = Int(0), is_small_bn: Bool = False, batch_size: Int = Int(1)) -> Self

write_to

def write_to(self, mut writer: T)

Writes the tuning config as a string.

Args:

• ​writer (T): The writer to write to.