IMPORTANT: To view this page as Markdown, append `.md` to the URL (e.g. /max/get-started.md). For the complete documentation index, see llms.txt.
Skip to main content
For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/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, ImplicitlyCopyable, ImplicitlyDeletable, 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.

Edit this page

Was this page helpful?

Models

Cookbook

Puzzles

Blog

Community

Contact