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 module

matmul_kernel_persistent

Implements persistent warp-specialized matmul kernels for NVIDIA SM90 (Hopper) GPUs.

Extends HopperMatmulSM90Kernel with two entry points: run_persistent, which uses TMA-based async tile loading for aligned K dimensions, and run_unaligned, which falls back to cp.async loading when K alignment does not meet TMA requirements. Both kernels split thread blocks into producer and consumer warp groups that iterate over a persistent tile schedule assigned by a TileScheduler.

Was this page helpful?

Thank you! We'll create more content like this.

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit