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 function

grouped_matmul_sm90

def grouped_matmul_sm90[c_type: DType, a_type: DType, b_type: DType, //, *, transpose_b: Bool = True, wgmma_shape: IndexList[Int(3)] = Index[Int, Int, Int](Int(64), Int(256), Int(16)), config: MatmulConfig[a_type, b_type, c_type, transpose_b] = default_config_sm90[a_type, b_type, c_type, transpose_b, wgmma_shape](), elementwise_lambda_fn: Optional[def[dtype: DType, width: SIMDLength, *, alignment: Int = Int(1)](IndexList[Int(2)], SIMD[dtype, width]) capturing thin -> None] = None](c: TileTensor[c_type, Storage=c.Storage, linear_idx_type=c.linear_idx_type], a: TileTensor[a_type, Storage=a.Storage, linear_idx_type=a.linear_idx_type], a_offsets: TileTensor[DType.uint32, Storage=a_offsets.Storage, linear_idx_type=a_offsets.linear_idx_type], max_num_tokens_per_expert: Int, b: TileTensor[b_type, Storage=b.Storage, linear_idx_type=b.linear_idx_type], expert_ids: TileTensor[DType.int32, Storage=expert_ids.Storage, linear_idx_type=expert_ids.linear_idx_type], num_active_experts: Int, ctx: DeviceContext)

Performs grouped GEMM for MoE routing on SM90 (Hopper) GPUs.

Dispatches a batched expert matmul where each expert has a variable number of tokens, stored contiguously in a at offsets given by a_offsets. Expert weight matrices are stacked along axis 0 in b. Uses TMA-based warp-specialized pipelining via HopperMatmulSM90Kernel.

Parameters:

c_type (DType): Output element type.
a_type (DType): A-matrix (activations) element type.
b_type (DType): B-matrix (expert weights) element type.
transpose_b (Bool): Whether B is stored transposed (must be True).
wgmma_shape (IndexList[Int(3)]): WGMMA instruction shape (M, N, K).
config (MatmulConfig[a_type, b_type, c_type, transpose_b]): Full SM90 kernel configuration.
elementwise_lambda_fn (Optional[def[dtype: DType, width: SIMDLength, *, alignment: Int = Int(1)](IndexList[Int(2)], SIMD[dtype, width]) capturing thin -> None]): Optional epilogue applied to each output tile.

Args:

c (TileTensor[c_type, Storage=c.Storage, linear_idx_type=c.linear_idx_type]): Output matrix [total_tokens, N].
a (TileTensor[a_type, Storage=a.Storage, linear_idx_type=a.linear_idx_type]): Activation matrix [total_tokens, K].
a_offsets (TileTensor[DType.uint32, Storage=a_offsets.Storage, linear_idx_type=a_offsets.linear_idx_type]): Per-expert token start offsets into a.
max_num_tokens_per_expert (Int): Maximum tokens for any single expert.
b (TileTensor[b_type, Storage=b.Storage, linear_idx_type=b.linear_idx_type]): Expert weight tensor [num_experts, N, K].
expert_ids (TileTensor[DType.int32, Storage=expert_ids.Storage, linear_idx_type=expert_ids.linear_idx_type]): Active expert indices [num_active_experts].
num_active_experts (Int): Number of experts with non-zero token count.
ctx (DeviceContext): Device context for kernel launch.