Kernel Backends Reference

Kernel Backends Reference

Documentation home · Quickstart · Lifecycle

On this page

• Overview
• Workflow at a glance
• Packed storage and execution concepts
• Backend → class → layout → device
• Exact declared quantization contracts
• Marlin and caller-owned scratch
• Machete
• Swordfish
• ExLlama v2 and v3
• Triton v2 and Triton utilities
• Torch families
• BitBLAS and target detection
• GGUF family
• FP8, FP4, and W4A activation families
• AWQ native and portable families
• QQQ, VPTQ, ParoQuant, and bitsandbytes
• MLX checkpoint holders and executable modules
• Packing utilities and extension packers
• Hooked modules during quantization
• Native C++/CUDA source map
• Kernel implementation and validation checklist
• Pitfalls
• Source files
• See also

Overview

This page describes the quantized linear implementations, packing boundaries, support declarations, and native extension sources in GPTQModel. Use Inference Backends to choose a backend; use this reference to understand exactly which class executes a layer, what it stores, and why validation can reject a seemingly compatible checkpoint.

The implementation baseline is commit c1976918cfe3703fbcf6d1ce2440a82b26394375. Optional dependencies, installed Torch operators, target architecture, tensor qtype, and per-layer dimensions remain part of the executable contract. Declaration tables are not a promise that every declared combination works without runtime validation.

Workflow at a glance

Reading the diagram: Independent entry routes. Initialization is kernel-specific; scratch is not universal. Dense or expanded weight caches may not survive their consuming operation.

Text version of the workflow

• Raw packing: Quantizer outputs → hf_pack_layer → Initialize if required → Layer forward.
• Existing packed source: Checkpoint-layout layer → hf_repack_layer → Initialize target → Layer forward.
• Changed packed state: State changes → Invalidate owned state → Reinitialize as needed → Next forward.

Open full-size SVG

Packed storage and execution concepts

Term	Meaning
Checkpoint layout	Serialization of quantized codes and metadata. GPTQ, AWQ GEMM, QQQ, VPTQ, GGUF, and EXL3 do not share one universal packed layout.
Runtime ABI	A kernel-private representation created after loading: tiled/prepacked matrices, scale permutations, native handles, or specialized INT4 operator inputs.
Group size -1	Channelwise grouping represented by the class's full input dimension rather than a fixed finite group length.
desc_act / g_idx	Activation-order grouping and input-to-group mapping. Some kernels reorder operands; others reject activation ordering.
pack_dtype	Integer word storage dtype, not the logical quantization bit width or activation dtype.
SUPPORTS_FORMATS priority	Positive values enter AUTO ordering; zero retains explicit support without generic AUTO selection.
post_init()	Runtime initialization after packed checkpoint buffers have been populated; not a substitute for loading missing weights.
Raw packing	Encoding quantizer outputs from an original layer; requires a pack API and is distinct from loading/repacking existing codes.

For standard aligned GPTQ INT4 with int32 words, input width I, output width O, and G groups, qweight is [I/8, O], qzeros is [G, O/8], scales is [G, O], and g_idx is [I]. These are checkpoint-side shapes; a runtime-transformed class may no longer expose that ABI after initialization. GPTQ v1/v2 zero-point interpretation must be preserved when converting; a raw copy of qzeros is not always a semantic conversion.

Backend → class → layout → device

Paths in this table are relative to gptqmodel/nn_modules/qlinear/. ALL means the class's portable device declaration, not that every runtime implements every required operation. Device-sensitive validators can narrow declarations. MLX rows are checkpoint holders, not directly executable Torch layers. EXL3 classes live outside qlinear and are listed separately afterward.

These are layer-level contracts, not a model-loader allowlist. For example, AwqMacheteLinear declares Marlin-format support, but the local model loader requires the matching Marlin backend or MLX for a serialized Marlin checkpoint. The same boundary applies to serialized BitBLAS checkpoints; see format-specific loader branches.

Native AWQ/QQQ device metadata also does not establish that CUDA/PTX extension sources run on ROCm. The AWQ and QQQ pages distinguish declarations from the native implementation's portability constraints.

BACKEND	Class and qlinear source	Checkpoint / runtime packing relationship	Declared device families
GPTQ_BITBLAS	BitblasLinear bitblas.py	GPTQ/BitBLAS → BitBLAS transformed weight/scales/zeros	CUDA
AWQ_BITBLAS	AWQBitBlasKernel bitblas_awq.py	AWQ GEMM/BitBLAS → BitBLAS transformed weight/scales/zeros	CUDA
BITSANDBYTES	BitsAndBytesLinear bitsandbytes.py	bitsandbytes codes plus QuantState / INT8 state	CPU, CUDA
GPTQ_EXLLAMA_V2	ExllamaV2Linear exllamav2.py	GPTQ packed buffers → extension matrix handle	CUDA
AWQ_EXLLAMA_V2	AwqExllamaV2Linear exllamav2_awq.py	AWQ GEMM → ExLlama v2-compatible packed buffers	CUDA
FP8_TORCH	TorchFP8Linear fp8.py	FP8 weight codes plus scales	CPU, CUDA, ROCM, XPU, MPS
AWQ_GEMM	AwqGEMMLinear gemm_awq.py	AWQ GEMM packed qweight/qzeros/scales	CUDA
AWQ_GEMM_TRITON	AwqGEMMTritonLinear gemm_awq_triton.py	AWQ GEMM packed qweight/qzeros/scales	CUDA, ROCM
AWQ_GEMV	AwqGEMVLinear gemv_awq.py	AWQ GEMV layout	CUDA, ROCM
AWQ_GEMV_FAST	AwqGEMVFastLinear gemv_fast_awq.py	GEMV_FAST or LLM_AWQ layout, according to class	CUDA, ROCM
AWQ_GEMV_FAST	LLMAwqLinear gemv_fast_awq.py	GEMV_FAST or LLM_AWQ layout, according to class	CUDA, ROCM
GGUF_TORCH	GGUFTorchLinear gguf.py	Native GGUF block bytes and qtype	ALL
GGUF_CPP_CPU	GGUFCppKernel gguf_cpp.py	Native GGUF block bytes; ggml bridge plans	CPU
GGUF_CPP_CUDA	GGUFCudaKernel gguf_cpp.py	Native GGUF block bytes; ggml bridge plans	CUDA
GGUF_TRITON	GGUFTritonKernel gguf_triton.py	Native GGUF block bytes and qtype	CUDA
GPTQ_MACHETE	MacheteLinear machete.py	GPTQ → Machete prepacked B	CUDA
AWQ_MACHETE	AwqMacheteLinear machete_awq.py	AWQ/Marlin → Machete prepacked B	CUDA
GPTQ_MARLIN	MarlinLinear marlin.py	GPTQ → Marlin tiles; permuted scales and activation-order metadata	CUDA
AWQ_MARLIN	AwqMarlinLinear marlin_awq.py	AWQ/Marlin → Marlin tiles; scale/zero permutations	CUDA
MLX	AWQBitBLASMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	AWQMarlinMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	AwqGemvFastMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	AwqGemvMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	AwqMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	BitsAndBytesMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	FP8MlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	GGUFMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	GPTQBitBLASMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	GPTQMarlinMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	LLMAwqMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	MlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	ParoMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
MLX	QQQMlxQuantLinear mlx.py	Torch checkpoint holder → method-specific MLX quantized module	MPS
PAROQUANT_CUDA	ParoLinear paroquant.py	Packed affine weights plus rotation/scaling metadata	ALL
PAROQUANT_TRITON	ParoQuantTritonLinear paroquant_triton.py	ParoQuant packed affine weights plus rotation metadata	CUDA
QQQ	QQQLinear qqq.py	QQQ packed B plus group/channel scales	CUDA, ROCM
QQQ_TORCH	QQQTorchLinear qqq.py	QQQ packed B plus group/channel scales	ALL
AWQ_SWORDFISH	AwqSwordfishLinear swordfish.py	GPTQ or AWQ → Swordfish prepacked B	CUDA
GPTQ_SWORDFISH	SwordfishLinear swordfish.py	GPTQ or AWQ → Swordfish prepacked B	CUDA
GPTQ_TORCH	TorchLinear torch.py	GPTQ continuous or GPTQ_P planar words	ALL
GPTQ_TORCH	TorchQuantEmbeddings torch.py	GPTQ continuous or GPTQ_P planar words	ALL
GPTQ_TORCH_ATEN	TorchAtenLinear torch_aten_kernel.py	GPTQ → CPU ATen int4pack + affine scale/zero offsets	CPU
AWQ_TORCH_ATEN	TorchAtenAwqLinear torch_aten_kernel_awq.py	AWQ → CPU ATen int4pack + affine scale/zero offsets	CPU
AWQ_TORCH	AwqTorchLinear torch_awq.py	AWQ GEMM words	ALL
GPTQ_TORCH_CPU_AVX512	TorchCpuAvx512Linear torch_cpu_avx512.py	GPTQ INT4 → native CPU packed layout	CPU
AWQ_TORCH_CPU_AVX512	TorchCpuAvx512AwqLinear torch_cpu_avx512_awq.py	AWQ INT4 → native CPU packed layout	CPU
GPTQ_TORCH_FUSED	TorchFusedLinear torch_fused.py	Checkpoint GPTQ → Torch fused INT4 ABI	CPU, XPU
AWQ_TORCH_FUSED	TorchFusedAwqLinear torch_fused_awq.py	AWQ → Torch fused INT4 ABI	CPU, XPU
GPTQ_TORCH_INT8	TorchInt8Linear torch_int8.py	Native 8-bit GPTQ → centered INT8 operand plus channel scales	CPU
AWQ_TORCH_INT8	TorchInt8AwqLinear torch_int8_awq.py	Retired; no executable runtime layout	CPU
GPTQ_TRITON	TritonV2Linear tritonv2.py	GPTQ packed qweight/qzeros/scales/g_idx	CUDA
VPTQ_CUDA	VPTQCudaLinear vptq_cuda.py	Packed vector indices and codebooks	CUDA
VPTQ_TORCH	VPTQLinear vptq_torch.py	Packed vector indices, codebooks and optional residual/outlier state	ALL
GPTQ_W4AFP8	W4AFP8Linear w4a_floatx.py	Saved GPTQ INT4 plus runtime E4M3 operand	CUDA
GPTQ_W4A_NVFP4	W4ANVFP4Linear w4a_nvfp4.py	Saved GPTQ INT4 plus runtime FP4/FP8 operands	CUDA

Method-specific classes outside the generic qlinear registry

Backend	Class / source	Packed state	Execution boundary
EXL3_EXLLAMA_V3	ExllamaV3Linear, gptqmodel/nn_modules/exllamav3.py	trellis, suh, svh, su, sv, optional bias, mcg, mul1	Native LinearEXL3 wrapper; loader requires CUDA/HIP.
EXL3_TORCH	ExllamaV3TorchLinear, gptqmodel/nn_modules/exllamav3_torch.py	EXL3 trellis/codebook and transform metadata	Torch decoding/matmul implementation, including chunked execution paths.
Native FP4 weight helper, no separate BACKEND member	TorchFP4Linear, qlinear/fp4.py	Packed NVFP4 weight, weight scale, block size, original dtype	TorchAO NVFP4Tensor on supported native paths; CPU decodes current weights for F.linear. Not the same as GPTQ W4A NVFP4.

Exact declared quantization contracts

This table transcribes class declarations, resolving inherited constants where applicable. Format names are FORMAT members; numbers after them are priorities. It deliberately separates bit/group/training declarations from the architecture and operator checks in subsequent sections. For example, a GGUF logical bit width does not identify the complete GGML block qtype, and a class declaring training support may use a different forward path in training mode.

Class	Format: AUTO priority	Bits	Group sizes	desc_act	sym	Training
BitblasLinear	BITBLAS: 30, GPTQ: 30, GPTQ_V2: 30	2, 4, 8	-1, 32, 64, 128	False	False, True	False
AWQBitBlasKernel	GEMM: 0, BITBLAS: 30	4	-1, 32, 64, 128	False, True	False, True	False
BitsAndBytesLinear	BITSANDBYTES: 40	4, 8	—	—	—	False
ExllamaV2Linear	GPTQ: 80, GPTQ_V2: 80	4	-1, 32, 64, 128	True, False	True, False	False
AwqExllamaV2Linear	GEMM: 80	4	-1, 32, 64, 128	True, False	True, False	False
TorchFP8Linear	FP8: 15	8	—	—	—	True
AwqGEMMLinear	GEMM: 60	4	-1, 16, 32, 64, 128	True, False	True, False	True
AwqGEMMTritonLinear	GEMM: 50	4	-1, 32, 64, 128	True, False	True, False	True
AwqGEMVLinear	GEMV: 40	4	-1, 64, 128	True, False	True, False	True
AwqGEMVFastLinear	GEMV_FAST: 30	4	-1, 16, 32, 64, 128	True, False	True, False	True
LLMAwqLinear	LLM_AWQ: 100	4	-1, 16, 32, 64, 128	True, False	True, False	True
GGUFTorchLinear	GGUF: 15	1, 2, 3, 4, 5, 6, 8	—	—	—	True
GGUFCppKernel	GGUF: 25	1, 2, 3, 4, 5, 6, 8	—	—	—	False
GGUFCudaKernel	GGUF: 35	1, 2, 3, 4, 5, 6, 8	—	—	—	False
GGUFTritonKernel	GGUF: 45	1, 4, 5, 6	—	—	—	False
MacheteLinear	GPTQ: 100	4, 8	-1, 64, 128	True, False	True, False	False
AwqMacheteLinear	GEMM: 100, MARLIN: 100	4, 8	-1, 32, 64, 128	False	True, False	False
MarlinLinear	GPTQ: 90, GPTQ_V2: 90, MARLIN: 90	4, 8	-1, 32, 64, 128	True, False	True	False
AwqMarlinLinear	GEMM: 90, MARLIN: 90	4, 8	-1, 32, 64, 128	True, False	True, False	False
AWQBitBLASMlxQuantLinear	BITBLAS: 0	2, 3, 4, 5, 6, 7, 8	-1, 32, 64, 128	True, False	True, False	False
AWQMarlinMlxQuantLinear	MARLIN: 0	4, 8	-1, 32, 64, 128	True, False	True, False	False
AwqGemvFastMlxQuantLinear	GEMV_FAST: 0	4	-1, 16, 32, 64, 128	True, False	True, False	False
AwqGemvMlxQuantLinear	GEMV: 0	4	-1, 64, 128	True, False	True, False	False
AwqMlxQuantLinear	GEMM: 0	2, 3, 4, 5, 6, 7, 8	-1, 16, 32, 64, 128	False	True, False	False
BitsAndBytesMlxQuantLinear	BITSANDBYTES: 0	4, 8	—	—	—	False
FP8MlxQuantLinear	FP8: 0	8	—	—	—	False
GGUFMlxQuantLinear	GGUF: 0	1, 2, 4, 5, 6, 8	—	—	—	False
GPTQBitBLASMlxQuantLinear	BITBLAS: 0	2, 4, 8	-1, 32, 64, 128	False	True, False	False
GPTQMarlinMlxQuantLinear	MARLIN: 0	4, 8	-1, 32, 64, 128	False	True	False
LLMAwqMlxQuantLinear	LLM_AWQ: 0	4	-1, 16, 32, 64, 128	True, False	True, False	False
MlxQuantLinear	GPTQ: 0, GPTQ_V2: 0, GPTQ_P: 0	2, 3, 4, 5, 6, 7, 8	-1, 16, 32, 64, 128, 256, 512, 1024	False	True, False	False
ParoMlxQuantLinear	PAROQUANT: 0	4	-1, 16, 32, 64, 128	True, False	True	False
QQQMlxQuantLinear	QQQ: 0	4	-1, 128	True, False	True	False
ParoLinear	PAROQUANT: 55	4	-1, 16, 32, 64, 128	True, False	True	True
ParoQuantTritonLinear	PAROQUANT: 0	4	-1, 16, 32, 64, 128	True, False	True	True
QQQLinear	QQQ: 100	4	-1, 128	False	True	False
QQQTorchLinear	QQQ: 90	4	-1, 128	False	True	False
AwqSwordfishLinear	GEMM: 101	4	-1, 32, 64, 128	False	True, False	False
SwordfishLinear	GPTQ: 101	4, 8	-1, 32, 64, 128	True, False	True, False	False
TorchLinear	GPTQ: 20, GPTQ_V2: 20, GPTQ_P: 20	2, 3, 4, 5, 6, 7, 8	-1, 16, 32, 64, 128, 256, 512, 1024	True, False	True, False	True
TorchQuantEmbeddings	GPTQ: 20, GPTQ_V2: 20, GPTQ_P: 20	2, 3, 4, 5, 6, 7, 8	-1, 16, 32, 64, 128, 256, 512, 1024	True, False	True, False	False
TorchAtenLinear	GPTQ: 110, GPTQ_V2: 110	4	32, 64, 128, 256	True, False	True, False	True
TorchAtenAwqLinear	GEMM: 110	4	32, 64, 128, 256	True, False	True, False	True
AwqTorchLinear	GEMM: 10	2, 3, 4, 5, 6, 7, 8	-1, 16, 32, 64, 128	True, False	True, False	True
TorchCpuAvx512Linear	GPTQ: 120, GPTQ_V2: 120	4	16, 32, 64, 128, 256, 512	True, False	True, False	False
TorchCpuAvx512AwqLinear	GEMM: 120	4	16, 32, 64, 128, 256, 512	True, False	True, False	False
TorchFusedLinear	GPTQ: 50, GPTQ_V2: 50	4	-1, 16, 32, 64, 128	True, False	True, False	True
TorchFusedAwqLinear	GEMM: 20	4	-1, 16, 32, 64, 128	True, False	True, False	True
TorchInt8Linear	GPTQ: 0, GPTQ_V2: 0	8	-1, 16, 32, 64, 128	True, False	True, False	False
TorchInt8AwqLinear	GEMM: 0	4	-1, 16, 32, 64, 128	True, False	True, False	False
TritonV2Linear	GPTQ: 40, GPTQ_V2: 40	2, 4, 8	-1, 16, 32, 64, 128, 256, 512, 1024	True, False	True, False	True
VPTQCudaLinear	VPTQ: 100	1, 2, 3, 4, 5, 6, 7, 8	—	—	—	False
VPTQLinear	VPTQ: 100	1, 2, 3, 4, 5, 6, 7, 8	—	—	—	False
W4AFP8Linear	GPTQ: 0, GPTQ_V2: 0	4	128	False	True	False
W4ANVFP4Linear	GPTQ: 0, GPTQ_V2: 0	4	128	False	True	False

The base implementation also validates platform, activation dtype, packing dtype, adapters, automatic padding, and input/output divisibility. AUTO passes the adapter into candidate validation; the explicit-backend selector branch currently does not, so explicit selection alone is not proof of adapter compatibility. Sharded-loading support uses SUPPORTS_SHARDED_LOAD when present, otherwise SUPPORTS_SHARDS. Swordfish explicitly disables sharded loading. Inspect these additional fields when adding a format or diagnosing a shape that passes the bit/group columns but fails selection.

Operand dtypes, dimension declarations, and sharding

The following are declarations, including inherited values. A dash means no concrete declaration was found for that field, not unrestricted support. Automatic padding and specialized post-initialization can impose further constraints or adapt the checkpoint dimensions. Private Torch operations can have a narrower actual compute dtype than the supported interface, as documented in the family sections. A declared sharding flag does not override method-specific loader restrictions.

Class	Activation dtypes	Pack word dtypes	Input divisibility	Output divisibility	Auto padding	Sharded load
BitblasLinear	float16, bfloat16	int32	16	16	False	True
AWQBitBlasKernel	float16, bfloat16	int32	16	16	False	True
BitsAndBytesLinear	float16, bfloat16, float32	int8, int16, int32, int64	1	1	False	True
ExllamaV2Linear	float16, bfloat16	int32	32	32	False	True
AwqExllamaV2Linear	float16, bfloat16	int32	32	32	False	True
TorchFP8Linear	float16, bfloat16, float32	int8, int16, int32, int64	1	1	False	True
AwqGEMMLinear	float16, bfloat16	int32	1	1	False	True
AwqGEMMTritonLinear	float16	int32	32	8	False	True
AwqGEMVLinear	float16	int32	1	64	False	True
AwqGEMVFastLinear	float16	int16	1	1	False	True
LLMAwqLinear	float16	int16	1	1	False	True
GGUFTorchLinear	float16, bfloat16, float32	int8, int16, int32	1	1	True	True
GGUFCppKernel	float16, bfloat16, float32	int8, int16, int32	1	1	True	True
GGUFCudaKernel	float16, bfloat16, float32	int8, int16, int32	1	1	True	True
GGUFTritonKernel	float16, bfloat16, float32	int8, int16, int32	1	1	True	True
MacheteLinear	float16, bfloat16	int32	64	128	False	True
AwqMacheteLinear	float16, bfloat16	int32	64	128	False	True
MarlinLinear	float16, bfloat16	int32	1	1	False	True
AwqMarlinLinear	float16, bfloat16	int32	1	1	False	True
AWQBitBLASMlxQuantLinear	float16	int32	16	16	False	True
AWQMarlinMlxQuantLinear	float16	int32	8	8	False	True
AwqGemvFastMlxQuantLinear	float16	int16	64	8	False	True
AwqGemvMlxQuantLinear	float16	int32	32	64	False	True
AwqMlxQuantLinear	float16	int32	8	8	False	True
BitsAndBytesMlxQuantLinear	float16	int8, int16, int32, int64	1	1	False	True
FP8MlxQuantLinear	float16	int8, int16, int32, int64	1	1	False	True
GGUFMlxQuantLinear	float16	int8, int16, int32	32	8	False	True
GPTQBitBLASMlxQuantLinear	float16	int32	32	32	False	True
GPTQMarlinMlxQuantLinear	float16	int32	32	32	False	True
LLMAwqMlxQuantLinear	float16	int16	64	8	False	True
MlxQuantLinear	float16	int32	32	32	False	True
ParoMlxQuantLinear	float16	int32	32	8	False	True
QQQMlxQuantLinear	float16	int32	64	64	False	True
ParoLinear	float16, bfloat16	int32	1	1	False	True
ParoQuantTritonLinear	float16, bfloat16	int32	1	1	False	True
QQQLinear	float16, bfloat16	int32	64	64	False	True
QQQTorchLinear	float16, bfloat16	int32	64	64	False	True
AwqSwordfishLinear	float16, bfloat16	int32	64	64	False	False
SwordfishLinear	float16, bfloat16	int32	64	64	False	False
TorchLinear	float16, bfloat16	int8, int16, int32	1	1	True	True
TorchQuantEmbeddings	float16, bfloat16	int8, int16, int32	1	1	True	True
TorchAtenLinear	float16, bfloat16	int32	1	1	True	True
TorchAtenAwqLinear	float16, bfloat16	int32	1	1	True	True
AwqTorchLinear	float16, bfloat16	int32	1	8	False	True
TorchCpuAvx512Linear	bfloat16	int32	2	32	False	True
TorchCpuAvx512AwqLinear	bfloat16	int32	2	32	False	True
TorchFusedLinear	float32, float16, bfloat16	int32	1	1	True	True
TorchFusedAwqLinear	float32, float16, bfloat16	int32	1	1	True	True
TorchInt8Linear	float16, bfloat16, float32	int32	1	1	True	True
TorchInt8AwqLinear	float16, bfloat16, float32	int32	1	1	False	True
TritonV2Linear	float16, bfloat16	int32, int16, int8	32	32	True	True
VPTQCudaLinear	float16, bfloat16, float32	int32	1	1	True	True
VPTQLinear	float16, bfloat16, float32	int32	1	1	True	True
W4AFP8Linear	float16, bfloat16	int32	128	32	False	True
W4ANVFP4Linear	float16, bfloat16	int32	128	32	False	True

Marlin and caller-owned scratch

qlinear/marlin.py and marlin_awq.py adapt GPTQ/AWQ tensors to Marlin's tile layout. utils/marlin.py owns extension access, device checks, padding helpers, repacking, scale/zero permutations, workspace allocation, and GEMM wrappers. GPTQ Marlin permits CUDA capability 7.5+, whereas AWQ Marlin's device floor is 8.0. ROCm is rejected. Input/output packing divisibility and activation-order alignment are validated even when broader declarations permit padding.

GPTQ Marlin reads GPTQMODEL_MARLIN_USE_FP32 with default True. Disabling it changes reduction precision; this is not a lossless memory-only toggle. The low-level helper defaults and the module's configured value must not be confused: the module forwards its own reduction choice to the GEMM wrapper.

Scratch API

gptq_marlin_gemm accepts caller-provided c_tmp and a_tmp; without them the native implementation retains its normal allocation behavior. marlin_scratch_sizes() returns element counts, not bytes. The query and GEMM share the native size calculation, including executed/padded K, activation-order permutation requirements, and reduction mode.

def __init__(
    self,
    device: torch.device | str | int,
    *,
    max_cached_bytes: int = 64 << 20,
) -> None:
    ...

import torch
from gptqmodel.utils.marlin_scratch import MarlinScratchContext

with torch.cuda.device("cuda:0"):
    with MarlinScratchContext("cuda:0", max_cached_bytes=64 << 20):
        # model must already be loaded with a compatible Marlin kernel.
        result = model(**inputs)

Scratch property	Contract
Ownership	One CUDA device, host thread, and CUDA stream; no nested/reentrant contexts or simultaneous leases.
Retained budget	Default 64 MiB total including the lock workspace; oversized requests use temporary storage. This is not a bound on outputs or allocator-reserved memory.
Buffers	Growing reduction and activation buffers plus a separate zero-initialized lock workspace; dtype changes isolate activation storage.
Supplied native scratch	Checked for device, dtype, contiguity, capacity, 16-byte alignment, and shared storage with GEMM operands/other scratch. Invalid buffers fail rather than silently falling back.
CUDA graphs	The eager context rejects capture. A low-level graph caller must preallocate sufficient scratch and locks and retain exclusive ownership for the whole graph lifetime.
Lifetime	No scratch is saved in module checkpoints. Leaving a with scope deactivates the context but does not clear a context object kept by the caller; bounded buffers can be reused on later compatible entries. clear() releases retained ownership with stream-lifetime tracking, without synchronizing.

The feature is opt-in and does not claim an established latency improvement or concurrent-forward support. docs/marlin_scratch.md explains capacity formulas and graph ownership in detail.

Machete

MacheteLinear and AwqMacheteLinear use CUTLASS-based prepacked matrices. GPTQ consumes GPTQ format; AWQ declares GEMM and MARLIN inputs. utils/machete.py manages extension generation/loading, prepacking, schedules, and supported quantization types.

The runtime is Hopper-only, capability exactly 9.0, with at least 204800 bytes of opt-in shared memory per block. Its sm_90a build is not forward-compatible evidence for Blackwell. The helper pins CUTLASS 4.7.1, validates the archive checksum/version, and caches generated sources. An import/JIT cache is separate from per-device runtime validation.

Loading a supported checkpoint still requires satisfying the class's shape, scale dtype, group size, symmetry, and activation-order rules. Repacking to Machete is a layout transform, not a new quantization pass. Do not serialize the private runtime B layout as ordinary GPTQ qweight without a matching format contract.

Swordfish

SwordfishLinear and AwqSwordfishLinear are in qlinear/swordfish.py; utils/swordfish.py exposes runtime checks, prepacking, quant-type queries, and swordfish_mm. Native sources are under gptqmodel_ext/swordfish/, including stable-libtorch bindings.

Swordfish's vendored code is licensed GNU AGPL v3.0+, as stated in its extension loader and gptqmodel_ext/swordfish/LICENSE / licenses/SWORDFISH. Review that dependency's license when distributing a build; it should not be inferred from another kernel's license.

Property	Current implementation
Architectures	Exact native build targets include sm_100a, sm_103a, sm_110a; runtime support is derived from architecture flags, including their PTX rules. Not a generic SM100+ predicate.
Torch requirement	Stable ABI target Torch 2.10; older or unparseable versions fail validation.
Device/platform	CUDA on Linux; ROCm rejected.
Shape	Prepacked block shape (64, 64); both feature dimensions divisible by 64.
Groups	-1, 32, 64, 128.
Activations	FP16 or BF16.
Weight quant types	With zero points: uint4; without zero points: uint4b8 or uint8b128. Zero-point INT8 is rejected.
Sharded load / training	Both disabled by the classes.

GPTQ accepts 4/8-bit and AWQ accepts 4-bit. Generic declaration support is followed by shape, extension, quant-type, and exact-device checks. Swordfish's Blackwell targets must not be conflated with W4A NVFP4's separate capability 12.1 floor.

ExLlama v2 and v3

ExLlama v2

ExllamaV2Linear consumes 4-bit GPTQ/GPTQ_V2, with group sizes -1, 32, 64, or 128. Group size 16 is deliberately absent. Both feature dimensions must be divisible by 32. The class declares CUDA rather than ROCm and does not support training.

post_init() prepares an extension matrix handle and FP16 scales, with scratch supplied by the loader. Forward casts input to FP16 for native execution; the interface declares FP16/BF16. AwqExllamaV2Linear first adapts AWQ GEMM tensors. Do not reuse an initialized native handle after replacing its underlying buffers without rebuilding runtime state.

Native v2 bindings and CUDA sources live under gptqmodel_ext/exllamav2/; utility setup is in gptqmodel/utils/exllamav2.py.

ExLlama v3 / EXL3

EXL3 is not GPTQ INT4 with a different GEMM. ExllamaV3Linear wraps the bundled EXL3 runtime and its trellis/transformation tensors. ExllamaV3TorchLinear provides a Torch implementation with its own codebook decoding and chunked matmul paths. They are loaded by the EXL3-specific branch rather than generic BaseQuantLinear discovery.

The bundled Python runtime is under gptqmodel/exllamav3/; native code is under gptqmodel_ext/exllamav3/. The loader chooses between native and Torch EXL3 for AUTO based on the requested device and CUDA/HIP availability. mlx_exl3.py implements a separate packed-trellis MLX path. See EXL3.

Triton v2 and Triton utilities

TritonV2Linear supports 2/4/8-bit GPTQ/GPTQ_V2 with CUDA declarations and both inference and training support. It consumes quantized words, scales, zeros, and group indexes directly through Triton kernels rather than maintaining a permanently dense weight matrix. The validator requires Triton 2.0.0+ and, when the GIL is disabled, 3.4.0+.

These are embedded validator thresholds, not the current packaging requirement: the triton extra in pyproject.toml requests Triton 3.8.0 or newer on Python below 3.15. See Inference Backends for the distinction between optional dependency metadata and runtime validation.

File under nn_modules/triton_utils/	Responsibility
kernels.py	Packed 2/4/8-bit quantized matmul, transpose matmul, autograd wrappers, and inference-only wrapper.
dequant.py	Packed-weight decoding, dtype mapping, and quantized matmul helper/autograd wrapper. A dequantization helper is not itself permission to retain expanded weights.
gptq_gemv.py	Specialized packed GPTQ matrix-vector execution.
gptq_gemm.py	Specialized small-M packed GPTQ matrix multiplication and shape limits.
custom_autotune.py	CustomizedTritonAutoTuner, autotune decoration, and matmul configuration pruning.
mixin.py	TritonModuleMixin; the current warmup(..., transpose=False, seqlen=2048) body is a no-op.

Triton helpers are also used by selected portable/ROCm paths. A helper's device context code does not expand the selectable devices of every class importing it. The declaration table and runtime validator jointly determine selection.

Torch families

Portable GPTQ and AWQ

TorchLinear supports 2–8-bit GPTQ-family weights, including GPTQ_P planar packing. AwqTorchLinear is the portable AWQ GEMM implementation. TorchQuantEmbeddings is a GPTQ packed embedding implementation, not another linear GEMM backend: it sets SUPPORTS_BACKEND_SELECTION=False and is excluded from general selector discovery despite sharing the GPTQ_TORCH declaration. Its rows above describe the role-specific class, not an additional AUTO candidate. These classes provide format-aware decoding and fallback math; "Torch" does not mean weights are permanently stored as a normal dense nn.Linear.weight.

torch.py includes a metadata shim for integrations that inspect weight attributes, device-aware unpacking (including NPU shift handling), and optional Triton paths. Runtime-cache keys must include the owning packed tensor identities and versions; see _runtime_cache.py and Inference Backends.

Fused INT4

TorchFusedLinear / TorchFusedAwqLinear declare CPU and XPU, 4-bit weights, and training support. The fused operator ABI is not the same on CPU and XPU.

docs/torch_fused_int4_transformations.md explains the checkpoint starting layout, XPU nibble unpack/reorder/transposition, and CPU int4pack conversion. Activation-order indices can require input permutation. AWQ first adapts its GEMM packing into the appropriate fused representation. Runtime code remains authoritative if historical line references in the note have shifted.

Training can use the Torch path and switch qzero format state; therefore a training-support declaration is not proof that the same fused inference operator has a backward implementation. Use the class's train() and normal initialization rather than manually toggling ABI buffers.

CPU ATen INT4

TorchAtenLinear and TorchAtenAwqLinear call installed ATen CPU int4pack conversion/matmul operators. GPTQ declares 4-bit, groups 32/64/128/256, automatic padding, and training support. validate_once() requires both _convert_weight_to_int4pack_for_cpu and _weight_int4pack_mm_for_cpu.

The affine zero-offset conversion is important. The CPU operator's signed-code form uses scale * (signed_code - 2^(bits-1)) + zero_offset; the adapter computes (2^(bits-1) - zero_code) * scale so it reproduces the checkpoint's scale * (code - zero_code). Supplying raw zero codes where the operator expects offsets changes the math.

CPU AVX512-BF16

TorchCpuAvx512Linear and TorchCpuAvx512AwqLinear use int4_qzero_group_scaled_cpu. Their validators check AVX512-BF16 and resolve native pack/GEMM ops. Declared GPTQ dimensional constraints include input width divisible by 2 and output width divisible by 32; other padding/packing rules still apply. The specialized fused input dtype is BF16, with the class controlling other interface behavior. Training is disabled.

Native INT8 and the retired AWQ expansion

TorchInt8Linear is an 8-bit-only, CPU, explicit-only backend. post_init() converts loaded 8-bit GPTQ tensors to a centered INT8 operand/channel-scale representation for _weight_int8pack_mm. It rejects raw packing and training. It is not a path for persistent INT4-to-INT8 expansion.

TorchInt8AwqLinear is a retired compatibility class. It has priority zero and SUPPORTS_RAW_PACKING=False; construction and validation raise a migration error explaining that persistent AWQ INT4-to-INT8 expansion is disabled. Its continued presence in the enum is not runtime support.

Lookahead

qlinear/lookahead.py retains configure_lookahead_chain() and configure_default_lookahead() for compatibility. They route through deprecated lookahead controls; dense lookahead is disabled with a warning, and no dense weight chain is retained. Do not document these helpers as a current inference acceleration feature.

BitBLAS and target detection

bitblas.py and bitblas_awq.py create BitBLAS operators and transform packed weights/scales/zeros into the runtime layout. The GPTQ class is BitblasLinear; the module also exposes a compatibility BitBLASLinear spelling used by utilities. Declared bits are 2/4/8 for GPTQ and 4 for AWQ. The module sets MINIMUM_BITBLAS_VERSION = "0.1.0.post1".

bitblas_target_detector.py obtains NVIDIA target tags from TVM, fuzzy-matches the GPU model with threshold 25, verifies target architecture compatibility, and falls back to "cuda" when appropriate. OEM model names pg506-230 and pg506-232 map to NVIDIA A100. The main module includes an SM90 target fallback limit for compatibility; target-string fallback is not independent validation that generated code supports every later device.

utils/bitblas.py distinguishes already-serialized BITBLAS from GPTQ-to-BitBLAS repacking. BITBLAS_ENABLE_TUNING can override tuning; absent that variable, repacking does not enable tuning by default, while non-repack conversion does. For the documented BitBLAS version, BF16 is an interface dtype with FP16 operator execution and boundary casting because the bundled low-bit BF16 code generation fails; accumulation remains FP32. Do not promise native BF16 compute from the support dtype list alone.

GGUF family

GGUF stores typed GGML quantization blocks, not generic GPTQ groups. GGUFTorchLinear reads native packed GGUF buffers and selects fused/tiled or current-call decoding paths according to qtype and workload. GGUFTritonKernel supports its own subset of qtypes and executes with FP16 work inputs before restoring the interface dtype. GGUFCppKernel uses the ggml CPU bridge with a Torch CPU alternative where its dispatch permits; GGUFCudaKernel keeps native packed buffers and caches CUDA execution plans.

Class	Important execution boundary
GGUFTorchLinear	Broad logical bit declaration is not a blanket qtype guarantee; inspect qtype-specific implementation.
GGUFTritonKernel	Declares logical bits 1/4/5/6 and CUDA; requires Triton and its free-threaded-Python version gate.
GGUFCppKernel	Requires the bridge's CPU availability; CPU input required.
GGUFCudaKernel	Requires the bridge's CUDA availability; CUDA input required; native matmul plan is reused by compatible call signature.

The model loader imposes additional native Q1_0/Q2_0 restrictions described in Inference Backends. These local kernel descriptions do not apply to the dense Q1_0 FAST import path: in from_pretrained() with no quantization configuration, native Q1_0 AUTO resolves to FAST and can bypass local packed execution entirely. LOW_MEMORY uses the quantized-loading redirection instead. Plan or byte-buffer caching is different from retaining decoded dense weights. See GGUF for the format itself.

FP8, FP4, and W4A activation families

Weight-format FP8 and FP4

TorchFP8Linear uses FP8 checkpoint codes plus scales. It checks for Torch FP8 dtypes, attempts _scaled_mm when its conditions permit, and otherwise uses current-call decoded matmul; a failed scaled-mm path disables that path for the module. Compiled CPU execution has an eager fallback. Bias and adapters are applied at the module boundary.

TorchFP4Linear holds NVFP4-packed checkpoint weights and constructs a TorchAO NVFP4Tensor view for native execution. CPU execution decodes the live FP4 buffers for the current F.linear. This helper has no independent BACKEND enum member and should not be confused with W4A, which starts from GPTQ INT4 weights and quantizes activations.

W4A kernels and constraints

W4AFP8Linear and W4ANVFP4Linear both have zero AUTO priority. They require explicit/activation-policy-directed selection, GPTQ/GPTQ_V2 INT4, group size 128, desc_act=False, symmetry, int32 packing, input width divisible by 128, output width divisible by 32, and no adapters. The centered-code path checks contiguous groups and zero point 8 after the appropriate GPTQ zero-format interpretation.

Implementation	Role
w4a_floatx.py	W4AFP8Linear: FP8 hardware floor (8, 9); constructs an E4M3 runtime operand from centered INT4 codes; calls group-aware FP8 execution.
w4a_triton.py	Per-token FP8 packing, E4M3 rounding, grouped FP8 GEMM, and prepacked activation consumption. Launches bind to the operand device.
w4a_nvfp4.py	W4ANVFP4Linear: floor (12, 1); native FP4/FP8 operands, integer-plane encoding, scale swizzling, block-scaled GEMM, and mixed FP8 consumers.
w4a_nvfp4_triton.py	NVFP4 pack/swizzle/decode, supported recipe variants, and exact low/high INT4-plane accumulation with GPTQ group scales.
w4a_activation.py	Immutable W4AActivation carrier: mode, packed codes, scales, logical shape, model dtype, optional global/token scales, recipe, rotation flag, and exact-reference tensor.
w4a_boundary.py	Producer-boundary policy, mixed MLP precision policy, scale completeness checks, and NVFP4BoundaryQuantizer.
w4a_llama_replay.py	Installation/removal of quantization replay hooks; reference QDQ and precision matching for the operands that deployed consumers actually see.
w4a_llama_stream.py	Encoded stream installation on complete selected Llama decoder layers, residual/norm/attention/MLP boundaries, head transition, and data-parallel rebinding.

W4AActivation.to() moves codes and all scale/reference tensors together. decode() reconstructs the quantized hardware operand; exact() returns the source compute-dtype reference when present. Residual arithmetic, norms, and the unquantized head must not accidentally consume only a low-bit approximation when the stream policy requires the reference value. Packing scales independently of their codes, or dropping token_scale, changes semantics.

NVFP4 boundary scales must be finite positive FP32 values and match the exact expected producer key set. Runtime copies use INT32 bit buffers to survive model dtype conversion without rounding the FP32 calibration scalar. Replay, producer calibration, and runtime installation must use the same mixed-precision policy. The Llama helpers are model-structure-specific, not universal hooks for arbitrary architectures.

Current storage-policy limitation: both implementations expand centered INT4 codes into a full-layer E4M3 _weight_e4m3 buffer in post_init() and retain it across forwards. NVFP4 additionally retains its two packed FP4 weight planes in _weight_both. Marking these buffers persistent=False excludes them from the state dict; it does not release their runtime storage. The retained INT4-to-E4M3 expansion violates the repository's no-expanded-weight-cache rule, so these paths must not be described as passing the persistent-storage acceptance gate. Account for all retained operands rather than estimating resident memory from checkpoint INT4 bytes alone.

AWQ native and portable families

AWQ GEMM, GEMV, GEMV_FAST, and LLM_AWQ are different packing contracts. The model loader, class declaration, and native entry point must agree.

Family	Runtime behavior
AwqGEMMLinear	Native CUDA extension; rejects ROCm in validation. Its forward handles device/dtype/contiguity checks. The autograd wrapper has a dense-current-call branch for flattened row counts greater than 1024.
AwqGEMMTritonLinear	AWQ GEMM via Triton; checks versions and packed tensor co-location, converts work input to FP16, and restores the interface boundary.
AwqGEMVLinear	GEMV layout; uses the GEMMv2 forward for more than 8 flattened rows, otherwise GEMV.
AwqGEMVFastLinear / LLMAwqLinear	Fast GEMV and legacy LLM_AWQ layouts; runtime zero/scales preparation is layout-specific. Loader forces FP16 for these formats.
AwqTorchLinear	Portable packed GEMM fallback, including bit widths outside the native INT4 kernels' declarations.
AWQ Marlin/Machete/Swordfish/ExLlama/BitBLAS	Repacking/adaptation classes described above; not new AWQ quantizers.
AWQ fused/ATen/AVX512	CPU/XPU ABI adapters where their declarations/operators permit.

Native family declarations that mention ROCm still require the relevant runtime; the native CUDA AWQ GEMM extension specifically cannot build on ROCm. awq_gemm_pack_backend() therefore chooses a portable Torch packer for non-compatible environments and weight-only paths instead of attempting a CUDA-only packer everywhere. See AWQ.

QQQ, VPTQ, ParoQuant, and bitsandbytes

Family	Representation and execution
QQQ	QQQLinear uses QQQ-packed B, group/channel scales, dynamic activation quantization, and native multiply. It declares 4-bit symmetric weights, no activation ordering, group size -1 or 128, feature widths divisible by 64, and no training. QQQTorchLinear uses the same QQQ format with portable execution. Native availability/capability validation still applies.
VPTQ	VPTQLinear stores vector codebooks/packed indices plus configured residual/outlier state; Torch forward decodes the current weight for F.linear. VPTQCudaLinear uses a fused native path for eligible row counts/dtypes and checks that activation/indices/codebooks share a CUDA device; other work follows its fallback path. This is not scalar GPTQ packing.
ParoQuant	ParoLinear stores packed affine codes plus pairwise rotation/scaling metadata. It rotates inputs before quantized matmul and retains separate adapter semantics. Despite the PAROQUANT_CUDA name, its declaration is portable and includes device-dependent native/Torch/CPU handling. ParoQuantTritonLinear shares rotation semantics but selects Triton plans, has zero AUTO priority, and requires CUDA/Triton.
bitsandbytes	BitsAndBytesLinear represents 4-bit quantization state/codes or INT8 state, with CPU/CUDA declarations and state-aware execution paths. It has native/packed CPU and current-call fallback logic; a module name alone does not guarantee a native CUDA launch.

Use QQQ, VPTQ, ParoQuant, and BitsAndBytes for quantization-specific configuration and checkpoint details.

MLX checkpoint holders and executable modules

qlinear/mlx.py declares method/format-specific Torch checkpoint holders with priority zero and MPS device declarations. Their forward() raises until the model is converted into MLX modules. The holder hierarchy covers GPTQ/GPTQ_P, AWQ variants, MARLIN, BITBLAS, ParoQuant, QQQ, GGUF, FP8, and bitsandbytes; the loader also has EXL3-specific conversion.

Source	Executable MLX implementation
mlx_gptq.py	MlxGPTQLinear, MlxGPTQGroup16Linear; packed affine/group-16 execution and output cast to input dtype.
mlx_awq.py	MlxAWQLinear, MlxAWQGroup16Linear; AWQ affine/group-16 execution.
mlx_group16.py	MlxGroup16Linear; exact alternating half-group scales through two group-32 affine matmuls; validates input alignment and supported affine bits.
mlx_exl3.py	MlxEXL3Linear; packed trellis matmul and block Hadamard transforms without a dense-weight fallback.
mlx_fp8.py	MlxFP8Linear / MlxFP8PackedLinear; FP8 scales/codes and packed Metal paths.
mlx_gguf.py	MlxGGUFLinear / MlxGGUFQ6KLinear; affine GGUF and specialized packed Q6_K execution.
mlx_paro.py	MlxParoLinear; pairwise rotations before quantized matmul.
mlx_qqq.py	MlxQQQLinear; per-token activation and per-channel output-scale semantics.
mlx_bitsandbytes.py	MlxBitsAndBytesLinear; packed 4-bit / INT8 execution and exact signed-to-affine INT8 repacking helper.

An MLX implementation is not a Torch MPS kernel. Install/use the MLX-compatible model path and verify the method-specific conversion rather than calling a Torch holder directly. See MLX and Apple Silicon.

Packing utilities and extension packers

Stable layer APIs

def hf_pack_layer(
    module: BaseQuantLinear,
    linear: nn.Module,
    scales: torch.Tensor,
    zeros: Optional[torch.Tensor] = None,
    g_idx: Optional[torch.Tensor] = None,
    *,
    pack_impl: Optional[str] = None,
    device=None,
    block_in: int = 8192,
    workers: int = 1,
    scales_extra: Optional[torch.Tensor] = None,
    post_init: bool = False,
    post_init_kwargs: Optional[Dict[str, Any]] = None,
    checkpoint_format: Optional[Union[str, FORMAT]] = None,
) -> BaseQuantLinear:
    ...

The scales/zero inputs use quantizer-native [out_features, num_groups] layouts. For symmetric modules, omitted zeros can be synthesized; asymmetric modules require zeros. pack_impl=None means original packing; auto prefers GPU packing where implemented, then block, then original. QQQ/AWQ single-packer branches ignore this implementation choice. GPU packing returns CPU-registered packed buffers; final module placement is separate.

def hf_repack_layer(
    module: BaseQuantLinear,
    backend: Optional[Union[str, BACKEND]] = None,
    *,
    device=None,
    dtype: Optional[torch.dtype] = None,
    post_init: bool = True,
    post_init_kwargs: Optional[Dict[str, Any]] = None,
) -> BaseQuantLinear:
    ...

Repacking selects the destination kernel by source method/format, then uses a native repack hook or copies compatible checkpoint tensors and performs destination initialization. Sources already converted to certain private runtime ABIs are rejected. An initialized runtime matrix is not necessarily a valid checkpoint-layout source.

hf_check_packing_feasibility(), hf_check_best_packing_format(), and hf_post_init() support integration checks and initialization. hf_select_quant_linear and its v2 wrapper are in utils/importer.py; converter.py currently has an empty MODULE_CONVERTER_MAP and is not a packing dispatcher.

Planar packing

utils/planar_packing.py implements GPTQ_P split-plane packing. It requires int32 words and aligned groups of 32 logical codes for its row/column helpers.

Logical bits	Plane widths, low to high
2	2
3	2 + 1
4	4
5	4 + 1
6	4 + 2
7	4 + 2 + 1
8	8

Each plane width divides 32; a field never straddles a word boundary. For every 32 codes, the plane words are adjacent in low-to-high order. Single-plane 2/4/8-bit layouts are bit-identical to continuous packing; planar 3-bit is not the same as the continuous cross-word 3-bit layout. The helpers are planar_pack_rows, planar_unpack_rows, planar_pack_cols, and planar_unpack_cols.

CPU block packing

qlinear/pack_block_ext.py::pack_block_cpu loads the block-packing extension and calls torch.ops.gptqmodel.pack_block_cpu with weight, scales, zeros, g_idx, bit width, word width, input block size, and worker count. Its native source is gptqmodel_ext/pack_block_cpu.cpp. This wrapper raises if the extension cannot load; higher-level pack selection determines whether an alternative is available.

Hooked modules during quantization

gptqmodel/nn_modules/hooked_linear.py provides HookedLinear plus Conv1D/Conv1d/Conv2d-compatible wrappers, StopForward, and legacy/tree-based replacement helpers. These intercept forwards for quantization processing and can materialize meta weights through the shell/lazy source and restore output placement. They are not entries in the inference BACKEND enum and do not replace a qlinear kernel's packing contract.

See Model Definitions and Module Tree, Looper Processors, and Memory, Offload, and Multi-GPU for their place in the quantization lifecycle.

Native C++/CUDA source map

gptqmodel_ext/__init__.py is only a support package; its empty __all__ is not the runtime kernel registry. Python loaders/utilities bind the relevant extension and torch.ops entry points.

Source path under gptqmodel_ext/	Role
marlin/	GPTQ/AWQ repack kernels; FP16/BF16 GEMM sources and Torch bindings; shared kernel/core code.
machete/	CUTLASS Machete core and PyTorch CUDA interface; generation/build driven by utils/machete.py.
swordfish/	Prepacked Blackwell implementation, bindings, core, and stable-libtorch support.
exllamav2/	Native GPTQ/AWQ matrix creation, GEMM bindings, CUDA/CPU support code.
exllamav3/	EXL3 bindings, Hadamard, GEMM, trellis/quantization native implementation.
awq/	AWQ Torch bindings, GEMM/GEMV/fast quantization kernel sources.
qqq/	QQQ binding and quantized GEMM source.
vptq/	Native vector-quantized inference binding/kernel.
paroquant/	Native rotation implementation.
int4_qzero_group_scaled_cpu/	Native INT4 packing/GEMM, AVX512 implementation, QQQ CPU support, and binding code.
pack_block_cpu.cpp	Blockwise GPTQ CPU packing.
floatx_cpu.cpp	Native packed/floatx CPU support.
gptq_block/	Native GPTQ block-processing implementation used during quantization, not a selectable inference backend.
hadamard/	Hadamard transform bindings and CUDA implementation.
cutlass_extensions/	Shared CUTLASS support utilities.

Architecture-specific source compilation requires an appropriate Torch build, compiler/toolkit, headers, and the extension's dependencies. A successful Python import cannot substitute for an exact-device capability check. Do not copy one extension's build flags into another backend merely because both use CUDA.

ROCm quantization BLAS is not an inference BACKEND

utils/rocm.py::quant_blas_backend() scopes the preferred Torch BLAS library during quantization and restores the previous selection afterward on a best-effort basis. The default request is "rocblas", translated to Torch's "hipblas" spelling. It only attempts the switch on ROCm when the preferred-BLAS API exists and the previous selection is hipBLASLt. An already selected different library is preserved.

Environment setting	Effect
Unset GPTQMODEL_ROCM_QUANT_BLAS	Request the scoped rocBLAS selection when the runtime checks permit.
"hipblaslt", "cublaslt", "default", "off", "0", or an empty string	Disable this override.
Another non-disabled value	Pass the requested backend spelling through the helper's alias map/API, with failure leaving the run on its current path.

docs/rocm-quant-blas-backend.md explains the original gfx950 cold-start investigation. Its measurements are hardware/run-specific, not a benchmark of the inference kernels above. BLAS selection can change floating-point rounding and quantization results; it is not an automatically bit-exact change. Setting this environment variable does not make a CUDA-only Marlin, Machete, Swordfish, or AWQ GEMM extension ROCm-compatible.

Kernel implementation and validation checklist

1. Declare method, backend, format/priority, bit widths, groups, desc_act, symmetry, training, shards, devices/platforms, packing/activation dtypes, divisibility, and adapters accurately. A positive priority opts the kernel into AUTO.
2. Keep exact-device validation separate from process-global extension availability. Multi-GPU selection checks every target ordinal; do not freeze the current GPU's capability in an import cache.
3. Preserve checkpoint zero-point conventions and layout semantics. Validate raw packing separately from loading/repacking; do not claim pack support solely because post_init() can transform loaded tensors.
4. Keep weights quantized between operations. Temporary decode is allowed for the consuming operation; dense or expanded caches must not persist. Nonpersistent Torch buffers still consume memory and are not an exemption.
5. Invalidate caches/handles on state changes, placement/dtype transitions, or mutation as required. _runtime_cache.tensor_cache_key() covers identity/version/device/dtype/shape, but inference-mode mutation can require explicit invalidation.
6. Follow the repository's independent Torch-oracle and persistent-storage requirements for kernel/packer/math changes. Test supported FP16/BF16 paths independently, check output dtype, separate arithmetic error from output-rounding error, and include repeated forwards/mutation/lifetime cases rather than one successful call.

Pitfalls

• Do not equate enum presence, format declaration, extension importability, or a broad logical-bit list with successful end-to-end execution.
• Do not save private runtime packed tensors under a standard checkpoint format without restoring its serialization contract.
• Do not assume a training-support declaration means the inference-native operator supplies backward; classes may dispatch differently in training.
• Do not confuse CUDA/HIP's shared Torch device namespace with binary compatibility of native CUDA kernels.
• Do not benchmark JIT compilation/repacking as steady-state GEMM, or present a declaration priority as a measured performance ranking.
• Do not estimate peak memory from on-disk weight bytes; repacking, native workspaces, scratch, activations, and temporary decode can coexist.

Source files

• pyproject.toml
• gptqmodel/models/loader.py
• gptqmodel/nn_modules/qlinear
• gptqmodel/nn_modules/qlinear/__init__.py
• gptqmodel/nn_modules/qlinear/_runtime_cache.py
• gptqmodel/nn_modules/triton_utils
• gptqmodel/nn_modules/exllamav3.py
• gptqmodel/nn_modules/exllamav3_torch.py
• gptqmodel/nn_modules/hooked_linear.py
• gptqmodel/utils/importer.py
• gptqmodel/utils/packing.py
• gptqmodel/utils/planar_packing.py
• gptqmodel/utils/marlin.py
• gptqmodel/utils/marlin_scratch.py
• gptqmodel/utils/machete.py
• gptqmodel/utils/swordfish.py
• gptqmodel/utils/bitblas.py
• gptqmodel/utils/rocm.py
• gptqmodel/nn_modules/qlinear/bitblas_target_detector.py
• gptqmodel_ext
• docs/marlin_scratch.md
• docs/torch_fused_int4_transformations.md
• docs/rocm-quant-blas-backend.md

See also

• Inference Backends
• Checkpoint Formats and Compatibility
• Developer Guide
• FP8 and NVFP4
• AWQ
• EXL3
• GGUF
• MLX and Apple Silicon
• Memory, Offload, and Multi-GPU