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
FusedSwiGLUMatmulConfig
struct FusedSwiGLUMatmulConfig[a_type: DType, b_type: DType, c_type: DType, transpose_b: Bool = True]
Compile-time configuration for the SM100 fused GEMM+SwiGLU kernel.
Encodes all parameters needed by the fused epilogue, including the path
choice (register_swiglu=True for register→GMEM, False for
SMEM→TMA) directly in the config so dispatch loops need not carry a
separate boolean.
Pipeline stages are computed automatically from the SMEM budget. For the SMEM→TMA path the half-output-tile overhead is subtracted before solving for the stage count; for the register→GMEM path no extra deduction is made.
Parameters
- a_type (
DType): DType of the A operand elements. - b_type (
DType): DType of the B operand elements. - c_type (
DType): DType of the C output elements. - transpose_b (
Bool): Whether the B operand is stored transposed (defaults toTrue).
Fields
- cta_group (
Int): - mma_shape (
IndexList[Int(3)]): - cluster_shape (
IndexList[Int(3)]): - AB_swapped (
Bool): - block_swizzle_size (
Int): - raster_order (
RasterOrder): - register_swiglu (
Bool): - k_group_size (
Int): - num_pipeline_stages (
Int): - num_clc_pipeline_stages (
Int): - num_accum_pipeline_stages (
Int): - use_bias (
Bool): - block_tile_shape (
IndexList[Int(3)]): - output_tile_shape (
IndexList[Int(2)]): - a_swizzle (
TensorMapSwizzle): - b_swizzle (
TensorMapSwizzle):
Implemented traits
AnyType,
Copyable,
Equatable,
Hashable,
ImplicitlyCopyable,
ImplicitlyDeletable,
Movable,
RegisterPassable,
TrivialRegisterPassable,
Writable
comptime members
accum_type
comptime accum_type = get_accum_type[a_type]()
Methods
__init__
def __init__(*, mma_shape: IndexList[Int(3)], cta_group: Int = Int(2), cluster_shape: IndexList[Int(3)] = Index[Int, Int, Int](Int(2), Int(1), Int(1)), AB_swapped: Bool = False, block_swizzle_size: Int = Int(0), raster_order: RasterOrder = RasterOrder.AlongM, k_group_size: Int = Int(1), num_pipeline_stages: Optional[Int] = None, num_clc_pipeline_stages: Int = Int(2), num_accum_pipeline_stages: Int = Int(2), use_bias: Bool = False, register_swiglu: Bool = False) -> Self
get_kernel_name
write_to
def write_to[W: Writer](self, mut writer: W)
write_repr_to
def write_repr_to(self, mut writer: T)
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!