Repository navigation
Inference Backends
Documentation home · Quickstart · Lifecycle
On this page
- Overview
- Workflow at a glance
- Concepts
- Loading and selecting a backend
- Complete BACKEND enum reference
- Hardware, dependency, and format requirements
- How AUTO selection works
- Packing, conversion, and runtime initialization
- Numeric epilogues and input gradients
- Strided tensor operands
- Runtime cache lifetime
- Pitfalls and diagnosis
- Source files
- See also
GPT-QModel separates the quantization method, the checkpoint format, and the inference backend. A method describes how the weights were quantized; a format describes the tensors stored on disk; a backend selects the implementation that consumes those tensors. Choosing a different backend does not, by itself, turn an AWQ checkpoint into GPTQ or change its quantization quality.
This reference describes the source at commit e17a1a29bf0e738880faf6883ee6babb0b1f7b86. Hardware gates below are implementation requirements, not general claims about what an architecture could theoretically execute. In particular, Machete is Hopper-only, Swordfish checks specific Blackwell targets, and W4A NVFP4 has its own different capability floor.
Use BACKEND.AUTO for normal local loading unless a particular kernel is required. AUTO selects the first compatible kernel in a declaration-derived ordering. It does not benchmark the model or guarantee the fastest backend for every batch size, prompt length, or device.
Reading the diagram: Alternative routes. Format-specific loader branches can precede local selection; AUTO is priority-based, not benchmarking. MLX has its own conversion route.
Text version of the workflow
- Local AUTO: Method and format → Rank declarations → Validate candidates → Select a valid class.
- Explicit local backend: Named BACKEND → Validate selection → Initialize or reject.
- vLLM or SGLang: Engine requested → Check engine format → Engine initialization → Engine generation.
| Concept | Source representation | Meaning |
|---|---|---|
| Quantization method | METHOD |
GPTQ, AWQ, QQQ, and other method families used to interpret quantization configuration. |
| Checkpoint format | FORMAT |
The serialization contract: for example GPTQ, GPTQ_V2, GPTQ_P, GEMM, MARLIN, GGUF, or EXL3. |
| Backend |
BACKEND, in gptqmodel/utils/backend.py
|
A local qlinear implementation or a model-level engine/conversion path. |
| Kernel declaration |
SUPPORTS_BACKENDS, SUPPORTS_METHODS, SUPPORTS_FORMATS
|
Connects a class to supported method/format/backend combinations; format values also carry AUTO priorities. |
| Runtime packing | Kernel post_init(), repack_from_gptq(), or repack_from_awq()
|
Converts loaded packed tensors into a kernel's private execution layout where required. |
| Profile |
PROFILE.AUTO = "auto", FAST = "fast", LOW_MEMORY = "low_memory"
|
A separate loading/runtime strategy selector, not a checkpoint format or backend name. |
Portable device declarations are still subject to operator and dtype availability. DEVICE.ALL means the class declares a general implementation; it is not evidence that every PyTorch operation used by that class is implemented on every installed accelerator runtime.
Hardware families are explicit: DEVICE.CUDA declares NVIDIA support and DEVICE.ROCM declares AMD support. On a HIP build, normalize_device() resolves "cuda", "cuda:N", Torch CUDA devices, integer accelerator ordinals, and even DEVICE.CUDA to DEVICE.ROCM; automatic selection and quantization config retain that AMD family. Runtime placement translates it back to Torch's cuda namespace only at execution boundaries, preserving concrete ordinals and current-device behavior for indexless families. Neither torch.cuda.is_available(), a synthetic HIP capability tuple, nor Triton availability enables a CUDA-only kernel on AMD. In particular, GPTQ_TRITON (TritonV2Linear) is NVIDIA-only; portable GPTQ_TORCH has separate eligible ROCm Triton fast paths.
In particular, native AWQ and QQQ declarations that mention ROCm are not proof of CUDA/PTX kernel portability. Check the method-specific AWQ and QQQ runtime caveats before choosing a native backend.
from gptqmodel import BACKEND, GPTQModel
# Replace with a compatible, already-quantized checkpoint.
model = GPTQModel.load("/models/my-gptq-checkpoint", backend=BACKEND.AUTO)
# Force the portable GPTQ implementation when diagnosing a native-kernel issue.
reference_model = GPTQModel.load(
"/models/my-gptq-checkpoint",
backend=BACKEND.GPTQ_TORCH,
device="cuda:0",
)The model-level loading signature is:
def load(
cls,
model_id_or_path: Optional[str],
quantize_config: Optional[QuantizeConfig | Dict] = None,
device_map: Optional[Union[str, Dict[str, Union[str, int]]]] = None,
device: Optional[Union[str, torch.device]] = None,
backend: Union[str, BACKEND] = BACKEND.AUTO,
profile: Union[str, int, PROFILE] = PROFILE.AUTO,
trust_remote_code: bool = False,
allow_remote_adapter: bool = False,
adapter_revision: Optional[str] = None,
adapter_sha256: Optional[str] = None,
**kwargs,
):
...String names can be normalized by normalize_backend. Prefer method-specific enum members in new code: GPTQ_TORCH or AWQ_TORCH is less ambiguous than the legacy TORCH spelling. In generic local qlinear selection, explicit selection is a requirement, not a hint: a failed hardware, format, dtype, or dependency check raises an error instead of silently choosing a different backend. Model-level loader branches can take a different path before this selector runs; the native Q1_0 GGUF FAST profile described below is one such exception.
For an AWQ GEMM checkpoint, the portable alternative is BACKEND.AWQ_TORCH, not GPTQ_TORCH. For GGUF, EXL3, FP8, VPTQ, QQQ, or bitsandbytes, use the matching family in the tables below. See Loading and Inference for generation and Packing Formats and Compatibility for serialization.
The following two tables together include every enum member. Values are exact strings from backend.py. Format numbers shown beside local classes are their declared AUTO priorities, not bit widths or version numbers. A priority of zero means explicit selection remains possible but that format/backend combination does not participate in generic AUTO selection. Device entries describe declarations; additional runtime gates are documented afterward.
| Member | String value | Dispatch / declared checkpoint support and device families |
|---|---|---|
AUTO |
auto |
Local candidate selection in descending declared format priority; then validation. |
AUTO_TRAINABLE |
auto_trainable |
AUTO plus the trainable validation requirement. |
GPTQ_TORCH_FUSED |
gptq_torch_fused |
TorchFusedLinear: GPTQ: 50, GPTQ_V2: 50; CPU, XPU
|
GPTQ_TORCH_INT8 |
gptq_torch_int8 |
TorchInt8Linear: GPTQ: 0, GPTQ_V2: 0; CPU
|
GPTQ_TORCH |
gptq_torch |
TorchLinear: GPTQ: 20, GPTQ_V2: 20, GPTQ_P: 20; ALL. TorchQuantEmbeddings shares these declarations but is role-specific and excluded from general backend selection. |
GPTQ_TRITON |
gptq_triton |
TritonV2Linear: GPTQ: 40, GPTQ_V2: 40; CUDA
|
BITSANDBYTES |
bitsandbytes |
BitsAndBytesLinear: BITSANDBYTES: 40; CPU, CUDA, ROCM
|
GPTQ_EXLLAMA_V2 |
gptq_exllama_v2 |
ExllamaV2Linear: GPTQ: 80, GPTQ_V2: 80; CUDA
|
MACHETE |
machete |
GPTQ MacheteLinear: GPTQ: 100; AWQ AwqMacheteLinear: GEMM: 100, MARLIN: 100; CUDA
|
SWORDFISH |
swordfish |
GPTQ SwordfishLinear: GPTQ: 101; AWQ AwqSwordfishLinear: GEMM: 101; CUDA
|
MARLIN |
marlin |
GPTQ MarlinLinear: GPTQ: 90, GPTQ_V2: 90, MARLIN: 90; AWQ AwqMarlinLinear: GEMM: 90, MARLIN: 90; CUDA
|
ALLSPARK |
allspark |
GPTQ AllSparkLinear: GPTQ: 0, GPTQ_V2: 0; AWQ AllSparkAwqLinear: GEMM: 0; CUDA, ROCM; explicit-only W8A16 |
GPTQ_W4AFP8 |
gptq_w4afp8 |
W4AFP8Linear: GPTQ: 0, GPTQ_V2: 0; CUDA
|
GPTQ_W4A_NVFP4 |
gptq_w4a_nvfp4 |
W4ANVFP4Linear: GPTQ: 0, GPTQ_V2: 0; CUDA
|
GPTQ_BITBLAS |
gptq_bitblas |
BitblasLinear: BITBLAS: 30, GPTQ: 30, GPTQ_V2: 30; CUDA
|
GPTQ_TORCH_ATEN |
gptq_torch_aten |
TorchAtenLinear: GPTQ: 110, GPTQ_V2: 110; CPU
|
GPTQ_TORCH_CPU_AVX512 |
gptq_torch_cpu_avx512 |
TorchCpuAvx512Linear: GPTQ: 120, GPTQ_V2: 120; CPU
|
QQQ |
qqq |
QQQLinear: QQQ: 100; CUDA, ROCM
|
QQQ_TORCH |
qqq_torch |
QQQTorchLinear: QQQ: 90; ALL
|
AWQ_GEMM |
awq_gemm |
AwqGEMMLinear: GEMM: 60; CUDA
|
AWQ_GEMM_TRITON |
awq_gemm_triton |
AwqGEMMTritonLinear: GEMM: 50; CUDA, ROCM
|
AWQ_GEMV |
awq_gemv |
AwqGEMVLinear: GEMV: 40; CUDA
|
AWQ_GEMV_TRITON |
awq_gemv_triton |
AwqGEMVTritonLinear: GEMV: 0; CUDA, ROCM; explicit-only, Linux, FP16 |
AWQ_GEMV_FAST |
awq_gemv_fast |
AwqGEMVFastLinear: GEMV_FAST: 30; CUDA; LLMAwqLinear: LLM_AWQ: 100; CUDA
|
AWQ_TORCH_INT8 |
awq_torch_int8 |
TorchInt8AwqLinear: GEMM: 0; CPU
|
AWQ_TORCH_FUSED |
awq_torch_fused |
TorchFusedAwqLinear: GEMM: 20; CPU, XPU
|
AWQ_TORCH_ATEN |
awq_torch_aten |
TorchAtenAwqLinear: GEMM: 110; CPU
|
AWQ_TORCH_CPU_AVX512 |
awq_torch_cpu_avx512 |
TorchCpuAvx512AwqLinear: GEMM: 120; CPU
|
AWQ_TORCH |
awq_torch |
AwqTorchLinear: GEMM: 10; ALL
|
AWQ_BITBLAS |
awq_bitblas |
AWQBitBlasKernel: GEMM: 0, BITBLAS: 30; CUDA
|
AWQ_EXLLAMA_V2 |
awq_exllama_v2 |
AwqExllamaV2Linear: GEMM: 80; CUDA
|
PAROQUANT_CUDA |
paroquant_cuda |
ParoLinear: PAROQUANT: 55; ALL
|
PAROQUANT_TRITON |
paroquant_triton |
ParoQuantTritonLinear: PAROQUANT: 0; CUDA, ROCM
|
FP8_TORCH |
fp8_torch |
TorchFP8Linear: FP8: 15; CPU, CUDA, ROCM, XPU, MPS
|
NF12_AD |
nf12_ad |
NF12ADLinear: NF12_AD (nf12): 100; ALL. BF16 native CUDA/RDNA4 HIP or portable codec/matmul fallback. See NF12 for geometry and dependencies. |
VPTQ_TORCH |
vptq_torch |
VPTQLinear: VPTQ: 100; ALL
|
VPTQ_CUDA |
vptq_cuda |
VPTQCudaLinear: VPTQ: 100; CUDA
|
GGUF_TORCH |
gguf_torch |
GGUFTorchLinear: GGUF: 15; ALL
|
GGUF_TRITON |
gguf_triton |
GGUFTritonKernel: GGUF: 45; CUDA, ROCM
|
GGUF_CPP_CPU |
gguf_cpp_cpu |
GGUFCppKernel: GGUF: 25; CPU
|
GGUF_CPP_CUDA |
gguf_cpp_cuda |
GGUFCudaKernel: GGUF: 35; CUDA
|
EXL3_EXLLAMA_V3 |
exl3_exllama_v3 |
EXL3 loader branch; ExllamaV3Linear; NVIDIA CUDA only. |
EXL3_TORCH |
exl3_torch |
EXL3 loader branch; ExllamaV3TorchLinear; portable Torch implementation. |
VLLM |
vllm |
External vLLM loader, not a local qlinear class. |
SGLANG |
sglang |
External SGLang loader, not a local qlinear class. |
MLX |
mlx |
Explicit MLX conversion path; see the holder table in Kernel Backends Reference. |
These are distinct enum values retained for old call sites and saved arguments, not Python enum aliases with the same string value. _LEGACY_BACKEND_BY_METHOD rewrites them only when the corresponding quantization method is known. A legacy member not mapped for a method is not automatically a valid backend for that method.
| Legacy member | String value | Method-dependent normalization |
|---|---|---|
TORCH_FUSED |
torch_fused |
gptq → GPTQ_TORCH_FUSED; awq → AWQ_TORCH_FUSED
|
TORCH_INT8 |
torch_int8 |
gptq → GPTQ_TORCH_INT8; awq → AWQ_TORCH_INT8
|
TORCH |
torch |
gptq → GPTQ_TORCH; awq → AWQ_TORCH; fp8 → FP8_TORCH; exl3 → EXL3_TORCH; vptq → VPTQ_TORCH; nf12 → NF12_AD
|
TRITON |
triton |
gptq → GPTQ_TRITON
|
EXLLAMA_V2 |
exllama_v2 |
gptq → GPTQ_EXLLAMA_V2; awq → AWQ_EXLLAMA_V2
|
EXLLAMA_V3 |
exllama_v3 |
exl3 → EXL3_EXLLAMA_V3
|
BITBLAS |
bitblas |
gptq → GPTQ_BITBLAS; awq → AWQ_BITBLAS
|
GEMM |
gemm |
awq → AWQ_GEMM
|
GEMM_TRITON |
gemm_triton |
awq → AWQ_GEMM_TRITON
|
GEMV |
gemv |
awq → AWQ_GEMV
|
GEMV_TRITON |
gemv_triton |
awq → AWQ_GEMV_TRITON
|
GEMV_FAST |
gemv_fast |
awq → AWQ_GEMV_FAST
|
TORCH_INT8_AWQ |
torch_int8_awq |
awq → AWQ_TORCH_INT8
|
TORCH_FUSED_AWQ |
torch_fused_awq |
awq → AWQ_TORCH_FUSED
|
TORCH_AWQ |
torch_awq |
awq → AWQ_TORCH
|
BITBLAS_AWQ |
bitblas_awq |
awq → AWQ_BITBLAS
|
PARO |
paroquant |
paroquant → PAROQUANT_CUDA
|
MARLIN, MACHETE, SWORDFISH, and ALLSPARK are shared canonical backends: method and checkpoint format choose the GPTQ or AWQ sub-kernel. Old strings gptq_marlin/awq_marlin, gptq_machete/awq_machete, gptq_swordfish/awq_swordfish, and gptq_allspark/awq_allspark normalize to the corresponding shared backend. They are string compatibility aliases, not method-qualified enum members; use BACKEND.MARLIN, for example, rather than the removed BACKEND.GPTQ_MARLIN.
normalize_backend(None) returns None; a blank or whitespace-only string also returns None. The selector resolves the absence of an explicit backend to AUTO. Nonempty strings are stripped and can match enum values case-insensitively, member names, or the shared-backend string aliases above. Unknown backend names raise ValueError; non-string, non-enum inputs raise TypeError. Method-aware normalization is why a legacy "torch" can mean GPTQ Torch, AWQ Torch, FP8 Torch, EXL3 Torch, or VPTQ Torch rather than one universal kernel.
| Family | Actual requirement / important restriction |
|---|---|
| GPTQ Marlin | CUDA, not ROCm; compute capability 7.5 or newer; extension availability and per-layer shape validation. SM75 has BF16 handling distinct from native BF16-capable devices. Bits 4 and 8 are declared. |
| AWQ Marlin | CUDA, not ROCm; compute capability 8.0 or newer; Marlin/AWQ shape, type, and packing checks still apply. |
| Machete | CUDA, not ROCm; capability exactly 9.0; at least 204800 bytes of opt-in shared memory per block; generated CUTLASS kernels target sm_90a. A newer Blackwell GPU is not accepted as a Hopper substitute. The helper pins CUTLASS 4.7.1 and verifies its source archive. |
| Swordfish | CUDA, not ROCm; supported capability targets derived from its build flags include 10.0, 10.3, and 11.0, rather than every device with capability at least 10.0. The runtime checks the Torch stable ABI requirement as well as extension availability. GPTQ supports 4/8-bit; AWQ supports 4-bit. Both require input/output widths divisible by 64 and reject sharded loading. |
| ExLlama v2 | The GPTQ/AWQ classes declare NVIDIA CUDA and 4-bit weights with FP16/BF16/FP32 activation dtypes. Validation, availability helpers, and registry loading reject ROCm before attempting JIT compilation, with guidance to use AUTO or the method-specific Torch backend (TORCH for GPTQ, AWQ_TORCH for AWQ). Native GEMM preserves activation/output dtype instead of unconditionally narrowing to FP16, and uses FP32 accumulation. Scales remain FP16 checkpoint/runtime metadata. Native extension availability, dimension constraints, and scratch initialization are required. AWQ has its own GEMM-to-ExLlama adaptation class. |
| BitBLAS | CUDA with BitBLAS/TVM runtime support; GPTQ declares 2/4/8-bit and AWQ declares 4-bit. Transformed packing and target detection are part of the runtime, not an interchangeable GPTQ tensor layout. |
| AllSpark | Explicit BACKEND.ALLSPARK / "allspark"; W8A16 only (bits=8, group_size=-1), int32 packing, K/N divisible by 16, FP16/BF16 activations. GPTQ requires sym=True, desc_act=False; AWQ consumes GEMM layout and permits symmetric/asymmetric zeros. CUDA gate is compute capability 8.0–12.0 inclusive; the separate ROCm path accepts only gfx1200/gfx1201, not all RDNA4/CDNA devices. No sharded loading, training, or adapters. GPTQ is load/repack-only, not a raw quantization packer. |
| Native AWQ GEMM | Native CUDA AWQ extension. validate_once() explicitly rejects ROCm even though several other AWQ implementations have ROCm declarations. |
| W4A FP8 | NVIDIA only; ROCm is rejected before compilation/launch. W4AFP8Linear.MIN_COMPUTE_CAPABILITY = (8, 9); Triton with the Gluon mma_v2 FP8 GEMM (first shipped in Triton 3.6.0). The class still imports and validates on older Triton builds — activation packing works — but the GEMM launch raises native_fp8_gemm_error(). This is more precise than the broad SM12x comment beside the backend enum. |
| W4A NVFP4 | NVIDIA only; ROCm is rejected before compilation/launch. W4ANVFP4Linear.MIN_COMPUTE_CAPABILITY = (12, 1); native torch.float4_e2m1fn_x2, torch.nn.functional.scaled_mm, ScalingType, SwizzleType, and the activation-packing Triton module (plain tl.jit — no Gluon needed for pure NVFP4; a mixed recipe's FP8 attention lane reports the Triton 3.6.0 floor when it runs). Do not infer support from Swordfish's SM100 support. |
Extension importability is not proof that the selected device is supported. Marlin, Machete, and Swordfish retain device-sensitive checks; the selector preserves exact GPU ordinals so the currently active CUDA device cannot stand in for another GPU in a device map.
AllSpark keeps the original packed checkpoint tensors registered after CUDA/HIP repacking, so state_dict() and get_state_dict_for_save() remain exportable after inference. Native packed buffers are non-persistent derived state and rebuild after checkpoint loads, device/dtype moves, or ordinary versioned source changes. Call clear_weight_cache() after unversioned inference-tensor edits. Bias joins the FP32 accumulator before the final FP16/BF16 cast on both device paths, and FP32 checkpoint bias remains FP32 through loading and module dtype moves. See the checkpoint-layout contract and bias epilogue.
Device-validation errors retain a kernel's specific NotImplementedError reason, for example an unsupported compute capability, missing kernel image, or CUDA-only implementation on HIP. A generic family mismatch is not duplicated; on ROCm it gains portable-backend guidance. Swordfish, Machete, and BitBLAS explicitly reject HIP with method-appropriate alternatives. Read the full rejection reason before changing checkpoint metadata or rebuilding extensions; this improves diagnostics, not hardware support.
| Target | Selection considerations |
|---|---|
| CPU INT4, AVX512-BF16 |
GPTQ_TORCH_CPU_AVX512 / AWQ_TORCH_CPU_AVX512 use the native int4_qzero_group_scaled_cpu implementation and require AVX512-BF16 support. An AVX512-branded CPU without the required BF16 feature is not enough. |
| CPU ATen INT4 |
GPTQ_TORCH_ATEN / AWQ_TORCH_ATEN require _convert_weight_to_int4pack_for_cpu and _weight_int4pack_mm_for_cpu. GPTQ group sizes are 32, 64, 128, or 256; the backend can validate on hosts that lack the specialized AVX512 path. |
| CPU Torch fused | The fused family declares CPU and XPU. Availability depends on the installed Torch/torchao/operator path; its runtime layout differs from checkpoint GPTQ/AWQ. |
| CPU native 8-bit |
GPTQ_TORCH_INT8 is explicit-only and declares 8-bit, not 4-bit-to-8-bit expansion. It uses _weight_int8pack_mm; raw packing is disabled. |
| Retired AWQ INT8 |
AWQ_TORCH_INT8 remains in the enum only to issue an actionable migration error. Its constructor and validate_once() reject use; select AUTO or AWQ_TORCH. |
| ROCm AllSpark W8A16 | Explicit ALLSPARK has separate HIP kernels for gfx1200/gfx1201 only (not gfx1250 or gfx950). Each selected target is validated; an indexless family selector checks every visible accelerator, not just the current GPU. On mixed AMD architectures, select an eligible indexed device explicitly. Warmup runs on the selected device. Use 8-bit per-channel GPTQ/GPTQ_V2 or AWQ GEMM with the AllSpark constraints above; the loader chooses HIP-specific repacking rather than CUDA's N32K16 layout. |
| Intel XPU | Torch fused is an explicit declared XPU candidate; generic Torch can provide a portable path. Do not read XPU-aware code in a module as proof that the module's selector declaration accepts XPU. The full declaration table is in Kernel Backends Reference. |
| ROCm | PyTorch uses the cuda device namespace for HIP; selection distinguishes DEVICE.ROCM, and ROCm >= 10.1 is the supported floor (older runtimes log an upgrade warning). gptqmodel/utils/arch.py classifies the device gcnArchName into cdna4 (gfx950) and rdna4 (gfx1200/gfx1201/gfx1250); only those registered families enable the fused GPTQ Triton GEMV/small-M GEMM, each with a per-architecture launch table (GPTQ_TORCH_ROCM_GEMV/GPTQ_TORCH_ROCM_GEMM opt-outs). CUDA-only native Marlin, Machete, Swordfish, AWQ GEMV/GEMV_FAST/LLM_AWQ, native QQQ, and VPTQ_CUDA do not become HIP kernels — native AWQ GEMV-family requests and QQQ raise NotImplementedError naming the portable alternative. Generic Torch, AWQ GEMM Triton, explicit AWQ GEMV Triton for compatible GEMV layouts, ParoQuant Triton (explicit selection only), GGUF_TRITON, GGUF_TORCH, VPTQ_TORCH, QQQ_TORCH, FP8 Torch, and EXL3_TORCH remain candidates where their own checks pass. BitsAndBytes also uses the HIP CUDA-facing device namespace, with eligible RDNA4 four-bit calls accelerated by Triton. |
| Apple Silicon / MLX | MLX is a model conversion/execution path with checkpoint-holder classes, not an ordinary Torch MPS qlinear forward. Holders must be converted into method-specific MLX modules before inference. Loader-level AUTO may choose MLX separately from the generic zero-priority holder declarations. |
docs/rocm-quant-blas-backend.md describes quantization-time BLAS selection, not a new inference backend. quant_blas_backend() and GPTQMODEL_ROCM_QUANT_BLAS control the hipBLASLt/rocBLAS quantization context; rocBLAS is the default on both CDNA and RDNA4 (RDNA4's hipBLASLt Tensile libraries exist but are untuned, so GPTQMODEL_ROCM_QUANT_BLAS=hipblaslt remains an explicit experiment). Changing that setting does not enable a CUDA-only inference extension on ROCm. The per-backend RDNA4 support matrix and its measurements live in docs/rocm-rdna4.md.
The tuned GPTQ GEMV/GEMM/dequant, AWQ GEMV, and ParoQuant GEMM helpers have separate Ada (SM 8.9) and RDNA4 launch tables. tuning_family() in utils/arch.py separates NVIDIA compute capability from AMD gfx architecture; this does not make CUDA-only native extensions portable.
-
GPTQ_TORCHcan use fused Triton GEMV on Ada and registered ROCm families. Ada's small-M GEMM is gated to 2–16 rows, K <= 2048, N <= 2048, FP16/BF16, plus the packed-word/alignment/no-gradient checks. RDNA4 keeps its FP16-only small-M GEMM gate.GPTQ_TORCH_ROCM_GEMVandGPTQ_TORCH_ROCM_GEMMstill default to1and now also control Ada's paths;GPTQ_TORCH_TRITON_DEQUANT=0disables their shared prerequisite. - AWQ GEMM
autoscheduling uses exact measured RTX 4090 or gfx1201 R9700 device/shape/FP16 rules, not a blanket family-wide performance threshold.AWQ_GEMV_TRITONis a separate, explicit-only GEMV-layout backend optimized for decode; see AWQ. -
GGUF_TRITONnow validates on ROCm with generic/autotuned launches and can win AUTO selection (priority 45 vs Torch 15). NVIDIA SM 8.9 fixed Q4_K/Q5_K/Q6_K launches remain capability/shape-gated; they never run on HIP. - BitsAndBytes NF4/FP4 uses fused Triton by default on eligible RDNA4 calls with M >= 16 and FP16/BF16 compute. Ada keeps the vendor default;
GPTQMODEL_BNB_TRITON=1enables forced A/B use without removing row/dtype gates. See BitsAndBytes.
These are dispatch policies, not promises of equal speed or coverage. Source-pinned hardware parity and RDNA4 measurements describe the tested configurations.
The GPTQ Triton v2, AWQ GEMM Triton, ParoQuant Triton, and GGUF Triton validators check Triton availability/version. Their shared minimum is 2.0.0; on free-threaded Python with the GIL disabled, the validators require 3.4.0 or newer. W4A validators additionally check the particular modules/operators needed by their FP8/FP4 paths; separately, the grouped FP8 GEMM needs the Gluon mma_v2 API introduced in Triton 3.6.0, checked at launch rather than at import. The name TritonV2Linear does not pin installation to Triton 2.x.
The package declares Python 3.10 or newer. Optional dependency requirements in pyproject.toml are separate from the minimum checks embedded in individual kernel validators; passing an older validator threshold does not mean it satisfies the current installation extra.
| Optional extra | Requirement declared at this revision |
|---|---|
triton |
triton>=3.8.0 when Python is below 3.15. The extra has no applicable Triton requirement on Python 3.15 or newer; that is not a runtime-support guarantee. |
bitblas |
bitblas==0.1.0.post1; the module's validator minimum is 0.1.0.post1. |
bitsandbytes |
bitsandbytes>=0.50.2. |
gguf |
gguf>=0.19.0 and llama-cpp-python>=0.3.36. |
mlx |
mlx>=0.32.4; mlx_lm>=0.31.3 below Python 3.11, otherwise mlx_lm>=0.32.2. Apple Silicon/Metal validation still applies. |
vllm |
vllm>=0.31.0 below Python 3.11; vllm>=0.32.0 on Python 3.11 through 3.14. There is no applicable requirement in this extra on Python 3.15 or newer. |
sglang |
sglang[srt]>=0.5.21 below Python 3.11; sglang[srt]>=0.5.22 on Python 3.11 or newer. |
The engine extras carry their own Torch/Transformers pins; use a dedicated environment for each engine rather than assuming it shares the local-kernel dependency set. Installing an extra does not waive architecture, operator, checkpoint-format, or native-build requirements.
| Checkpoint / path | Loader behavior |
|---|---|
| Serialized Marlin, local loader | AUTO becomes MARLIN; method and format select its GPTQ/AWQ sub-kernel. Explicit local selection must use MARLIN, or use MLX conversion. A class declaring Marlin-format support does not bypass this loader restriction. |
| Serialized BitBLAS, local loader | AUTO becomes GPTQ_BITBLAS or AWQ_BITBLAS according to the method. Explicit local selection must match that backend, or use MLX conversion. |
| EXL3 | Accepts AUTO, EXL3_EXLLAMA_V3, EXL3_TORCH, or MLX. AUTO chooses native ExLlama v3 only when torch.cuda.is_available() is true and the selector family is NVIDIA CUDA; on ROCm (CDNA and RDNA4 alike) it resolves to EXL3_TORCH, and other non-CUDA targets also choose Torch. This branch does not probe native-extension build success before choosing the native backend, so an extension failure is not an automatic retry with Torch. Explicit EXL3_EXLLAMA_V3 inference rejects ROCm/CPU with guidance to use EXL3_TORCH; HIP quantization is a separate path. |
| bitsandbytes | Accepts AUTO, BITSANDBYTES, or MLX. AUTO becomes BITSANDBYTES in this branch. The local class explicitly declares CPU, CUDA, and ROCm, 4/8-bit. |
| Native GGUF Q1_0 | Loader restricts selection to AUTO, GGUF_TORCH, GGUF_TRITON, or MLX. A generic GGUF class's broad bit list is not sufficient to override this qtype-level rule. |
| Native GGUF Q2_0 | AUTO becomes GGUF_TORCH; explicit selection must be GGUF_TORCH or MLX. |
| AWQ GEMV / GEMV_FAST / LLM_AWQ | GEMV_FAST/LLM_AWQ force model dtype torch.float16; these formats are not interchangeable with GEMM. Native GEMV-family readers remain CUDA-only. ROCm rejects these layouts unless AWQ_GEMV_TRITON is explicitly selected; that backend reads GEMV only, not GEMV_FAST/LLM_AWQ, and still applies its format/shape/dtype checks. AUTO does not opt into it. |
| vLLM | External-loader format allowlist: GPTQ and GEMM. Runtime support also depends on the installed vLLM version and its own kernels. |
| SGLang | External-loader format allowlist: GPTQ, GPTQ_V2, GEMM, and MARLIN. Engine initialization and generation are delegated to the SGLang integration. |
The external-engine paths return a wrapper with qlinear_kernel=None before the local Marlin/BitBLAS restrictions above. They use their own format allowlists, not generic local qlinear selection; selecting them does not promise that the engine implements every GPT-QModel local format. See Serving.
normalize_profile() accepts enum members, case-insensitive names/values, and the integer aliases below. String names may use spaces or hyphens in place of underscores. None and an empty/whitespace string resolve to AUTO; an unknown integer/name or an unsupported type raises an error.
| Profile | Value | Integer alias | Native Q1_0 behavior in from_pretrained() with no quantization configuration |
|---|---|---|---|
PROFILE.AUTO |
"auto" |
0 |
Resolves to FAST. |
PROFILE.FAST |
"fast" |
1 |
Uses the dense Hugging Face GGUF import path. A non-AUTO backend is logged as ignored on this path. |
PROFILE.LOW_MEMORY |
"low_memory" |
2 |
Does not take the dense FAST branch; native GGUF detection redirects to quantized loading. |
This is a loader-path decision, not a universal speed/memory switch on every qlinear class. GPTQModel.load() forwards profile through its from_pretrained() branch; its already-quantized from_quantized() branch does not forward this parameter. In particular, Q1_0 FAST can materialize dense weights even though the source file is quantized. Use LOW_MEMORY when entering the native GGUF detection path to request packed execution, and still obey the qtype-specific backend restrictions above. See GGUF.
gptqmodel/utils/importer.py imports non-private modules under gptqmodel.nn_modules.qlinear. Import failures caused by ImportError or OSError can exclude optional modules. It walks BaseQuantLinear subclasses, retaining selectable classes with format declarations, then builds:
| Mapping | Purpose |
|---|---|
AUTO_BACKEND_KERNEL_MAPPING |
Method → format → backend/class candidates in AUTO order. |
BACKEND_TO_METHOD_FORMAT_MAPPING |
Method/format/backend support lookup, including explicit-only declarations. |
SUPPORTS_FORMATS maps format enums to integer priorities. Positive priorities participate in AUTO; zero or negative priorities do not. Candidates are sorted by descending priority and backend string value; duplicate backend entries are removed in the resulting order. AUTO is therefore format-specific, not one global ranking of backend names.
For example, CPU AVX512 has GPTQ priority 120 and CPU ATen 110, while Swordfish has GPTQ priority 101, Machete 100, Marlin 90, ExLlama v2 80, Torch fused 50, Triton 40, BitBLAS 30, and Torch 20. Device, dtype, shape, and availability checks eliminate incompatible candidates. These numbers do not say that a CPU backend runs on CUDA, or that any candidate is a measured speed winner. The complete declarations, including AWQ, GGUF, and method-specific priorities, are in Kernel Backends Reference.
def select_quant_linear(
bits,
group_size: int,
desc_act: bool,
sym: bool,
device: Optional[SelectorDevices],
backend: BACKEND = BACKEND.AUTO,
format: FORMAT = FORMAT.GPTQ,
quant_method: METHOD = METHOD.GPTQ,
pack: bool = False,
allow_marlin: bool = True, # TODO: remove this after marlin padding is fixed
dynamic=None,
pack_dtype: torch.dtype = None,
dtype: Optional[torch.dtype] = None,
multi_select: bool = False, # return all valid kernels
adapter: Optional[Adapter] = None,
is_sharded: bool = False,
) -> Union[Type[BaseQuantLinear], List[Type[BaseQuantLinear]]]:
...The selector performs the following operations:
- Normalize format and method strings to enums, resolve method-dependent backend names, normalize devices without losing ordinals, and normalize the requested bit width.
- Reject unknown method/format combinations and incompatible device-family/capability combinations.
- For AUTO or AUTO_TRAINABLE, iterate the generated candidates. Respect
allow_marlin, declared device families, and sharded-load restrictions. - Validate each candidate against bits, group size, symmetry, activation ordering, pack dtype, activation dtype, dynamic overrides, trainability, and adapter support. Exact targets in a multi-device selector are checked individually.
- When
pack=True, additionally require a supported raw packing API.SUPPORTS_RAW_PACKING=Falseexcludes classes even if they can load an already-packed checkpoint. - Return the first valid class, or all qualifying classes when
multi_select=True. If nothing validates, propagate the last validation error where available, otherwise raise a no-valid-kernel error.
For multi-select with dynamic overrides, the code enumerates effective base/dynamic contracts and retains a class if it can handle an effective contract. This allows different classes to serve different layer contracts later; it does not mean every returned class supports every layer. Single-select retains a model-wide validation contract with the full dynamic map.
Adapter validation has an implementation asymmetry: AUTO forwards adapter into candidate validation, but the explicit-backend branch of select_quant_linear() does not forward it to validate_candidate(). A successful explicit selector call alone therefore does not establish adapter compatibility; check the class's adapter declarations and construction/runtime checks.
AUTO_TRAINABLE sets trainable=True during validation. It selects implementations that declare training support; it is not a promise that packed integer weights become ordinary trainable floating parameters. Adapter handling and the class's actual autograd implementation remain relevant. See EoRA.
On CPU-only GPTQ/AWQ loads that are not raw packing, plain AUTO also prefers a trainable kernel: HF/optimum load checkpoints with BACKEND.AUTO and PEFT or the Trainer then switch the model to training mode, so an inference-only CPU kernel picked by priority order alone would regress fine-tuning. AUTO falls back to an inference-only CPU kernel only when no trainable candidate validates. Accelerator targets and non-GPTQ/AWQ methods keep their priority order. Switching an inference-only kernel to training mode raises NotImplementedError pointing at backend="auto_trainable".
hf_select_quant_linear and hf_select_quant_linear_v2 are integration-facing wrappers around this selection machinery. Use the wrapper signature appropriate to the integration; do not maintain a separate hard-coded backend ranking in external code.
validate_quant_linear() uses a bounded, thread-safe cache with a maximum of 1024 entries. Keys incorporate the class's validation implementation, immutable contract values, and device descriptors/capability. Cached failures retain exception type/message metadata rather than mutable exception instances with traceback references; callers receive newly constructed exceptions.
The selector deduplicates device entries in first-seen order. For a concrete multi-device selector, it validates every retained exact ordinal. hf_normalize_device_device_map() ignores Accelerate "disk" entries and, when accelerator entries exist, excludes CPU offload entries from accelerator-kernel selection; an all-CPU map retains CPU as its target. Symbolic Accelerate placement keywords resolve to a selector device rather than an already-expanded list of final ordinals. Thus selection must not be described as validating a GPU kernel for every CPU/disk offload entry. Accelerator-family mixtures and CUDA/ROCm devices with different compute capabilities are rejected by the relevant device-family check; a kernel valid on the first GPU is not assumed valid on the others. This is distinct from quantization-time device pools described in Memory, Offload, and Multi-GPU.
Projection Fusion combines eligible post-quantization QKV
or gate/up calls only after backend selection. It is disabled by default; use
plan = model.fuse() or GPTQModel.load(..., fuse=True) with a qualified
definition. AUTO's selected exact native class is preserved, not replaced by
Torch to make a group eligible. Standalone backend support is broader than
fusion support: native layout, architecture, dtype, effective per-module
settings, and parent-forward checks still apply. See the fusion guide's
native backend limits and pinned
complete inventory for tested classes and explicit exclusions.
gptqmodel/nn_modules/converter.py currently defines an empty MODULE_CONVERTER_MAP. It includes a text-decoder-config helper and notes architectures handled by Defuser during model loading. It is not the authoritative backend dispatch table, and it is not a general checkpoint-format conversion service.
Backend selection is handled by utils/importer.py; layer packing/repacking is handled by utils/packing.py and kernel methods. Model-architecture conversion and quantized tensor packing are separate operations.
hf_pack_layer() receives an already-instantiated kernel module and the quantizer's original linear layer, scales, zeros, and optional activation-order index. It dispatches the module's pack, pack_original, pack_block, or pack_gpu implementation. QQQ and AWQ follow their own single-pack() branches.
| Packing control | Behavior |
|---|---|
pack_impl=None |
Normalizes to "original" for this API. |
"original", "pack_original"
|
Original packer. |
"auto", "default"
|
Prefer GPU packing when CUDA/HIP is available and the module has pack_gpu; otherwise block packing when present, then original. Embeddings do not take the GPU-packing path. |
"cpu", "block", "pack_block"
|
Block packer where implemented; otherwise log a warning and fall back to original. |
"gpu" |
GPU packing where available; otherwise warn and use original. Packed output buffers are registered on CPU; this option accelerates packing math rather than setting the final inference placement. |
Missing zeros
|
Symmetric modules can synthesize 2 ** (bits - 1); asymmetric packing without zeros raises. |
hf_repack_layer() selects the destination class using the source method/format and either calls a native repack hook or copies a shared checkpoint layout before destination post_init(). It returns the source unchanged when the requested backend/class already matches. A source already transformed into certain private runtime ABIs is rejected; repack from the checkpoint-layout holder instead.
Runtime initialization can allocate workspace, prepack tiles, shuffle scales/zeros, establish native handles, or create hardware operand buffers. This is why directly copying qweight into a different class and immediately calling forward() is not a supported general conversion procedure. ExLlama v2 initialization additionally needs its scratch-space arguments.
The shared linear helper preserves wider or incomparable stored bias for FP16/BF16 activations: it reduces and adds bias in FP32 with autocast disabled, then casts once to the activation dtype. A narrower bias can be promoted losslessly. FP4 and VPTQ may keep their cheaper path only when a value check proves the stored bias narrows exactly; otherwise they bypass native/fused execution for that call. VPTQ checkpoint bias is kept in FP32 even with FP16/BF16 codebooks. These checks do not authorize persistent dense weight caches. See the bias-width contract.
On ROCm, batched FP16/BF16 Torch linear calls with bias flatten to a 2D fused matmul/bias epilogue before restoring the leading shape. Bias is no longer added after an already-rounded matmul result, preserving cancellation and avoiding premature output overflow. Unbiased dispatch and packed checkpoint representation are unchanged.
CPU portable GPTQ with FP16/BF16 input and bias uses addmm rather than a separately rounded matmul followed by bias addition. When decoded weights or bias have a wider dtype, the epilogue uses FP32 and disables CPU autocast before the final activation-dtype cast. Native 8-bit GPTQ_TORCH_INT8 similarly adds bias to its FP32 reduction before output rounding. GGUF CPU dense fallback and the C++ wrappers finish bias before narrowing too; low-precision ggml output with bias can require a corrective FP32-decode/matmul. FP8 requests FP32 biased scaled-mm output on NVIDIA and ROCm and uses a wide fallback when bias dtype or autocast requires it. See GGUF arithmetic and FP8 storage and forward contracts.
On NVIDIA CUDA, portable GPTQ, Torch AWQ, Torch FP8, and Torch GGUF paths keep biased matmul output in FP32 through bias addition before rounding to the activation dtype. GPTQ streaming also keeps unbiased tile reductions in FP32 so a narrow split reduction cannot overflow prematurely. The shared matmul_fp32_reduce helper supplies a custom first-order backward for FP32-output mm/addmm when gradients are required; it computes requested input/weight gradients with wide GEMMs and bias gradients with an FP32 reduction. This restores activation/adapter training through these paths, not floating-point training of packed integer codes.
| Backend path | Bias and gradient behavior |
|---|---|
GPTQ_TRITON (TritonV2Linear) |
Passes bias into the shared wide dequantize/matmul epilogue; its custom backward returns the correct argument count for input gradients. |
AWQ_GEMM |
Preserves input and bias autograd in evaluation mode, including empty-row gradients. Packed integer weights remain non-trainable. |
AWQ_GEMM_TRITON |
Biased fused calls request FP32 output even when fp32_accum=False, add bias, then cast once to FP16. Dense fallback uses the shared helper; backward supplies input and bias gradients, including zero bias gradient for empty rows. |
AWQ_GEMV_TRITON |
Adds bias to an FP32 kernel output before casting. When input or bias gradients are needed, uses differentiable temporary dequantization plus matmul rather than the inference-only fused kernel. |
PAROQUANT_TRITON |
Decode/prefill fused kernels add bias before final conversion. Input/rotation or bias gradient requirements force the dense differentiable plan instead of autotuning a detached fused route. |
QQQ_TORCH |
Adds bias inside the shared wide epilogue after its existing activation quantize/dequantize step; this does not remove QQQ's activation quantization. |
GGUF_TRITON |
BF16/FP32 inputs use row-scaled FP16 staging when unbiased, or FP32 activations/dots and wide K-quant metadata when biased. Bias and restored row scales apply before output conversion; see GGUF arithmetic. The class remains inference-only. |
These are backend-specific changes: final results still must fit the requested dtype, and temporary dense fallback weights are not retained between forwards. Native AWQ GEMV/GEMV_FAST still have an FP16 kernel-result limit; do not confuse their wrapper-side wide bias addition with the Triton GEMV epilogue. The source arithmetic guide contains per-path flows and recorded device measurements.
Native and Triton AWQ GEMM do not enter internal inference mode when input/bias gradients are needed or an adapter is attached. Built-in Lora.apply() returns a new sum instead of updating the base output in place, preserving custom-autograd views and adapter gradients. This does not make every adapter-compatible backend trainable; the backend's own gradient contract still applies. See EoRA runtime loading.
Selected low-level helpers now preserve logical values for non-contiguous views rather than assuming every supplied tensor has dense storage. This is operand normalization, not a new backend, packed format, or expansion of hardware/shape eligibility.
| Helper / path | Strided operand handling |
|---|---|
GPTQ triton_utils/dequant.py::dequant()
|
Makes qweight, scales, qzeros, and g_idx contiguous before dense-offset decoding. |
| ParoQuant Triton dequantization / GEMM | Normalizes packed weights, scales, and zero points; fused GEMM also normalizes activations. |
| BitsAndBytes Triton NF4/FP4 | Normalizes packed bytes, absmax scales, codebook, GEMM activations, and bias. Global nibble parity still handles odd input widths. |
QQQ direct Triton dynamic_quant()
|
Normalizes activation rows. The public class still selects Triton only for already-contiguous eligible inputs. |
| W4A-FP8 prepacked GEMM | Normalizes activation codes, token scales, weight operands, group scales, and bias. |
| W4A INT4-to-E4M3 unpack | Uses the actual packed-input and destination row/column strides rather than forcing a contiguous destination. |
| W4A-NVFP4 pack / decode | Normalizes input activations for every packing recipe, and packed codes/swizzled scales for decoding. |
| NF12 HIP linear | Normalizes flattened activation rows and restores all caller leading dimensions on every execution route. |
| EXL3 wide-input Torch helper | Reshapes non-contiguous BF16/FP32 activation views; the native wrapper routes these dtypes away from FP16 reconstruction. |
Contiguous copies can add transient memory and work for a strided input; they do not authorize retaining expanded weights between forwards. Helpers that require contiguous input and are not listed here retain their existing contracts. See kernel utilities, W4A transport, and NF12 for method-specific details.
GPTQ reorder metadata derives ret_idx with one stable argsort; fused Torch inference skips the input gather when that permutation is identity. This does not change activation-order semantics or introduce a user flag.
Runtime availability is checked before shape validation, including Machete. ExLlama v2 handles and GGUF CUDA execution plans are invalidated when their packed source tensors change; the latter uses an eight-plan LRU with immediate native-plan release on eviction. Fused CPU training/device transitions restore the saved packed checkpoint state before leaving the transformed inference layout. See availability and runtime-state lifecycle.
Native 8-bit GPTQ INT8 keeps its original qweight, qzeros, scales, and g_idx as persistent checkpoint state after post_init(). The derived centered INT8 codes/channel scales are non-persistent and refresh after source edits, loads, or moves; call clear_weight_cache() after unversioned inference-tensor mutation. This retains both packed source and native runtime storage, not a dense cache or a lossless grouped-GPTQ repack. ExLlama v2 keys act-order handles after native permutation metadata is populated so unchanged forwards reuse the completed handle. ParoQuant cache clear also recomputes the identity-rotation shortcut, and EXL3 tracks sign/codebook metadata edits while keeping its wide-input helper out of checkpoint keys. See native INT8 lifecycle, ParoQuant, and EXL3.
Current inference paths also protect intermediate numeric range: CUDA FP16/BF16 Torch matmul/linear helpers request FP32 reduction output before casting back; FP8/FP4 scale application and VPTQ residual/normalization arithmetic widen before the final cast. CPU ATen INT4 paths retain FP32 scale/offset metadata when narrowing would lose range or affine identity. Marlin avoids atomic partial-output reduction when FP32 reduction is selected. These are backend-specific fixes, not a new quantization format or a guarantee that an unrepresentable final output becomes finite. See GGUF chunk precision and QQQ activation scaling for method-specific behavior.
The Triton dequantize-plus-matmul helper uses the same CUDA FP32-reduction helper, including transpose mode. On HIP, batched low-precision Torch linear calls with bias flatten to a 2D fused epilogue so the matmul is not rounded before a cancelling bias is added. AWQ GEMV/GEMV_FAST add their FP16 kernel result and bias in FP32 before restoring the activation dtype; this avoids premature BF16 rounding, not the kernels' FP16 compute limit.
TorchFusedLinear and TorchAtenLinear refresh activation-order indexing when g_idx changes and rebuild derived scale/zero metadata after checkpoint loads, device/dtype moves, or observable source-buffer changes. They keep the transformed packed-weight layout rather than reconstructing dense weights. ExLlama v2 GPTQ/AWQ handles no longer retain dense-weight scratch (temp_dq_size() returns zero); row-count fallback reconstruction is temporary for the current call, so peak memory can still include it.
gptqmodel/nn_modules/qlinear/_runtime_cache.py supplies tensor_cache_key() for safely identifying tensor-backed cached state. The key accounts for tensor identity, version where available, data pointer, device, dtype, shape, and inference-tensor status. A shape/data-pointer-only key would miss in-place updates or allocator reuse.
The helper is not a global model cache and is not a disk checkpoint cache. Its source documents invalidation on state-dict load, device/dtype movement, and training transitions in the classes that own packed-weight caches. Inference tensors can be mutated inside InferenceMode without a normal version counter; callers that do so must call the module's clear_weight_cache() explicitly. Kernel-specific buffers and initialization rules still apply; the helper does not rebuild a private runtime ABI automatically.
For kernel development, distinguish permitted packed quantized state from a dense or expanded weight cache. The repository's storage policy forbids retaining expanded/dequantized weights between operations; an implementation using temporary dequantization must release it after its consuming operation. Inspect a backend's actual memory behavior rather than assuming that persistent=False means it occupies no memory. W4A FP8 and NVFP4's mixed FP8 lane now materialize their E4M3 weight operands per forward; NVFP4 retains only its packed alternate FP4 planes and their scales. Checkpoint INT4 size alone is still not a runtime memory estimate; see W4A storage details.
The shared extension loader now reports the current load failure rather than reusing an older JIT error after a later load attempt. On ROCm, extension.load("all") skips CUDA-only AWQ/QQQ/Marlin extensions; requesting one explicitly still fails. Building all RDNA4 family targets does not qualify AllSpark's unimplemented gfx1250 WMMA path. Native NF12 does have a non-packed-dot fallback there; see NF12.
| Symptom / assumption | What to check |
|---|---|
| Explicit backend fails but AUTO loads | The explicit class may require a different architecture, format, dtype, group size, extension, or unsharded checkpoint. Read the validation error before reinstalling the package. |
| A newer GPU cannot run Machete | Current Machete kernels target Hopper sm_90a, not an open-ended SM90+ range. Use a compatible alternative on Blackwell. |
| NVFP4 fails on a Swordfish-compatible GPU | W4A NVFP4 checks capability 12.1+, while Swordfish targets 10.x/11.0 variants. These are different kernels. |
W4AFP8 native FP8 GEMM needs Triton >= 3.6.0 (Gluon mma_v2) |
The grouped FP8 GEMM uses the Gluon mma_v2 API; upgrade Triton or select a different backend. Activation packing and pure W4A-NVFP4 do not need it. |
AWQ_TORCH_INT8 raises immediately |
This backend is retired intentionally, not a missing optional dependency. |
| A class appears to support the bit width but loading fails | Check checkpoint format, tensor qtype, shapes, group size, symmetry, desc_act, dtype, pack dtype, and loader-specific restrictions. |
| A class can load but cannot pack | Loading/repacking and raw float-weight packing have different contracts. Inspect SUPPORTS_RAW_PACKING and the pack API. |
| A sharded model fails with Swordfish | Its classes explicitly set SUPPORTS_SHARDED_LOAD=False. |
| Optional backend absent from AUTO | The module may not import, have priority zero, fail runtime validation, or be excluded by the active format/device. |
| A custom mutation produces stale results | Invalidate the owning module's weight cache and re-run any required post-initialization; inference-tensor mutation cannot always be detected automatically. |
| Multi-GPU inference works differently from multi-GPU quantization | Inference device_map and quantization dense/MoE calculation pools are different mechanisms. |
-
gptqmodel/utils/arch.pyandgptqmodel/nn_modules/qlinear/torch.py— architecture families and fused GPTQ gates. -
gptqmodel/nn_modules/qlinear/gemv_awq.pyandgptqmodel/quantization/awq/modules/triton/gemv.py— explicit portable GEMV layout reader. -
gptqmodel/nn_modules/qlinear/bitsandbytes_triton.py— fused four-bit inference and opt-in/default dispatch. -
gptqmodel/nn_modules/qlinear/allspark.py,allspark_awq.py— AllSpark method contracts and repacking. -
gptqmodel/utils/allspark.py,allspark_hip.py— device gates and extension dispatch. pyproject.tomlgptqmodel/utils/backend.pygptqmodel/utils/importer.pygptqmodel/models/auto.pygptqmodel/models/loader.pygptqmodel/nn_modules/converter.pygptqmodel/nn_modules/qlinear/__init__.py-
gptqmodel/nn_modules/qlinear/_numerics.py— operation-local FP32 reduction helpers. gptqmodel/nn_modules/qlinear/_runtime_cache.pygptqmodel/utils/packing.pygptqmodel/utils/marlin.pygptqmodel/utils/machete.pygptqmodel/utils/swordfish.pygptqmodel/nn_modules/qlinear/torch_int8_awq.pygptqmodel/nn_modules/qlinear/w4a_floatx.pygptqmodel/nn_modules/qlinear/w4a_nvfp4.pydocs/torch_fused_int4_transformations.mddocs/rocm-quant-blas-backend.md
Documentation home · Quickstart · Troubleshooting · Repository · PyPI
Source baseline: GPT-QModel 7.9.0 / 40948a6. Maintained by ModelCloud.ai.
Getting started
- What is Quantization?
- Installation and Quickstart
- Quantization Lifecycle
- QuantizeConfig Reference
- Packing Formats and Compatibility
Quantization methods
Complementary configs
Using a model
- Loading and Inference
- Projection Fusion
- Inference Backends
- Kernel Backends Reference
- Serving
- Evaluation
- Troubleshooting and FAQ
Platforms
- Memory, Offload, Multi-GPU
- MLX (Apple Silicon)
- Hugging Face Integration
- Supported Models
- Multimodal and MoE Models
Internals