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 module

block_scaled_matmul_small_bn

Implements the small-BN block-scaled matmul kernel for SM100 (Blackwell B200).

Provides the warp-specialized TMA + UMMA kernel and its host-side launch helpers for block-scaled (MXFP8 / NVFP4) matrix multiplication targeting the small-N MMA shapes that require a cp.async SFB load path instead of TMA.

Structs

• ​B200BlockScaledMatmulSmem: Owns the shared-memory buffers and mbarrier pools for one kernel block.
• ​WarpRole: Encodes the warp role assignments used by the warp-specialized kernel.

Functions

• ​blackwell_block_scaled_matmul_tma_umma_warp_specialized: Launch small-BN block-scaled FP8 matmul kernel on SM100.
• ​blackwell_block_scaled_tma_umma_warp_specialized_kernel:
• ​consumer_main_loop: TileTensor overload of consumer_main_loop for block-scaled MMA (small BN).
• ​load_AB_SFA: Issues multicast TMA loads for one K-group of A, B, and SFA into SMEM.