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 package

apple

Provides the Apple silicon GPU backend implementations for matmuls.

Modules

• ​fp4_dequant: Portable NVFP4 -> bf16 weight-only dequant for the Apple M5 matmul (W4A16).
• ​fp4_gemv: Apple M5 weight-only NVFP4 (W4A16) GEMV: out = x @ dequant(W)^T at M=1.
• ​fp4_matmul: Apple M5 weight-only NVFP4 (W4A16) matmul: out = activation @ weight^T.
• ​fp8_gemv: Apple M5 weight-only FP8 (W8A16) GEMV: out = x @ dequant(W)^T at M=1.
• ​int8_matmul: Apple M5 int8 W8A8 matmul: out = dequant(A_int8 @ B_int8^T).
• ​matmul2d_fp4: Apple M5 weight-only NVFP4 (W4A16) matmul on the 16x32x16 matmul2d tile.
• ​matmul2d_fp8: Apple M5 weight-only FP8 (W8A16) tiled matmul: bf16 A x fp8 B -> fp32.
• ​matmul_8x8: 8x8 simdgroup_matrix GEMM kernel for Apple Silicon GPUs (M1-M4).
• ​matmul_kernel: Structured Apple M5 simdgroup-tiled matmul (Metal 4 hardware MMA).

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