vLLM Machete weight-only GEMM kernel source (machete_mm_kernel.cuh)

Fields

id

sources/kern-machete-vllm-source

URL

https://github.com/vllm-project/vllm/blob/main/csrc/libtorch_stable/quantization/machete/machete_mm_kernel.cuh

Publisher

vllm-project

Kind

repo

Accessed

2026-10-03

notes

Kernel source file (Apache-2.0). Hard evidence for the arch gate and the offline-prepacking mechanism: 'using ArchTag = cutlass::arch::Sm90;' with OperatorClass OpClassTensorOp, TileShapeK = 128*8/sizeof_bits<MmaType>, Alignments derived from 128-bit; LayoutB is declared ColumnMajor with the in-file comment 'not actually used since B has the prepacked layout'. Machete is therefore a Hopper-only WGMMA-tensor-core GEMM consuming a pre-tiled int4 B operand plus group scales and group zeros, not a hardware int4 path - on other archs there is no fallback in this file.

Notes

Kernel source file (Apache-2.0). Hard evidence for the arch gate and the offline-prepacking mechanism: 'using ArchTag = cutlass::arch::Sm90;' with OperatorClass OpClassTensorOp, TileShapeK = 128*8/sizeof_bits<MmaType>, Alignments derived from 128-bit; LayoutB is declared ColumnMajor with the in-file comment 'not actually used since B has the prepacked layout'. Machete is therefore a Hopper-only WGMMA-tensor-core GEMM consuming a pre-tiled int4 B operand plus group scales and group zeros, not a hardware int4 path - on other archs there is no fallback in this file.

Cross-references

This record references no other records.