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 function

matmul_sm100_fallback

def matmul_sm100_fallback[c_type: DType, a_type: DType, b_type: DType, *, transpose_b: Bool, umma_shape: IndexList[Int(3)], block_tile_shape: IndexList[Int(3)], a_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, b_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, 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[Storage=c.Storage, address_space=c.address_space, linear_idx_type=c.linear_idx_type], a: TileTensor[Storage=a.Storage, address_space=a.address_space, linear_idx_type=a.linear_idx_type], b: TileTensor[Storage=b.Storage, address_space=b.address_space, linear_idx_type=b.linear_idx_type], ctx: DeviceContext)

Launches the SM100 fallback matmul kernel for unsupported shapes or dtypes.

Uses a simple non-pipelined kernel without CLC/TMEM overhead, computing shared memory from the actual tile sizes rather than the hardware maximum. Only transposed-B, bfloat16, and float8_e4m3fn inputs are supported.

Parameters:

• ​c_type (DType): Output element dtype of the result matrix.
• ​a_type (DType): Input element dtype of the LHS matrix; must be bfloat16 or float8_e4m3fn.
• ​b_type (DType): Input element dtype of the RHS matrix; must equal a_type.
• ​transpose_b (Bool): Whether B is stored transposed as (N, K). Must be True.
• ​umma_shape (IndexList[Int(3)]): Tensor core MMA instruction shape as a 3-element IndexList of (M, N, K).
• ​block_tile_shape (IndexList[Int(3)]): CTA block tile shape as a 3-element IndexList of (BM, BN, BK) giving the tile dimensions along M, N, and K.
• ​a_swizzle (TensorMapSwizzle): TMA swizzle mode for A tensor loads (defaults to SWIZZLE_128B).
• ​b_swizzle (TensorMapSwizzle): TMA swizzle mode for B tensor loads (defaults to SWIZZLE_128B).
• ​elementwise_lambda_fn (Optional[def[dtype: DType, width: SIMDLength, *, alignment: Int = Int(1)](IndexList[Int(2)], SIMD[dtype, width]) capturing thin -> None]): Optional epilogue lambda applied after the matmul (defaults to None).