diff --git a/.github/workflows/behavior-check.yml b/.github/workflows/behavior-check.yml index 7c5ecacdc..45845e11b 100644 --- a/.github/workflows/behavior-check.yml +++ b/.github/workflows/behavior-check.yml @@ -13,6 +13,8 @@ on: '**/*.hpp', '**/*.c', '**/*.cpp', + '**/*.py', + '**/requirements*.txt', ] pull_request: @@ -24,7 +26,9 @@ on: '**/*.h', '**/*.hpp', '**/*.c', - '**/*.cpp' + '**/*.cpp', + '**/*.py', + '**/requirements*.txt', ] cache-mode: none @@ -69,7 +73,7 @@ jobs: run: | python3 -m venv .venv source .venv/bin/activate - pip install jinja2==3.1.6 + pip install -r requirements/requirements-convert_hf_to_gguf.txt - name: ccache-buckets-restore uses: ./.github/actions/ccache-buckets @@ -134,7 +138,7 @@ jobs: run: | python3 -m venv .venv source .venv/Scripts/activate - pip install jinja2==3.1.6 + pip install -r requirements/requirements-convert_hf_to_gguf.txt - name: Build id: cmake_build diff --git a/conversion/base.py b/conversion/base.py index 066be61b2..6347d7fd3 100644 --- a/conversion/base.py +++ b/conversion/base.py @@ -16,6 +16,7 @@ from pathlib import Path from hashlib import sha256 from typing import TYPE_CHECKING, Any, Callable, ContextManager, Iterable, Iterator, Literal, Sequence, TypeVar, cast from itertools import chain +from math import prod from transformers import AutoConfig import numpy as np @@ -155,6 +156,7 @@ class ModelBase: self.fuse_gate_up_exps = fuse_gate_up_exps self._gate_exp_buffer: dict[int, Tensor] = {} self._up_exp_buffer: dict[int, Tensor] = {} + self._unfusable_gate_up_layers: set[int] = set() self.fuse_qkv = fuse_qkv self._q_buffer: dict[int, Tensor] = {} self._k_buffer: dict[int, Tensor] = {} @@ -175,7 +177,9 @@ class ModelBase: self._nvfp4_layer_algo: dict[str, str | None] = {} # per-layer quant_algo, keyed by HF module path self._prec_a4: dict[str, bool] = {} # gguf tensor name -> can use 4-bit (A4) activations self._fp8_as_q8 = fp8_as_q8 + self._preserve_fp8 = ftype == gguf.LlamaFileType.GUESSED and not fp8_as_q8 self._fp8_dequantized: set[str] = set() + self._fp8_e4m3_preserved: set[str] = set() # Apply heuristics to figure out typical tensor encoding based on first tensor's dtype # NOTE: can't use field "torch_dtype" in config.json, because some finetunes lie. @@ -326,7 +330,7 @@ class ModelBase: def dequant_model(self): # If all quantized tensors were already handled (e.g. pure NVFP4), skip - if self._is_nvfp4 and not any(k.endswith((".weight_scale", ".weight_scale_inv")) for k in self.model_tensors): + if self._is_nvfp4 and not any(k.endswith((".weight_scale", ".weight_scale_inv", ".input_scale", ".activation_scale", "_activation_scale", ".k_scale", ".v_scale")) for k in self.model_tensors): return tensors_to_remove: list[str] = [] @@ -446,7 +450,7 @@ class ModelBase: return (unpacked * scale.unsqueeze(-1).float()).reshape(shape) - def dequant_fp8() -> None: + def dequant_fp8(fp8_configs: list[dict[str, Any]] | None = None) -> None: for name in self.model_tensors.keys(): if name.endswith(".weight_scale"): weight_name = name.removesuffix("_scale") @@ -455,14 +459,29 @@ class ModelBase: continue w = self.model_tensors[weight_name] s = self.model_tensors[name] + block_size = None is_fp8_weight = False - if self._fp8_as_q8: - is_fp8_weight = w().dtype in (torch.float8_e4m3fn, torch.float8_e5m2) - self.model_tensors[weight_name] = lambda w=w, s=s: dequant_simple(w(), s(), None) + if self._fp8_as_q8 or fp8_configs is not None: + weight = w() + is_fp8_weight = weight.dtype in (torch.float8_e4m3fn, torch.float8_e5m2) + if fp8_configs is not None: + if not is_fp8_weight: + raise ValueError(f"Expected remaining weight {weight_name!r} to be FP8, got {weight.dtype}") + scale = s() + is_channel_scale = scale.ndim >= 1 and scale.ndim <= weight.ndim and scale.shape[0] == weight.shape[0] and all(size == 1 for size in scale.shape[1:]) + if prod(scale.shape) != 1 and not is_channel_scale: + block_configs = [c for c in fp8_configs if c.get("strategy") == "block"] + if not block_configs or any(not c.get("block_structure") or c.get("group_size") is not None for c in block_configs): + raise ValueError(f"FP8 weight {weight_name!r} with scale shape {list(scale.shape)} requires a block layout") + block_sizes = {tuple(c["block_structure"]) for c in block_configs} + if len(block_sizes) != 1: + raise ValueError(f"FP8 weight {weight_name!r} has ambiguous block layouts") + block_size = next(iter(block_sizes)) + self.model_tensors[weight_name] = lambda w=w, s=s, bs=block_size: dequant_simple(w(), s(), bs) tensors_to_remove.append(name) - if is_fp8_weight: + if self._fp8_as_q8 and is_fp8_weight: self._fp8_dequantized.add(weight_name) - if name.endswith((".input_scale", ".k_scale", ".v_scale")): + if name.endswith((".input_scale", ".activation_scale", "_activation_scale", ".k_scale", ".v_scale")): tensors_to_remove.append(name) if quant_method == "bitnet": @@ -527,29 +546,37 @@ class ModelBase: nvfp4_compressed_tensors = self._is_nvfp4_compressed_tensors(quant_method, quant_format, groups) if nvfp4_compressed_tensors: - dequant_fp8() + fp8_configs = [ + g["weights"] for g in groups.values() + if isinstance(g, dict) and g.get("format") == "float-quantized" + and g.get("weights", {}).get("type") == "float" + and g.get("weights", {}).get("num_bits") == 8 + ] + dequant_fp8(fp8_configs) elif len(groups) > 1: raise NotImplementedError("Can't handle multiple config groups for compressed-tensors yet") elif quant_format == "float-quantized" or quant_format == "int-quantized" or quant_format == "naive-quantized": weight_config = tuple(groups.values())[0]["weights"] block_size = weight_config.get("block_structure", None) strategy = weight_config.get("strategy") - assert strategy == "channel" or strategy == "block" - assert weight_config.get("group_size") is None # didn't find a model using this yet is_fp8 = ( quant_format == "float-quantized" and weight_config.get("type") == "float" and weight_config.get("num_bits") == 8 ) + assert strategy in ("channel", "block") or (is_fp8 and strategy == "tensor") + assert weight_config.get("group_size") is None # didn't find a model using this yet for name in self.model_tensors.keys(): if name.endswith(".weight_scale"): weight_name = name.removesuffix("_scale") w = self.model_tensors[weight_name] s = self.model_tensors[name] - self.model_tensors[weight_name] = lambda w=w, s=s: dequant_simple(w(), s(), block_size) + self.model_tensors[weight_name] = lambda w=w, s=s, bs=block_size: dequant_simple(w(), s(), bs) tensors_to_remove.append(name) if self._fp8_as_q8 and is_fp8: self._fp8_dequantized.add(weight_name) + if is_fp8 and name.endswith((".input_scale", ".activation_scale", "_activation_scale", ".k_scale", ".v_scale")): + tensors_to_remove.append(name) elif quant_format == "pack-quantized": weight_config = tuple(groups.values())[0]["weights"] assert weight_config.get("strategy") == "group" @@ -587,6 +614,148 @@ class ModelBase: for name, value in new_tensors.items(): self.model_tensors[name] = value + def _map_fp8_weight_names(self, name: str) -> tuple[str, ...]: + new_name = self.tensor_map.get_name(name, try_suffixes=(".weight",)) + return (new_name,) if new_name is not None else () + + def _prepare_fp8_e4m3_tensors(self): + if not self._preserve_fp8: + return + + scale_tensors: dict[str, list[tuple[int, float]] | np.ndarray] = {} + input_scale_tensors: dict[str, list[tuple[int, float]] | np.ndarray] = {} + consumed: list[str] = [] + candidates: dict[str, tuple[str, Tensor, tuple[str, ...], torch.Size]] = {} + weight_dtypes: dict[str, torch.dtype] = {} + scaled_weights: set[str] = set() + expert_pattern = re.compile(r"(\.[^.]*experts)\.(\d+)\.") + scale_types = { + gguf.MODEL_TENSOR.ATTN_Q, gguf.MODEL_TENSOR.ATTN_K, gguf.MODEL_TENSOR.ATTN_V, + gguf.MODEL_TENSOR.ATTN_OUT, gguf.MODEL_TENSOR.ATTN_QKV, gguf.MODEL_TENSOR.ATTN_GATE, + gguf.MODEL_TENSOR.FFN_GATE, gguf.MODEL_TENSOR.FFN_UP, gguf.MODEL_TENSOR.FFN_DOWN, + gguf.MODEL_TENSOR.FFN_GATE_SHEXP, gguf.MODEL_TENSOR.FFN_UP_SHEXP, gguf.MODEL_TENSOR.FFN_DOWN_SHEXP, + gguf.MODEL_TENSOR.FFN_GATE_EXP, gguf.MODEL_TENSOR.FFN_UP_EXP, gguf.MODEL_TENSOR.FFN_DOWN_EXP, + gguf.MODEL_TENSOR.SSM_IN, gguf.MODEL_TENSOR.SSM_OUT, gguf.MODEL_TENSOR.SSM_ALPHA, gguf.MODEL_TENSOR.SSM_BETA, + gguf.MODEL_TENSOR.NEXTN_EH_PROJ, gguf.MODEL_TENSOR.NEXTN_SHARED_HEAD_HEAD, gguf.MODEL_TENSOR.OUTPUT, + } + + for scale_name in list(self.model_tensors.keys()): + weight_name = None + if scale_name.endswith(".weight_scale"): + weight_name = scale_name.removesuffix("_scale") + elif scale_name.endswith("_scale_inv"): + weight_name = scale_name.removesuffix("_scale_inv") + elif scale_name.endswith(".qscale_weight"): + weight_name = scale_name.removesuffix("qscale_weight") + "weight" + + if weight_name is None or weight_name not in self.model_tensors: + continue + + scaled_weights.add(weight_name) + weight = self.model_tensors[weight_name]() + weight_dtypes[weight_name] = weight.dtype + if weight.dtype != torch.float8_e4m3fn or weight.ndim < 2: + continue + + scale = self.model_tensors[scale_name]() + if prod(scale.shape) != 1: + continue + + merged_name = expert_pattern.sub(r"\1.", weight_name) + new_names = self._map_fp8_weight_names(merged_name) + if not new_names: + continue + if any((mapped := self.tensor_map.get_type_and_name(new_name, try_suffixes=(".weight",))) is None or mapped[0] not in scale_types for new_name in new_names): + continue + bid = next((int(part) for part in new_names[0].split(".") if part.isdecimal()), None) + self._fp8_e4m3_preserved.add(weight_name) + if any(new_name.endswith("_norm.weight") or self.tensor_force_quant(weight_name, new_name, bid, weight.ndim) != gguf.GGMLQuantizationType.F8_E4M3 for new_name in new_names): + self._fp8_e4m3_preserved.discard(weight_name) + continue + candidates[weight_name] = (scale_name, scale, new_names, weight.shape) + + expert_groups: dict[str, list[str]] = {} + for name in self.model_tensors: + if (match := expert_pattern.search(name)) and name.endswith(".weight"): + expert_groups.setdefault(name[:match.end(1)], []).append(name) + for names in expert_groups.values(): + for name in names: + if name not in weight_dtypes: + weight_dtypes[name] = self.model_tensors[name]().dtype + dtypes = {weight_dtypes[name] for name in names} + if torch.float8_e4m3fn not in dtypes: + continue + if dtypes == {torch.float8_e4m3fn} and not any(name in scaled_weights for name in names): + continue + # The converter merges expert projections together. + if not all(name in candidates for name in names): + for name in names: + candidates.pop(name, None) + self._fp8_e4m3_preserved.discard(name) + if name not in scaled_weights: + gen = self.model_tensors[name] + self.model_tensors[name] = lambda gen=gen: gen().float() + + for weight_name, (scale_name, scale, new_names, shape) in candidates.items(): + scale = LazyTorchTensor.to_eager(scale).float().flatten() + + weight_prefix = weight_name.removesuffix(".weight") + # Transformers fine-grained FP8 uses activation_scale while ModelOpt uses input_scale. + # Ref: https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/models/mistral3.py#L357-L361 + input_scale_name = next((name for name in ( + weight_prefix + ".input_scale", + weight_prefix + ".activation_scale", + weight_prefix + "_activation_scale", + weight_prefix + ".qscale_act", + ) if name in self.model_tensors), None) + input_scale = None + if input_scale_name is not None: + input_scale = self.model_tensors[input_scale_name]() + if prod(input_scale.shape) != 1: + raise ValueError(f"FP8 input scale {input_scale_name!r} must be a scalar") + input_scale = LazyTorchTensor.to_eager(input_scale).float().flatten() + consumed.append(input_scale_name) + + consumed.append(scale_name) + + expert_match = expert_pattern.search(weight_name) + if expert_match: + expert_id = int(expert_match.group(2)) + for new_name in new_names: + target_name = new_name.replace(".weight", ".scale") + entries = scale_tensors.setdefault(target_name, []) + assert isinstance(entries, list) + cast(list[tuple[int, float]], entries).append((expert_id, float(scale[0]))) + if input_scale is not None: + target_name = new_name.replace(".weight", ".input_scale") + entries = input_scale_tensors.setdefault(target_name, []) + assert isinstance(entries, list) + cast(list[tuple[int, float]], entries).append((expert_id, float(input_scale[0]))) + else: + for new_name in new_names: + mapped = self.tensor_map.get_type_and_name(new_name, try_suffixes=(".weight",)) if len(shape) == 3 else None + is_packed_expert = mapped is not None and mapped[0] in ( + gguf.MODEL_TENSOR.FFN_GATE_EXP, gguf.MODEL_TENSOR.FFN_UP_EXP, gguf.MODEL_TENSOR.FFN_DOWN_EXP, + ) + # Packed expert weights need one scale value per expert. + n_scales = shape[0] if is_packed_expert else 1 + scale_tensors[new_name.replace(".weight", ".scale")] = np.repeat(scale.numpy(), n_scales) + if input_scale is not None: + input_scale_tensors[new_name.replace(".weight", ".input_scale")] = np.repeat(input_scale.numpy(), n_scales) + + for name in consumed: + self.model_tensors.pop(name, None) + + for name, values in chain(scale_tensors.items(), input_scale_tensors.items()): + if isinstance(values, list): + values.sort(key=lambda item: item[0]) + scale = np.array([item[1] for item in values], dtype=np.float32) + else: + scale = values.astype(np.float32) + if not np.allclose(scale, 1.0, atol=1e-6): + logger.info(f" + {name} (FP8 scale, shape [{scale.size}])") + self.gguf_writer.add_tensor(name, scale) + @classmethod def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: name, gen = item @@ -631,6 +800,26 @@ class ModelBase: raise ValueError(f"Can not map tensor {name!r}") return new_name + def prepare_gate_up_fusion(self) -> None: + self._unfusable_gate_up_layers.clear() + if not self.fuse_gate_up_exps or self._fp8_as_q8: + return + + for name, gen in self.model_tensors.items(): + merged_name = re.sub(r"(\.experts)\.\d+\.", r"\1.", name) + mapped = self.tensor_map.get_type_and_name(merged_name, try_suffixes=(".weight",)) + if mapped is None: + continue + tensor_type, new_name = mapped + if tensor_type not in (gguf.MODEL_TENSOR.FFN_GATE_EXP, gguf.MODEL_TENSOR.FFN_UP_EXP, gguf.MODEL_TENSOR.FFN_GATE_UP_EXP) or not new_name.endswith(".weight"): + continue + if gen().dtype == torch.float8_e4m3fn: + bid = next(int(part) for part in new_name.split(".") if part.isdecimal()) + self._unfusable_gate_up_layers.add(bid) + + for bid in sorted(self._unfusable_gate_up_layers): + logger.warning(f"Skipping --fuse-gate-up-exps for layer {bid}: fusion of preserved FP8 expert projections is not supported") + def prepare_qkv_fusion(self) -> None: self._fusable_qkv_weight_layers.clear() self._fusable_qkv_bias_layers.clear() @@ -644,6 +833,7 @@ class ModelBase: } weights: dict[int, set[gguf.MODEL_TENSOR]] = {} biases: dict[int, set[gguf.MODEL_TENSOR]] = {} + fp8_layers: set[int] = set() for name in self.model_tensors: mapped = self.tensor_map.get_type_and_name(name, try_suffixes=(".weight", ".bias")) @@ -658,12 +848,19 @@ class ModelBase: continue if new_name.endswith(".weight"): weights.setdefault(bid, set()).add(tensor_type) + if not self._fp8_as_q8 and self.model_tensors[name]().dtype == torch.float8_e4m3fn: + fp8_layers.add(bid) elif new_name.endswith(".bias"): biases.setdefault(bid, set()).add(tensor_type) for bid, weight_types in weights.items(): bias_types = biases.get(bid, set()) if weight_types == qkv_types and (not bias_types or bias_types == qkv_types): + # NVFP4 weights use separate reordering and repacking hooks and are exported before this fusion pass. + # Separate FP8 scales cannot be represented by one fused QKV scale. + if bid in fp8_layers: + logger.warning(f"Skipping --fuse-qkv for layer {bid}: fusion of preserved FP8 projections is not supported") + continue self._fusable_qkv_weight_layers.add(bid) if bias_types: self._fusable_qkv_bias_layers.add(bid) @@ -852,7 +1049,8 @@ class ModelBase: new_name = self.map_tensor_name(name) # Handle gate/up expert tensor fusion if enabled - if self.fuse_gate_up_exps and bid is not None: + # Preserved FP8 gate/up scales must stay with their separate projections. + if self.fuse_gate_up_exps and bid is not None and bid not in self._unfusable_gate_up_layers: if self.match_model_tensor_name(new_name, gguf.MODEL_TENSOR.FFN_GATE_EXP, bid): self._gate_exp_buffer[bid] = data_torch elif self.match_model_tensor_name(new_name, gguf.MODEL_TENSOR.FFN_UP_EXP, bid): @@ -909,8 +1107,46 @@ class ModelBase: return [(new_name, data_torch)] + def _tensor_requires_f32(self, new_name: str, bid: int | None, n_dims: int) -> bool: + # Conditions should closely match those in llama_model_quantize_internal in llama.cpp. + return n_dims <= 1 or new_name.endswith("_norm.weight") or ( + any( + self.match_model_tensor_name(new_name, key, bid) + for key in ( + gguf.MODEL_TENSOR.FFN_GATE_INP, + gguf.MODEL_TENSOR.FFN_GATE_INP_SHEXP, + gguf.MODEL_TENSOR.POS_EMBD, + gguf.MODEL_TENSOR.TOKEN_TYPES, + gguf.MODEL_TENSOR.SSM_CONV1D, + gguf.MODEL_TENSOR.SHORTCONV_CONV, + gguf.MODEL_TENSOR.TIME_MIX_FIRST, + gguf.MODEL_TENSOR.TIME_MIX_W1, + gguf.MODEL_TENSOR.TIME_MIX_W2, + gguf.MODEL_TENSOR.TIME_MIX_DECAY_W1, + gguf.MODEL_TENSOR.TIME_MIX_DECAY_W2, + gguf.MODEL_TENSOR.TIME_MIX_LERP_FUSED, + gguf.MODEL_TENSOR.POSNET_NORM1, + gguf.MODEL_TENSOR.POSNET_NORM2, + gguf.MODEL_TENSOR.V_ENC_EMBD_POS, + gguf.MODEL_TENSOR.A_ENC_EMBD_POS, + gguf.MODEL_TENSOR.ALTUP_CORRECT_COEF, + gguf.MODEL_TENSOR.ALTUP_PREDICT_COEF, + # Kimi KDA conv weights should be F32 + gguf.MODEL_TENSOR.SSM_CONV1D_Q, + gguf.MODEL_TENSOR.SSM_CONV1D_K, + gguf.MODEL_TENSOR.SSM_CONV1D_V, + # DSA indexer weights should be F32 + gguf.MODEL_TENSOR.INDEXER_PROJ, + ) + ) + or new_name[-7:] not in (".weight", ".lora_a", ".lora_b") + ) + def tensor_force_quant(self, name: str, new_name: str, bid: int | None, n_dims: int) -> gguf.GGMLQuantizationType | bool: - del new_name, bid # unused + if name in self._fp8_e4m3_preserved and n_dims >= 2: + if self._tensor_requires_f32(new_name, bid, n_dims): + return gguf.GGMLQuantizationType.F32 + return gguf.GGMLQuantizationType.F8_E4M3 # Force FP8-original tensors to Q8_0 when requested; Q8_0 is faster than F16/BF16. if self._fp8_as_q8 and name in self._fp8_dequantized and n_dims >= 2: return gguf.GGMLQuantizationType.Q8_0 @@ -1052,14 +1288,17 @@ class ModelBase: input_scale_name = name.replace(".weight", ".input_scale") if scale_name not in self.model_tensors: continue - # Force eager materialization of lazy tensors - weight = LazyTorchTensor.to_eager(self.model_tensors[name]()) - scale = LazyTorchTensor.to_eager(self.model_tensors[scale_name]()) + weight = self.model_tensors[name]() + scale = self.model_tensors[scale_name]() - # Skip non-NVFP4 tensors(e.g. 1D scale, or float8 weight) - if scale.ndim < 2 or weight.dtype in (torch.float8_e4m3fn, torch.float8_e5m2): + # Leave FP8 weights and their scales for _prepare_fp8_e4m3_tensors. + if weight.dtype in (torch.float8_e4m3fn, torch.float8_e5m2) or scale.ndim < 2: continue + if weight.ndim != 2 or scale.shape[0] != weight.shape[0] or scale.shape[1] * 8 != weight.shape[1]: + raise ValueError(f"NVFP4 weight {name!r} has incompatible shapes: weight {list(weight.shape)}, scale {list(scale.shape)}") + weight = LazyTorchTensor.to_eager(weight) + scale = LazyTorchTensor.to_eager(scale) scale2 = LazyTorchTensor.to_eager(self.model_tensors.get(scale2_name, lambda: torch.tensor(1.0))()) input_scale = LazyTorchTensor.to_eager(self.model_tensors.get(input_scale_name, lambda: torch.tensor(1.0))()) @@ -1215,8 +1454,10 @@ class ModelBase: self.model_tensors[input_scale_name] = inverse_scale(self.model_tensors.pop(name)) self._generate_nvfp4_tensors() + self._prepare_fp8_e4m3_tensors() self.dequant_model() + self.prepare_gate_up_fusion() self.prepare_qkv_fusion() # Handle empty tensor_map for models with block_count=0 (like MobileNetV5) @@ -1230,10 +1471,21 @@ class ModelBase: if name.endswith((".attention.masked_bias", ".attention.bias", ".rotary_emb.inv_freq")): continue + # FP8 transforms must only reorder values; norms and vectors use numeric F32 transforms. + preserve_fp8_e4m3 = self._preserve_fp8 and data_torch.dtype == torch.float8_e4m3fn and data_torch.ndim >= 2 and not name.endswith(("norm.weight", "norm_f.weight")) + if preserve_fp8_e4m3: + for new_name in self._map_fp8_weight_names(name): + mapped_bid = next((int(part) for part in new_name.split(".") if part.isdecimal()), None) + if self._tensor_requires_f32(new_name, mapped_bid, data_torch.ndim): + preserve_fp8_e4m3 = False + break + if preserve_fp8_e4m3: + self._fp8_e4m3_preserved.add(name) + old_dtype = data_torch.dtype # convert any unsupported data types to float32 - if data_torch.dtype not in (torch.float16, torch.float32): + if data_torch.dtype not in (torch.float16, torch.float32) and name not in self._fp8_e4m3_preserved: data_torch = data_torch.to(torch.float32) # use the first number-like part of the tensor name as the block id @@ -1244,51 +1496,14 @@ class ModelBase: break for new_name, data_torch in (self.modify_tensors(data_torch, name, bid)): - # TODO: why do we squeeze here? - # data = data_torch.squeeze().numpy() - data = data_torch.numpy() - - n_dims = len(data.shape) + n_dims = len(data_torch.shape) data_qtype: gguf.GGMLQuantizationType | bool = self.tensor_force_quant(name, new_name, bid, n_dims) # Most of the codebase that takes in 1D tensors or norms only handles F32 tensors if n_dims <= 1 or new_name.endswith("_norm.weight"): data_qtype = gguf.GGMLQuantizationType.F32 - # Conditions should closely match those in llama_model_quantize_internal in llama.cpp - # Some tensor types are always in float32 - if data_qtype is False and ( - any( - self.match_model_tensor_name(new_name, key, bid) - for key in ( - gguf.MODEL_TENSOR.FFN_GATE_INP, - gguf.MODEL_TENSOR.FFN_GATE_INP_SHEXP, - gguf.MODEL_TENSOR.POS_EMBD, - gguf.MODEL_TENSOR.TOKEN_TYPES, - gguf.MODEL_TENSOR.SSM_CONV1D, - gguf.MODEL_TENSOR.SHORTCONV_CONV, - gguf.MODEL_TENSOR.TIME_MIX_FIRST, - gguf.MODEL_TENSOR.TIME_MIX_W1, - gguf.MODEL_TENSOR.TIME_MIX_W2, - gguf.MODEL_TENSOR.TIME_MIX_DECAY_W1, - gguf.MODEL_TENSOR.TIME_MIX_DECAY_W2, - gguf.MODEL_TENSOR.TIME_MIX_LERP_FUSED, - gguf.MODEL_TENSOR.POSNET_NORM1, - gguf.MODEL_TENSOR.POSNET_NORM2, - gguf.MODEL_TENSOR.V_ENC_EMBD_POS, - gguf.MODEL_TENSOR.A_ENC_EMBD_POS, - gguf.MODEL_TENSOR.ALTUP_CORRECT_COEF, - gguf.MODEL_TENSOR.ALTUP_PREDICT_COEF, - # Kimi KDA conv weights should be F32 - gguf.MODEL_TENSOR.SSM_CONV1D_Q, - gguf.MODEL_TENSOR.SSM_CONV1D_K, - gguf.MODEL_TENSOR.SSM_CONV1D_V, - # DSA indexer weights should be F32 - gguf.MODEL_TENSOR.INDEXER_PROJ, - ) - ) - or new_name[-7:] not in (".weight", ".lora_a", ".lora_b") - ): + if data_qtype is False and self._tensor_requires_f32(new_name, bid, n_dims): data_qtype = gguf.GGMLQuantizationType.F32 if data_qtype is False and any( @@ -1326,16 +1541,22 @@ class ModelBase: else: raise ValueError(f"Unknown file type: {self.ftype.name}") + if data_torch.dtype == torch.float8_e4m3fn: + data = data_torch.view(torch.uint8).numpy() if data_qtype == gguf.GGMLQuantizationType.F8_E4M3 else data_torch.float().numpy() + else: + data = data_torch.numpy() + # a chunked tensor quantizes as one chunk at a time, while it is written quantize = data.quantize if isinstance(data, gguf.LazyChunkedTensor) else ( lambda qtype, d=data: gguf.quants.quantize(d, qtype)) - try: - data = quantize(data_qtype) - except gguf.QuantError as e: - logger.warning("%s, %s", e, "falling back to F16") - data_qtype = gguf.GGMLQuantizationType.F16 - data = quantize(data_qtype) + if not (data_qtype == gguf.GGMLQuantizationType.F8_E4M3 and data.dtype == np.uint8): + try: + data = quantize(data_qtype) + except gguf.QuantError as e: + logger.warning("%s, %s", e, "falling back to F16") + data_qtype = gguf.GGMLQuantizationType.F16 + data = quantize(data_qtype) shape = gguf.quant_shape_from_byte_shape(data.shape, data_qtype) if data.dtype == np.uint8 else data.shape @@ -1376,6 +1597,8 @@ class ModelBase: self.ftype = gguf.LlamaFileType.MOSTLY_NVFP4 elif self._is_mxfp4: self.ftype = gguf.LlamaFileType.MOSTLY_MXFP4_MOE + elif self._fp8_e4m3_preserved: + self.ftype = gguf.LlamaFileType.MOSTLY_F8_E4M3 # Generate parameter weight class (useful for leader boards) if not yet determined if self.metadata.size_label is None and total_params > 0: diff --git a/conversion/deepseek.py b/conversion/deepseek.py index 0f8e2ffee..581111da5 100644 --- a/conversion/deepseek.py +++ b/conversion/deepseek.py @@ -236,6 +236,12 @@ class DeepseekModel(TextModel): class DeepseekV2Model(TextModel): model_arch = gguf.MODEL_ARCH.DEEPSEEK2 + def _map_fp8_weight_names(self, name: str) -> tuple[str, ...]: + # MLA projections do not have runtime weight-scale support. + if name.endswith((".q_a_proj.weight", ".q_b_proj.weight", ".kv_a_proj_with_mqa.weight", ".kv_b_proj.weight")): + return () + return super()._map_fp8_weight_names(name) + # TODO @ngxson : remove this when we support MTP for deepseek models skip_mtp = True diff --git a/conversion/nemotron.py b/conversion/nemotron.py index d76a45d82..371f5b616 100644 --- a/conversion/nemotron.py +++ b/conversion/nemotron.py @@ -433,6 +433,11 @@ class NemotronHModel(GraniteHybridModel): _ATTN_LAYER_TYPES = {"attention", "full_attention"} _MLP_LAYER_TYPES = {"moe"} + def _map_fp8_weight_names(self, name: str) -> tuple[str, ...]: + if ".mixer.experts." in name: + name = name.replace("backbone.", "model.").replace(".mixer.experts.", ".mlp.experts.") + return super()._map_fp8_weight_names(name) + def __init__(self, *args, **kwargs): # We have to determine the correct model architecture (MoE vs non-MoE) before # calling the parent __init__. This is because the parent constructor diff --git a/conversion/qwen.py b/conversion/qwen.py index 5ec4ab43c..e6456e5a5 100644 --- a/conversion/qwen.py +++ b/conversion/qwen.py @@ -83,6 +83,15 @@ class Qwen2Model(TextModel): class Qwen2MoeModel(TextModel): model_arch = gguf.MODEL_ARCH.QWEN2MOE + def _map_fp8_weight_names(self, name: str) -> tuple[str, ...]: + if name.removesuffix(".weight").endswith(".mlp.experts.gate_up_proj"): + bid = next(int(part) for part in self.map_tensor_name(name).split(".") if part.isdecimal()) + return ( + self.format_tensor_name(gguf.MODEL_TENSOR.FFN_GATE_EXP, bid), + self.format_tensor_name(gguf.MODEL_TENSOR.FFN_UP_EXP, bid), + ) + return super()._map_fp8_weight_names(name) + def set_gguf_parameters(self): super().set_gguf_parameters() if (moe_intermediate_size := self.hparams.get("moe_intermediate_size")) is not None: @@ -382,6 +391,16 @@ class _QwenMtpMixin: class Qwen3NextModel(_QwenMtpMixin, Qwen2MoeModel): model_arch = gguf.MODEL_ARCH.QWEN3NEXT + def _map_fp8_weight_names(self, name: str) -> tuple[str, ...]: + if name.endswith(".linear_attn.in_proj_qkvz.weight"): + bid = next(int(part) for part in self.map_tensor_name(name).split(".") if part.isdecimal()) + # Both split projections keep the original per-tensor scales. + return ( + self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_QKV, bid), + self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_GATE, bid), + ) + return super()._map_fp8_weight_names(name) + def set_gguf_parameters(self): super().set_gguf_parameters() self.gguf_writer.add_ssm_conv_kernel(self.hparams["linear_conv_kernel_dim"]) diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index 57ba49dbc..b863bccc5 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -58,7 +58,7 @@ def parse_args() -> argparse.Namespace: ) parser.add_argument( "--outtype", type=str, choices=["f32", "f16", "bf16", "q8_0", "tq1_0", "tq2_0", "auto"], default="auto", - help="output format - use f32 for float32, f16 for float16, bf16 for bfloat16, q8_0 for Q8_0, tq1_0 or tq2_0 for ternary, and auto for the highest-fidelity 16-bit float type", + help="output format - use f32 for float32, f16 for float16, bf16 for bfloat16, q8_0 for Q8_0, tq1_0 or tq2_0 for ternary, and auto to preserve supported FP8 weights and otherwise select the highest-fidelity 16-bit float type", ) parser.add_argument( "--bigendian", action="store_true", diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index d62a63104..b188c0bf9 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -381,6 +381,8 @@ extern "C" { GGML_API void ggml_fp32_to_bf16_row_ref(const float *, ggml_bf16_t *, int64_t); GGML_API void ggml_fp32_to_bf16_row(const float *, ggml_bf16_t *, int64_t); + typedef struct { uint8_t bits; } ggml_fp8_e4m3_t; + struct ggml_object; struct ggml_context; struct ggml_cgraph; @@ -430,7 +432,8 @@ extern "C" { GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, GGML_TYPE_Q2_0 = 42, - GGML_TYPE_COUNT = 43, + GGML_TYPE_F8_E4M3 = 43, + GGML_TYPE_COUNT = 44, }; // [TAG_GGML_PREC] @@ -486,6 +489,7 @@ extern "C" { GGML_FTYPE_MOSTLY_NVFP4 = 26, // except 1d tensors GGML_FTYPE_MOSTLY_Q1_0 = 27, // except 1d tensors GGML_FTYPE_MOSTLY_Q2_0 = 28, // except 1d tensors + GGML_FTYPE_MOSTLY_F8_E4M3 = 29, // except 1d tensors }; // available tensor operations: diff --git a/ggml/src/ggml-cpu/common.h b/ggml/src/ggml-cpu/common.h index abbadc359..2c62e6976 100644 --- a/ggml/src/ggml-cpu/common.h +++ b/ggml/src/ggml-cpu/common.h @@ -43,6 +43,14 @@ static inline float f32_to_f32(float x) { return x; } +static inline float f8_e4m3_to_f32(ggml_fp8_e4m3_t x) { + return ggml_f8_e4m3_to_fp32(x.bits); +} + +static inline ggml_fp8_e4m3_t f32_to_f8_e4m3(float x) { + return { ggml_fp32_to_f8_e4m3(x) }; +} + // TODO - merge this into the traits table, after using row-based conversions template struct type_conversion_table; @@ -65,6 +73,12 @@ struct type_conversion_table { static constexpr ggml_bf16_t (*from_f32)(float) = f32_to_bf16; }; +template <> +struct type_conversion_table { + static constexpr float (*to_f32)(ggml_fp8_e4m3_t) = f8_e4m3_to_f32; + static constexpr ggml_fp8_e4m3_t (*from_f32)(float) = f32_to_f8_e4m3; +}; + template <> struct type_conversion_table { static constexpr float (*to_f32)(int32_t) = i32_to_f32; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 53d39d249..f4eaea5df 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -86,6 +86,9 @@ float ggml_table_f32_e8m0_half[1 << 8]; // precomputed f32 table for ue4m3 (1 KB) (simd-mappings.h) float ggml_table_f32_ue4m3[1 << 8]; +// Signed FP8 needs its own table: UE4M3 ignores the sign, halves values for NVFP4, and maps NaN to zero. +float ggml_table_f32_f8_e4m3[1 << 8]; + #if defined(__ARM_ARCH) struct ggml_arm_arch_features_type { int sve_cnt; @@ -296,6 +299,12 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_0, .nrows = 1, }, + [GGML_TYPE_F8_E4M3] = { + .from_float = quantize_row_f8_e4m3, + .vec_dot = ggml_vec_dot_f8_e4m3_f32, + .vec_dot_type = GGML_TYPE_F32, + .nrows = 1, + }, [GGML_TYPE_Q2_K] = { .from_float = quantize_row_q2_K, .vec_dot = ggml_vec_dot_q2_K_q8_K, @@ -2876,7 +2885,7 @@ struct ggml_cplan ggml_graph_plan( case GGML_OP_ADD_ID: case GGML_OP_ADD1: { - if (ggml_is_quantized(node->src[0]->type)) { + if (ggml_is_quantized(node->src[0]->type) || node->src[0]->type == GGML_TYPE_F8_E4M3) { cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks; } } break; @@ -2924,7 +2933,7 @@ struct ggml_cplan ggml_graph_plan( } break; case GGML_OP_OUT_PROD: { - if (ggml_is_quantized(node->src[0]->type) || + if (ggml_is_quantized(node->src[0]->type) || node->src[0]->type == GGML_TYPE_F8_E4M3 || node->src[0]->type == GGML_TYPE_F16) { cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks; } @@ -3908,6 +3917,7 @@ void ggml_cpu_init(void) { // initialize UE4M3 table (256 entries) for (int i = 0; i < (1 << 8); ++i) { ggml_table_f32_ue4m3[i] = ggml_ue4m3_to_fp32(i); + ggml_table_f32_f8_e4m3[i] = ggml_f8_e4m3_to_fp32(i); } const uint64_t t_end = ggml_time_us(); UNUSED(t_end); diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 85e17018c..d3a156b12 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -537,26 +537,36 @@ void ggml_compute_forward_dup( switch (src0->type) { case GGML_TYPE_F16: { - /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F8_E4M3) ggml_compute_forward_dup_flt(params, dst); else ggml_compute_forward_dup_to_q(params, dst); } break; case GGML_TYPE_BF16: { - /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F8_E4M3) ggml_compute_forward_dup_flt(params, dst); else ggml_compute_forward_dup_to_q(params, dst); } break; case GGML_TYPE_F32: { - /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); - else if (dst->type == GGML_TYPE_I32) ggml_compute_forward_dup_flt(params, dst); + /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F8_E4M3) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_I32) ggml_compute_forward_dup_flt(params, dst); else ggml_compute_forward_dup_to_q(params, dst); } break; + case GGML_TYPE_F8_E4M3: + { + /**/ if (dst->type == GGML_TYPE_F16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_BF16) ggml_compute_forward_dup_flt(params, dst); + else if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); + else GGML_ABORT("not implemented"); + } break; case GGML_TYPE_I32: { if (dst->type == GGML_TYPE_F32) ggml_compute_forward_dup_flt(params, dst); @@ -605,8 +615,9 @@ static void ggml_compute_forward_add_q_f32( GGML_ASSERT(nb1 <= nb2); GGML_ASSERT(nb2 <= nb3); - GGML_ASSERT(ggml_is_quantized(src0->type)); + GGML_ASSERT(ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3); GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(ne00 % ggml_blck_size(type) == 0); // rows per thread const int dr = (nr + nth - 1)/nth; @@ -636,8 +647,6 @@ static void ggml_compute_forward_add_q_f32( float * src1_row = (float *)((char *) src1->data + (i11*nb11 + i12*nb12 + i13*nb13)); void * dst_row = (void *) ((char *) dst->data + ( i1*nb1 + i2*nb2 + i3*nb3)); - assert(ne00 % 32 == 0); - // unquantize row from src0 to temp buffer dequantize_row_q(src0_row, wdata, ne00); // add src1 @@ -673,6 +682,7 @@ void ggml_compute_forward_add( case GGML_TYPE_Q8_0: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -955,7 +965,7 @@ static void ggml_compute_forward_add1_q_f32( GGML_ASSERT(nb1 <= nb2); GGML_ASSERT(nb2 <= nb3); - GGML_ASSERT(ggml_is_quantized(src0->type)); + GGML_ASSERT(ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3); GGML_ASSERT(dst->type == src0->type); GGML_ASSERT(src1->type == GGML_TYPE_F32); @@ -1130,6 +1140,7 @@ void ggml_compute_forward_add1( case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -1261,6 +1272,7 @@ void ggml_compute_forward_acc( case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -4804,6 +4816,7 @@ void ggml_compute_forward_out_prod( case GGML_TYPE_Q8_0: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -5129,6 +5142,7 @@ void ggml_compute_forward_set( case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -5354,6 +5368,7 @@ void ggml_compute_forward_get_rows( case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -6114,6 +6129,7 @@ void ggml_compute_forward_clamp( case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_F8_E4M3: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 5e36459f8..44ea9fae5 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -3,6 +3,7 @@ #include "ggml-cpu-impl.h" #include "simd-mappings.h" +#include "vec.h" #include "ggml-quants.h" #include "quants.h" @@ -62,6 +63,10 @@ void quantize_row_nvfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, i quantize_row_nvfp4_ref(x, y, k); } +void quantize_row_f8_e4m3(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_f8_e4m3_ref(x, y, k); +} + // // 2-6 bit quantization in super-blocks // @@ -362,6 +367,49 @@ void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, *s = sumf; } +void ggml_vec_dot_f8_e4m3_f32(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const ggml_fp8_e4m3_t * GGML_RESTRICT x = vx; + const float * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + int pos = 0; +#if defined(GGML_SIMD) && !defined(__ARM_FEATURE_SVE) && !defined(__riscv_v_intrinsic) + GGML_F32_VEC sum[GGML_F32_ARR] = { GGML_F32_VEC_ZERO }; + const int np = n - n % GGML_F32_STEP; + for (; pos < np; pos += GGML_F32_STEP) { + for (int j = 0; j < GGML_F32_ARR; ++j) { + const int offset = pos + j*GGML_F32_EPR; + GGML_F32_VEC ax; +#if defined(__AVX512F__) && defined(__AVX512BW__) + const __m128i codes = _mm_loadu_si128((const __m128i *) (x + offset)); + ax = _mm512_i32gather_ps(_mm512_cvtepu8_epi32(codes), ggml_table_f32_f8_e4m3, sizeof(float)); +#elif defined(__AVX2__) && !defined(__AVX512F__) + const __m128i codes = _mm_loadl_epi64((const __m128i *) (x + offset)); + ax = _mm256_i32gather_ps(ggml_table_f32_f8_e4m3, _mm256_cvtepu8_epi32(codes), sizeof(float)); +#else + float values[GGML_F32_EPR]; + for (int k = 0; k < GGML_F32_EPR; ++k) { + values[k] = ggml_table_f32_f8_e4m3[x[offset + k].bits]; + } + ax = GGML_F32_VEC_LOAD(values); +#endif + sum[j] = GGML_F32_VEC_FMA(sum[j], ax, GGML_F32_VEC_LOAD(y + offset)); + } + } + GGML_F32_VEC_REDUCE(sumf, sum); +#endif + for (; pos < n; ++pos) { + sumf += ggml_table_f32_f8_e4m3[x[pos].bits] * y[pos]; + } + *s = sumf; +} + void ggml_vec_dot_q5_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index 93ea7eeff..f04be52f5 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -23,6 +23,7 @@ void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_mxfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_nvfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_f8_e4m3(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q2_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q3_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -40,6 +41,7 @@ void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, // Dot product void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_f8_e4m3_f32(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q5_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); diff --git a/ggml/src/ggml-cpu/simd-mappings.h b/ggml/src/ggml-cpu/simd-mappings.h index b7b390db4..7b86e8ad9 100644 --- a/ggml/src/ggml-cpu/simd-mappings.h +++ b/ggml/src/ggml-cpu/simd-mappings.h @@ -127,6 +127,8 @@ extern float ggml_table_f32_e8m0_half[1 << 8]; // defined in ggml-cpu.c, initialized in ggml_cpu_init() extern float ggml_table_f32_ue4m3[1 << 8]; +extern float ggml_table_f32_f8_e4m3[1 << 8]; + // Use lookup table for E8M0 on x86 (faster than bit manipulation) #if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) #define GGML_CPU_E8M0_TO_FP32_HALF(x) ggml_table_f32_e8m0_half[(uint8_t)(x)] diff --git a/ggml/src/ggml-impl.h b/ggml/src/ggml-impl.h index ae26e0c23..4dc55e78e 100644 --- a/ggml/src/ggml-impl.h +++ b/ggml/src/ggml-impl.h @@ -512,6 +512,7 @@ static inline float ggml_e8m0_to_fp32_half(uint8_t x) { // UE4M3: unsigned, 4 exp bits (bias=7), 3 mantissa bits // Returns value * 0.5 to match kvalues_mxfp4 convention (kvalues = 2 * E2M1_float) static inline float ggml_ue4m3_to_fp32(uint8_t x) { + // TODO: reject invalid unsigned E4M3 scales instead of decoding NaN as zero or ignoring the sign bit. if (x == 0 || x == 0x7F) { return 0.0f; } @@ -527,6 +528,7 @@ static inline float ggml_ue4m3_to_fp32(uint8_t x) { } static inline uint8_t ggml_fp32_to_ue4m3(float x) { + // TODO: use ties-to-even rounding, allow subnormal carry, and preserve finite values with exponent field 15. if (!(x > 0.0f)) { return 0; } @@ -564,6 +566,62 @@ static inline uint8_t ggml_fp32_to_ue4m3(float x) { return (uint8_t) ((ue4m3_exp << 3) | ue4m3_man); } +static inline float ggml_f8_e4m3_to_fp32(uint8_t x) { + const uint8_t ax = x & 0x7F; + if (ax == 0x7F) { + // OCP FP8 defines both 0x7F and 0xFF as NaN; the sign need not be preserved. + return NAN; + } + + const int exp = (ax >> 3) & 0xF; + const int man = ax & 0x7; + // OCP FP8 rev. 1.1, section 5.1: subnormal magnitude = M * 2^-9; normal magnitude = (1 + M / 8) * 2^(E - 7). + const float value = exp == 0 + ? ldexpf((float) man, -9) + : ldexpf(1.0f + (float) man / 8.0f, exp - 7); + return x & 0x80 ? -value : value; +} + +static inline int ggml_round_to_nearest_even(float x) { + const int value = (int) floorf(x); + const float fraction = x - value; + return fraction > 0.5f || (fraction == 0.5f && (value & 1)) ? value + 1 : value; +} + +// Supports saturating conversion only; infinities and overflow clamp to +/-448. +static inline uint8_t ggml_fp32_to_f8_e4m3(float x) { + const uint8_t sign = signbit(x) ? 0x80 : 0; + x = fabsf(x); + + if (isnan(x)) { + // OCP FP8 leaves NaN sign handling to the implementation; use canonical 0x7F. + return 0x7F; + } + if (x == 0.0f) { + return sign; + } + if (isinf(x) || x >= 448.0f) { + return sign | 0x7E; + } + if (x < 0.015625f) { + // OCP FP8 rev. 1.1, section 5.1: subnormal magnitude = M * 2^(1 - bias - m) = M / 512. + return sign | (uint8_t) ggml_round_to_nearest_even(x * 512.0f); + } + + int exp; + const float mantissa = frexpf(x, &exp) * 2.0f; + int encoded_exp = exp + 6; + int encoded_man = ggml_round_to_nearest_even((mantissa - 1.0f) * 8.0f); + if (encoded_man == 8) { + encoded_man = 0; + encoded_exp++; + } + if (encoded_exp > 15 || (encoded_exp == 15 && encoded_man > 6)) { + return sign | 0x7E; + } + return sign | (uint8_t) (encoded_exp << 3) | (uint8_t) encoded_man; +} + /** * Converts brain16 to float32. * diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 451f65299..f57baa3ca 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1516,6 +1516,7 @@ static bool ggml_metal_supports_mul_mat_op( bool src0_f16_has_mv, bool mm_path) { if (!has_simdgroup_reduction || + op->src[0]->type == GGML_TYPE_F8_E4M3 || op->src[0]->type == GGML_TYPE_NVFP4 || op->src[0]->type == GGML_TYPE_TQ1_0) { return false; @@ -1933,7 +1934,8 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te } case GGML_OP_GET_ROWS: return op->src[0]->type != GGML_TYPE_NVFP4 && - op->src[0]->type != GGML_TYPE_TQ1_0; + op->src[0]->type != GGML_TYPE_TQ1_0 && + op->src[0]->type != GGML_TYPE_F8_E4M3; case GGML_OP_SET_ROWS: { if (op->src[0]->type == GGML_TYPE_F16) { diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 7adec4237..3586f67ae 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -413,6 +413,12 @@ void quantize_row_nvfp4_ref(const float * GGML_RESTRICT x, block_nvfp4 * GGML_RE } } +void quantize_row_f8_e4m3_ref(const float * GGML_RESTRICT x, ggml_fp8_e4m3_t * GGML_RESTRICT y, int64_t k) { + for (int64_t i = 0; i < k; ++i) { + y[i].bits = ggml_fp32_to_f8_e4m3(x[i]); + } +} + void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { static const int qk = QK1_0; @@ -608,6 +614,12 @@ void dequantize_row_nvfp4(const block_nvfp4 * GGML_RESTRICT x, float * GGML_REST } } +void dequantize_row_f8_e4m3(const ggml_fp8_e4m3_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + for (int64_t i = 0; i < k; ++i) { + y[i] = ggml_f8_e4m3_to_fp32(x[i].bits); + } +} + // // 2-6 bit quantization in super-blocks // @@ -2309,6 +2321,12 @@ size_t quantize_nvfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, return nrow * ggml_row_size(GGML_TYPE_NVFP4, n_per_row); } +size_t quantize_f8_e4m3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + GGML_UNUSED(quant_weights); + quantize_row_f8_e4m3_ref(src, dst, (int64_t)nrow*n_per_row); + return nrow * ggml_row_size(GGML_TYPE_F8_E4M3, n_per_row); +} + // ====================== Ternary (de)-quantization (BitNet b1.58 and TriLMs) void quantize_row_tq1_0_ref(const float * GGML_RESTRICT x, block_tq1_0 * GGML_RESTRICT y, int64_t k) { @@ -5533,10 +5551,20 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte } break; case GGML_TYPE_NVFP4: { - // UE4M3 scales are uint8_t — all byte values are valid + // TODO: validate each unsigned E4M3 scale; NaN encodings and scales with the sign bit set are invalid. GGML_UNUSED(data); GGML_UNUSED(nb); } break; + case GGML_TYPE_F8_E4M3: + { + const ggml_fp8_e4m3_t * q = (const ggml_fp8_e4m3_t *) data; + for (size_t i = 0; i < nb; ++i) { + if ((q[i].bits & 0x7F) == 0x7F) { + fprintf(stderr, "%s: found NaN at index %zu in row of %zu F8_E4M3 values\n", __func__, i, nb); + return false; + } + } + } break; case GGML_TYPE_Q2_K: { VALIDATE_ROW_DATA_DM_F16_IMPL(block_q2_K, data, nb, d, dmin); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 75188f1af..ca36d1369 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -25,6 +25,7 @@ GGML_API void quantize_row_q8_1_ref(const float * GGML_RESTRICT x, block_q8_1 * GGML_API void quantize_row_mxfp4_ref(const float * GGML_RESTRICT x, block_mxfp4 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_nvfp4_ref(const float * GGML_RESTRICT x, block_nvfp4 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_f8_e4m3_ref(const float * GGML_RESTRICT x, ggml_fp8_e4m3_t * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q2_K_ref(const float * GGML_RESTRICT x, block_q2_K * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q3_K_ref(const float * GGML_RESTRICT x, block_q3_K * GGML_RESTRICT y, int64_t k); @@ -54,6 +55,7 @@ GGML_API void dequantize_row_q8_0(const block_q8_0 * GGML_RESTRICT x, float * GG GGML_API void dequantize_row_mxfp4(const block_mxfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_nvfp4(const block_nvfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_f8_e4m3(const ggml_fp8_e4m3_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q2_K(const block_q2_K * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q3_K(const block_q3_K * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -104,6 +106,7 @@ GGML_API size_t quantize_q8_0(const float * GGML_RESTRICT src, void * GGML_RESTR GGML_API size_t quantize_mxfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_nvfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_f8_e4m3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API void iq2xs_init_impl(enum ggml_type type); GGML_API void iq2xs_free_impl(enum ggml_type type); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 5d8449a88..e00f8e23a 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -773,6 +773,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_nvfp4, .from_float_ref = (ggml_from_float_t)quantize_row_nvfp4_ref, }, + [GGML_TYPE_F8_E4M3] = { + .type_name = "f8_e4m3", + .blck_size = 1, + .type_size = sizeof(ggml_fp8_e4m3_t), + .is_quantized = false, + .to_float = (ggml_to_float_t) dequantize_row_f8_e4m3, + .from_float_ref = (ggml_from_float_t) quantize_row_f8_e4m3_ref, + }, [GGML_TYPE_Q2_K] = { .type_name = "q2_K", .blck_size = QK_K, @@ -1450,6 +1458,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; case GGML_FTYPE_MOSTLY_MXFP4: wtype = GGML_TYPE_MXFP4; break; case GGML_FTYPE_MOSTLY_NVFP4: wtype = GGML_TYPE_NVFP4; break; + case GGML_FTYPE_MOSTLY_F8_E4M3: wtype = GGML_TYPE_F8_E4M3; break; case GGML_FTYPE_MOSTLY_Q2_K: wtype = GGML_TYPE_Q2_K; break; case GGML_FTYPE_MOSTLY_Q3_K: wtype = GGML_TYPE_Q3_K; break; case GGML_FTYPE_MOSTLY_Q4_K: wtype = GGML_TYPE_Q4_K; break; @@ -8184,6 +8193,7 @@ size_t ggml_quantize_chunk( case GGML_TYPE_Q8_0: result = quantize_q8_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_MXFP4: result = quantize_mxfp4 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_NVFP4: result = quantize_nvfp4 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_F8_E4M3: result = quantize_f8_e4m3(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q2_K: result = quantize_q2_K (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q3_K: result = quantize_q3_K (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_K: result = quantize_q4_K (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 7884b9841..1d67dc8cc 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -6526,6 +6526,7 @@ class GGMLQuantizationType(IntEnum): NVFP4 = 40 Q1_0 = 41 Q2_0 = 42 + F8_E4M3 = 43 class ExpertGatingFuncType(IntEnum): @@ -6582,6 +6583,7 @@ class LlamaFileType(IntEnum): MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors MOSTLY_Q2_0 = 41 # except 1d tensors + MOSTLY_F8_E4M3 = 42 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -6753,6 +6755,7 @@ GGML_QUANT_SIZES: dict[GGMLQuantizationType, tuple[int, int]] = { GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), GGMLQuantizationType.Q2_0: (64, 2 + 16), + GGMLQuantizationType.F8_E4M3: (1, 1), } diff --git a/gguf-py/gguf/quants.py b/gguf-py/gguf/quants.py index 359e279f0..c7f97e512 100644 --- a/gguf-py/gguf/quants.py +++ b/gguf-py/gguf/quants.py @@ -818,6 +818,23 @@ class NVFP4(__Quant, qtype=GGMLQuantizationType.NVFP4): return (d * vals.astype(np.float32)).reshape(n_super, 64) +class F8_E4M3(__Quant, qtype=GGMLQuantizationType.F8_E4M3): + @classmethod + def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: + bits = blocks.astype(np.uint8) + sign = np.where(bits & 0x80, np.float32(-1.0), np.float32(1.0)) + magnitude = bits & 0x7F + exponent = magnitude >> 3 + mantissa = magnitude & 0x07 + values = np.where( + exponent == 0, + np.ldexp(mantissa.astype(np.float32), -9), + np.ldexp(1.0 + mantissa.astype(np.float32) / 8.0, exponent.astype(np.int32) - 7), + ) + values = np.where(magnitude == 0x7F, np.nan, sign * values) + return values.astype(np.float32, copy=False) + + class IQ2_XXS(__Quant, qtype=GGMLQuantizationType.IQ2_XXS): ksigns: bytes = ( b"\x00\x81\x82\x03\x84\x05\x06\x87\x88\x09\x0a\x8b\x0c\x8d\x8e\x0f" diff --git a/gguf-py/gguf/scripts/gguf_convert_endian.py b/gguf-py/gguf/scripts/gguf_convert_endian.py index 3b20998c6..7e8b2f1be 100755 --- a/gguf-py/gguf/scripts/gguf_convert_endian.py +++ b/gguf-py/gguf/scripts/gguf_convert_endian.py @@ -40,6 +40,7 @@ byteswap_tensors: dict[gguf.GGMLQuantizationType, tuple[tuple[int, type[np.unsig gguf.GGMLQuantizationType.TQ2_0: ((64, np.uint16),), # d gguf.GGMLQuantizationType.MXFP4: (), # noop gguf.GGMLQuantizationType.NVFP4: (), # noop + gguf.GGMLQuantizationType.F8_E4M3: (), # noop } diff --git a/gguf-py/tests/test_quants.py b/gguf-py/tests/test_quants.py index 0d7729f95..054e019c7 100755 --- a/gguf-py/tests/test_quants.py +++ b/gguf-py/tests/test_quants.py @@ -70,6 +70,7 @@ class GGMLQuants: "mxfp4", "nvfp4", "q1_0", "q2_0", + "f8_e4m3", "iq2_xxs", "iq2_xs", "iq2_s", "iq3_xxs", "iq3_s", "iq1_s", "iq1_m", "iq4_nl", "iq4_xs", ): @@ -181,6 +182,19 @@ def do_test(libggml_path: Path, quick: bool = False, user_type: GGMLQuantization logger.info(f"Testing {qtype.name}") + if qtype == GGMLQuantizationType.F8_E4M3: + encodings = np.arange(256, dtype=np.uint8) + pydq = gguf.dequantize(encodings, qtype) + ggdq = ggml_quants.dequantize(encodings, qtype) + assert pydq.dtype == np.float32 + nan_mask = (encodings & 0x7F) == 0x7F + # Both 0x7F and 0xFF are NaNs; check classification, not FP32 NaN sign or payload. + np.testing.assert_array_equal(np.isnan(pydq), nan_mask) + np.testing.assert_array_equal(np.isnan(ggdq), nan_mask) + # Compare finite values by bits to also check signed zeros. + np.testing.assert_array_equal(pydq[~nan_mask].view(np.uint32), ggdq[~nan_mask].view(np.uint32)) + logger.info("All 256 FP8 E4M3 encodings match C") + rc = r.copy(order="C") pyq = None diff --git a/include/llama.h b/include/llama.h index ae54e661b..3877b69e2 100644 --- a/include/llama.h +++ b/include/llama.h @@ -158,6 +158,7 @@ extern "C" { LLAMA_FTYPE_MOSTLY_NVFP4 = 39, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q1_0 = 40, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors + LLAMA_FTYPE_MOSTLY_F8_E4M3 = 42, // except 1d tensors LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 45bed2a42..7a942663d 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1965,12 +1965,16 @@ ggml_tensor * llm_graph_context::build_ffn( return false; }; - GGML_ASSERT(!up_s || !up_b || !up || up->type != GGML_TYPE_NVFP4); - GGML_ASSERT(!gate_s || !gate_b || !gate || gate->type != GGML_TYPE_NVFP4); - GGML_ASSERT(!down_s || !down_b || !down || down->type != GGML_TYPE_NVFP4); - GGML_ASSERT(!up_s || !up || up->type != GGML_TYPE_NVFP4 || !has_lora(up)); - GGML_ASSERT(!gate_s || !gate || gate->type != GGML_TYPE_NVFP4 || !has_lora(gate)); - GGML_ASSERT(!down_s || !down || down->type != GGML_TYPE_NVFP4 || !has_lora(down)); + auto is_scaled_low_precision = [](ggml_tensor * w) { + return w && (w->type == GGML_TYPE_NVFP4 || w->type == GGML_TYPE_F8_E4M3); + }; + + GGML_ASSERT(!up_s || !up_b || !is_scaled_low_precision(up)); + GGML_ASSERT(!gate_s || !gate_b || !is_scaled_low_precision(gate)); + GGML_ASSERT(!down_s || !down_b || !is_scaled_low_precision(down)); + GGML_ASSERT(!up_s || !is_scaled_low_precision(up) || !has_lora(up)); + GGML_ASSERT(!gate_s || !is_scaled_low_precision(gate) || !has_lora(gate)); + GGML_ASSERT(!down_s || !is_scaled_low_precision(down) || !has_lora(down)); ggml_tensor * tmp = up ? build_lora_mm(up, cur) : cur; cb(tmp, "ffn_up", il); diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index a42992673..485483ba1 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -40,6 +40,7 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_BF16: name = LLAMA_FTYPE_PREFIX "BF16"; break; case LLAMA_FTYPE_MOSTLY_Q1_0: name = LLAMA_FTYPE_PREFIX "Q1_0"; break; case LLAMA_FTYPE_MOSTLY_Q2_0: name = LLAMA_FTYPE_PREFIX "Q2_0"; break; + case LLAMA_FTYPE_MOSTLY_F8_E4M3: name = LLAMA_FTYPE_PREFIX "F8_E4M3"; break; case LLAMA_FTYPE_MOSTLY_Q4_0: name = LLAMA_FTYPE_PREFIX "Q4_0"; break; case LLAMA_FTYPE_MOSTLY_Q4_1: name = LLAMA_FTYPE_PREFIX "Q4_1"; break; case LLAMA_FTYPE_MOSTLY_Q5_0: name = LLAMA_FTYPE_PREFIX "Q5_0"; break; @@ -783,6 +784,7 @@ llama_model_loader::llama_model_loader( case GGML_TYPE_NVFP4: ftype = LLAMA_FTYPE_MOSTLY_NVFP4; break; case GGML_TYPE_Q1_0: ftype = LLAMA_FTYPE_MOSTLY_Q1_0; break; case GGML_TYPE_Q2_0: ftype = LLAMA_FTYPE_MOSTLY_Q2_0; break; + case GGML_TYPE_F8_E4M3: ftype = LLAMA_FTYPE_MOSTLY_F8_E4M3; break; default: { LLAMA_LOG_WARN("%s: unknown type %s\n", __func__, ggml_type_name(type_max)); diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 5fa3dfad9..83aed0144 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1874,13 +1874,15 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { } } // output scales - if (output && output->type == GGML_TYPE_NVFP4) { + if (output) { // weight scale - if (!output_s) { + if (!output_s && (ml.get_weight(tn(LLM_TENSOR_OUTPUT, "scale").str().c_str()) || + gguf_find_tensor(ml.metadata, tn(LLM_TENSOR_OUTPUT, "scale").str().c_str()) >= 0)) { output_s = create_tensor(tn(LLM_TENSOR_OUTPUT, "scale"), {1}, TENSOR_NOT_REQUIRED); } // input scale - if (!output_in_s) { + if (!output_in_s && (ml.get_weight(tn(LLM_TENSOR_OUTPUT, "input_scale").str().c_str()) || + gguf_find_tensor(ml.metadata, tn(LLM_TENSOR_OUTPUT, "input_scale").str().c_str()) >= 0)) { output_in_s = create_tensor(tn(LLM_TENSOR_OUTPUT, "input_scale"), {1}, TENSOR_NOT_REQUIRED); } } @@ -1894,12 +1896,11 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { can_prefetch.insert(per_layer_tok_embd); } - // Tied NVFP4 output is valid when no separate LM-head scale tensors are present. - // If sidecar scales exist, the output weight must be an actual output tensor. - GGML_ASSERT(!(output && tok_embd && + if (output && tok_embd && strcmp(output->name, tok_embd->name) == 0 && - output->type == GGML_TYPE_NVFP4 && - (output_s || output_in_s))); + (output_s || output_in_s)) { + throw std::runtime_error("output scales require a separate output weight"); + } // populate tensors_by_name for (auto & [_, ctx_ptr] : ml.ctx_map) { for (auto * cur = ggml_get_first_tensor(ctx_ptr.get()); cur != NULL; cur = ggml_get_next_tensor(ctx_ptr.get(), cur)) { diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 9f8aa8984..85da3831e 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -233,7 +233,7 @@ static void llama_tensor_dequantize_impl( const size_t nelements, const int nthread ) { const ggml_type_traits * qtype = ggml_get_type_traits(type); - if (ggml_is_quantized(type)) { + if (ggml_is_quantized(type) || type == GGML_TYPE_F8_E4M3) { if (qtype->to_float == NULL) { throw std::runtime_error(format("type %s unsupported for integer quantization: no dequantization available", ggml_type_name(type))); } @@ -247,7 +247,7 @@ static void llama_tensor_dequantize_impl( ggml_fp16_to_fp32_row((const ggml_fp16_t *)data, f32_output, nelements); } else if (type == GGML_TYPE_BF16) { ggml_bf16_to_fp32_row((const ggml_bf16_t *)data, f32_output, nelements); - } else if (ggml_is_quantized(type)) { + } else if (ggml_is_quantized(type) || type == GGML_TYPE_F8_E4M3) { qtype->to_float(data, f32_output, nelements); } else { GGML_ABORT("fatal error"); // unreachable diff --git a/src/models/qwen3next.cpp b/src/models/qwen3next.cpp index 98fafd917..fd082cc7f 100644 --- a/src/models/qwen3next.cpp +++ b/src/models/qwen3next.cpp @@ -321,11 +321,11 @@ std::pair llama_model_qwen3next::graph::build_qkvz if (model.layers[il].wqkv) { // optimized path - ggml_tensor * qkv_mixed = build_lora_mm(model.layers[il].wqkv, input); + ggml_tensor * qkv_mixed = build_lora_mm(model.layers[il].wqkv, input, model.layers[il].wqkv_s); qkv_mixed = ggml_reshape_3d(ctx0, qkv_mixed, qkv_mixed->ne[0], n_seq_tokens, n_seqs); cb(qkv_mixed, "linear_attn_qkv_mixed", il); - ggml_tensor * z = build_lora_mm(model.layers[il].wqkv_gate, input); + ggml_tensor * z = build_lora_mm(model.layers[il].wqkv_gate, input, model.layers[il].wqkv_gate_s); cb(z, "z", il); return { qkv_mixed, z }; diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 2e768cb38..44cc6e5dd 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -226,6 +226,17 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS) endif() llama_build(test-llama-archs.cpp) +if (GGML_CPU AND (NOT WIN32 OR NOT BUILD_SHARED_LIBS)) + llama_build(test-fp8.cpp) + set(FP8_MODEL_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-models-fp8") + file(MAKE_DIRECTORY "${FP8_MODEL_DIR}") + llama_test(test-llama-archs NAME test-generate-fp8-models LABEL "main behavior backend" + ARGS -a llama -s 1 -o "${FP8_MODEL_DIR}") + set_tests_properties(test-generate-fp8-models PROPERTIES FIXTURES_SETUP fp8-models) + llama_test(test-fp8 NAME test-fp8-requantize LABEL "main behavior backend" + ARGS --models "${FP8_MODEL_DIR}") + set_tests_properties(test-fp8-requantize PROPERTIES FIXTURES_REQUIRED fp8-models) +endif() set(MODEL_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-models/") file(MAKE_DIRECTORY "${MODEL_DIR}") @@ -277,6 +288,12 @@ llama_build(test-fusion.cpp) llama_build_and_test(test-chat-peg-parser.cpp peg-parser/simple-tokenize.cpp LABEL behavior) llama_build_and_test(test-jinja.cpp LABEL behavior) llama_test(test-jinja NAME test-jinja-py ARGS -py LABEL "python behavior") +llama_test_cmd(python3 + NAME test-conversion-fp8 + LABEL "python behavior" + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + ARGS -m unittest discover -s tests -p test_conversion_fp8.py -v +) llama_build_and_test(test-json-schema.cpp LABEL behavior) llama_build_and_test(test-chat-auto-parser.cpp LABEL behavior WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}) llama_build_and_test(test-chat-template.cpp LABEL behavior) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index a7c1d61db..f75ff11b9 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -88,7 +88,7 @@ static void init_tensor_uniform(ggml_tensor * tensor, float min = -1.0f, float m if (tensor->type == GGML_TYPE_F32 || tensor->type == GGML_TYPE_I32) { ggml_backend_tensor_set(tensor, data.data(), 0, nels * sizeof(float)); - } else if (ggml_is_quantized(tensor->type) || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) { + } else if (ggml_is_quantized(tensor->type) || tensor->type == GGML_TYPE_F8_E4M3 || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) { GGML_ASSERT(nels % ggml_blck_size(tensor->type) == 0); // dummy importance matrix @@ -262,7 +262,7 @@ static std::vector tensor_to_float(const ggml_tensor * t) { const auto * tt = ggml_get_type_traits(t->type); size_t bs = ggml_blck_size(t->type); std::vector vq(ggml_blck_size(t->type)); - bool quantized = ggml_is_quantized(t->type); + bool uses_type_conversion = ggml_is_quantized(t->type) || t->type == GGML_TYPE_F8_E4M3; // access elements by index to avoid gaps in views for (int64_t i3 = 0; i3 < t->ne[3]; i3++) { @@ -284,7 +284,7 @@ static std::vector tensor_to_float(const ggml_tensor * t) { tv.push_back((float)*(int16_t *) &buf[i]); } else if (t->type == GGML_TYPE_I8) { tv.push_back((float)*(int8_t *) &buf[i]); - } else if (quantized) { + } else if (uses_type_conversion) { tt->to_float(&buf[i], vq.data(), bs); tv.insert(tv.end(), vq.begin(), vq.end()); } else { @@ -5411,6 +5411,9 @@ struct test_mul_mat : public test_case { backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { return 2e-2; } + if (type_a == GGML_TYPE_F8_E4M3 && backend_has_feature(backend, "NATIVE_FP8")) { + return 5e-3; + } return max_nmse_err(); } @@ -5670,6 +5673,9 @@ struct test_mul_mat_id : public test_case { backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { return 2e-2; } + if (type_a == GGML_TYPE_F8_E4M3 && backend_has_feature(backend, "NATIVE_FP8")) { + return 5e-3; + } return max_nmse_err(); } @@ -9978,6 +9984,7 @@ static const ggml_type all_types[] = { GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, GGML_TYPE_MXFP4, GGML_TYPE_NVFP4, + GGML_TYPE_F8_E4M3, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, @@ -9997,6 +10004,7 @@ static const ggml_type base_types[] = { GGML_TYPE_Q4_1, // for I8MM tests GGML_TYPE_Q4_K, GGML_TYPE_MXFP4, GGML_TYPE_NVFP4, // TODO: or "other" + GGML_TYPE_F8_E4M3, GGML_TYPE_IQ2_XXS }; @@ -10730,6 +10738,19 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {-1,-1,-1,-1}, {1, 0, 2, 3})); // cpy not-contiguous } } + for (ggml_type type_src : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16}) { + test_cases.emplace_back(new test_cpy(type_src, GGML_TYPE_F8_E4M3, {32, 2, 3, 4}, {-1,-1,-1,-1}, {0, 0, 0, 0}, {0, 0, 0, 0}, false, {64, 2, 3, 4})); + test_cases.emplace_back(new test_cpy(type_src, GGML_TYPE_F8_E4M3, {32, 2, 3, 4}, {-1,-1,-1,-1}, {1, 0, 2, 3}, {1, 0, 2, 3})); + test_cases.emplace_back(new test_cpy(type_src, GGML_TYPE_F8_E4M3, {32, 2, 3, 4}, {16, 4, 3, 4}, {1, 0, 2, 3}, {0, 2, 1, 3})); + } + + for (ggml_type type_dst : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16}) { + test_cases.emplace_back(new test_cpy(GGML_TYPE_F8_E4M3, type_dst, {33, 2, 3, 4})); + test_cases.emplace_back(new test_cpy(GGML_TYPE_F8_E4M3, type_dst, {32, 2, 3, 4}, {-1,-1,-1,-1}, {0, 0, 0, 0}, {0, 0, 0, 0}, false, {64, 2, 3, 4})); + test_cases.emplace_back(new test_cpy(GGML_TYPE_F8_E4M3, type_dst, {32, 2, 3, 4}, {-1,-1,-1,-1}, {1, 0, 2, 3}, {1, 0, 2, 3})); + test_cases.emplace_back(new test_cpy(GGML_TYPE_F8_E4M3, type_dst, {32, 2, 3, 4}, {16, 4, 3, 4}, {1, 0, 2, 3}, {0, 2, 1, 3})); + } + // quant block count not a multiple of the kernel block size test_cases.emplace_back(new test_cpy(GGML_TYPE_F32, GGML_TYPE_Q4_0, {96, 1, 1, 1})); test_cases.emplace_back(new test_cpy(GGML_TYPE_Q4_0, GGML_TYPE_F32, {96, 1, 1, 1})); @@ -10861,6 +10882,12 @@ static std::vector> make_test_cases_eval() { //add_test_bin_bcast(type, {3, 3, 2560, 1280}, {2, 1, 1, 1}); } + for (auto op : {ggml_add, ggml_add_inplace}) { + for (int64_t n : {1, 7, 33}) { + test_cases.emplace_back(new test_bin_bcast(op, GGML_TYPE_F8_E4M3, {n, 5, 2, 1}, {1, 1, 1, 1}, 1, false, false, GGML_TYPE_F32)); + } + } + for (auto op : {ggml_add, ggml_mul}) { test_cases.emplace_back(new test_bin_bcast(op, GGML_TYPE_F32, {1, 2, 2, 2}, {17, 1, 1, 1}, 1, false, false, GGML_TYPE_COUNT, true)); } @@ -12378,7 +12405,7 @@ static std::vector> make_test_cases_eval() { continue; } for (bool with_lane_scale : {false, true}) { - if (with_lane_scale && type != GGML_TYPE_NVFP4) { + if (with_lane_scale && type != GGML_TYPE_NVFP4 && type != GGML_TYPE_F8_E4M3) { continue; } test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, diff --git a/tests/test-fp8.cpp b/tests/test-fp8.cpp new file mode 100644 index 000000000..a92787836 --- /dev/null +++ b/tests/test-fp8.cpp @@ -0,0 +1,379 @@ +#include "common.h" +#include "log.h" +#include "ggml.h" +#include "ggml-backend.h" +#include "ggml-cpu.h" +#include "ggml-cpp.h" +#include "gguf.h" +#include "llama.h" +#include "llama-cpp.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// normalized mean squared error = mse(a, b) / mse(a, 0) +static double nmse(const std::vector & a, const std::vector & b) { + GGML_ASSERT(a.size() == b.size()); + double mse_a_b = 0.0; + double mse_a_0 = 0.0; + + for (size_t i = 0; i < a.size(); i++) { + float a_i = a[i]; + float b_i = b[i]; + + mse_a_b += (a_i - b_i) * (a_i - b_i); + mse_a_0 += a_i * a_i; + } + + return mse_a_b / mse_a_0; +} + +static std::vector get_tokens(const uint32_t n_tokens, const uint32_t n_vocab, const size_t seed) { + std::mt19937 gen(seed); + std::uniform_int_distribution<> dis(0, n_vocab - 1); + std::vector ret; + ret.reserve(n_tokens); + for (uint32_t i = 0; i < n_tokens; i++) { + ret.push_back(dis(gen)); + } + return ret; +} + +static std::vector get_logits( + llama_model * model, llama_context * lctx, const std::vector & tokens) { + const uint32_t n_vocab = llama_vocab_n_tokens(llama_model_get_vocab(model)); + const uint32_t n_ctx = llama_n_ctx(lctx); + const uint32_t n_tokens = tokens.size(); + common_batch batch(lctx); + GGML_ASSERT(n_tokens <= n_ctx); + for (uint32_t pos = 0; pos < n_tokens; pos++) { + batch.add(tokens[pos], pos, 0, true); + } + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch.get())) { + throw std::runtime_error("failed to decode batch"); + } + + std::vector ret; + ret.reserve(n_tokens*n_vocab); + for (uint32_t i = 0; i < n_tokens; i++) { + const float * logits_ith = llama_get_logits_ith(lctx, i); + for (uint32_t j = 0; j < n_vocab; j++) { + ret.push_back(logits_ith[j]); + } + } + return ret; +} + +static const char * FP8_TEST_WEIGHTS[] = { + "output.weight", "blk.0.attn_q.weight", + "blk.0.ffn_gate_exps.weight", "blk.0.ffn_up_exps.weight", "blk.0.ffn_down_exps.weight", +}; + +static int64_t prepare_models(const std::string & path, const std::string & input, const std::string & reference) { + ggml_context * ctx_raw = nullptr; + gguf_context_ptr source(gguf_init_from_file(path.c_str(), {false, &ctx_raw})); + ggml_context_ptr source_ctx(ctx_raw); + if (!source || !source_ctx) { + throw std::runtime_error("failed to read FP8 test fixture"); + } + ggml_tensor * attn_q = ggml_get_tensor(source_ctx.get(), "blk.0.attn_q.weight"); + if (!attn_q || !ggml_get_tensor(source_ctx.get(), "output.weight")) { + throw std::runtime_error("FP8 test fixture must have attention Q and output weights"); + } + gguf_context_ptr fp8(gguf_init_empty()); + gguf_context_ptr folded(gguf_init_empty()); + gguf_set_kv(fp8.get(), source.get()); + gguf_set_kv(folded.get(), source.get()); + for (int64_t i = 0; i < gguf_get_n_tensors(source.get()); ++i) { + ggml_tensor * tensor = ggml_get_tensor(source_ctx.get(), gguf_get_tensor_name(source.get(), i)); + gguf_add_tensor(fp8.get(), tensor); + gguf_add_tensor(folded.get(), tensor); + } + + ggml_context_ptr scale_ctx(ggml_init({16 * ggml_tensor_overhead(), nullptr, true})); + // GGUF keeps pointers into these buffers until the files are written. + std::vector> fp8_data; + std::vector> folded_data; + std::vector> scale_data; + for (const char * name : FP8_TEST_WEIGHTS) { + ggml_tensor * tensor = ggml_get_tensor(source_ctx.get(), name); + if (!tensor) { + continue; + } + std::vector values(ggml_nelements(tensor)); + if (tensor->type == GGML_TYPE_F32) { + memcpy(values.data(), tensor->data, ggml_nbytes(tensor)); + } else { + const auto * traits = ggml_get_type_traits(tensor->type); + if (!traits->to_float) { + throw std::runtime_error("unsupported FP8 test fixture tensor type"); + } + traits->to_float(tensor->data, values.data(), values.size()); + } + const auto * traits = ggml_get_type_traits(GGML_TYPE_F8_E4M3); + fp8_data.emplace_back(values.size()); + traits->from_float_ref(values.data(), fp8_data.back().data(), values.size()); + gguf_set_tensor_type(fp8.get(), name, GGML_TYPE_F8_E4M3); + gguf_set_tensor_data(fp8.get(), name, fp8_data.back().data()); + + traits->to_float(fp8_data.back().data(), values.data(), values.size()); + const size_t expert_size = tensor->ne[0] * tensor->ne[1]; + for (size_t i = 0; i < values.size(); ++i) { + values[i] *= i < expert_size ? 2.0f : 4.0f; + } + folded_data.emplace_back(values.size()); + ggml_fp32_to_fp16_row(values.data(), folded_data.back().data(), values.size()); + gguf_set_tensor_type(folded.get(), name, GGML_TYPE_F16); + gguf_set_tensor_data(folded.get(), name, folded_data.back().data()); + + const std::string prefix(name, strlen(name) - strlen("weight")); + for (const char * suffix : {"scale", "input_scale"}) { + const std::string scale_name = prefix + suffix; + ggml_tensor * scale = ggml_get_tensor(source_ctx.get(), scale_name.c_str()); + if (scale) { + if (scale->type != GGML_TYPE_F32 || ggml_nelements(scale) != tensor->ne[2]) { + throw std::runtime_error("unsupported FP8 test fixture scale shape or type"); + } + scale_data.emplace_back(tensor->ne[2], 1.0f); + gguf_set_tensor_data(folded.get(), scale_name.c_str(), scale_data.back().data()); + } else { + scale = ggml_new_tensor_1d(scale_ctx.get(), GGML_TYPE_F32, tensor->ne[2]); + ggml_set_name(scale, scale_name.c_str()); + gguf_add_tensor(fp8.get(), scale); + } + // Scales above one avoid subnormal Q6_K block scales in the folded reference. + scale_data.emplace_back(tensor->ne[2], 4.0f); + if (strcmp(suffix, "scale") == 0) { + scale_data.back()[0] = 2.0f; + } + gguf_set_tensor_data(fp8.get(), scale_name.c_str(), scale_data.back().data()); + } + } + if (!gguf_write_to_file(fp8.get(), input.c_str(), false) || + !gguf_write_to_file(folded.get(), reference.c_str(), false)) { + throw std::runtime_error("failed to write FP8 test models"); + } + return attn_q->ne[0]; +} + +static bool silent_model_load_progress(float /*progress*/, void * /*user_data*/) { + return true; +} + +static llama_model_ptr load_model(const std::string & path, const llama_model_params & params, FILE * file = nullptr) { + llama_model_ptr model(file ? llama_model_load_from_file_ptr(file, params) : + llama_model_load_from_file(path.c_str(), params)); + if (!model) { + throw std::runtime_error("failed to load FP8 test model"); + } + return model; +} + +static std::vector load_logits(const std::string & path, const llama_model_params & model_params, bool split = false) { + std::unique_ptr file(split ? nullptr : fopen(path.c_str(), "rb"), fclose); + if (!split && !file) { + throw std::runtime_error("failed to open FP8 test model"); + } + llama_model_ptr model = load_model(path, model_params, file.get()); + auto ctx_params = llama_context_default_params(); + ctx_params.n_ctx = 0; + ctx_params.n_ubatch = 64; + ctx_params.n_threads = ctx_params.n_threads_batch = 4; + llama_context_ptr lctx(llama_init_from_model(model.get(), ctx_params)); + if (!lctx) { + throw std::runtime_error("failed to create FP8 test context"); + } + const auto tokens = get_tokens(8, llama_vocab_n_tokens(llama_model_get_vocab(model.get())), 1); + return get_logits(model.get(), lctx.get(), tokens); +} + +static void check_model_save(const std::string & path, const llama_model_params & model_params) { + const std::string saved = path + ".saved.gguf"; + const auto logits = load_logits(path, model_params); + auto save_params = model_params; + // Packed CPU buffers do not support reading weights back for saving. + save_params.use_extra_bufts = false; + save_params.load_mode = LLAMA_LOAD_MODE_NONE; + { + llama_model_ptr model = load_model(path, save_params); + llama_model_save_to_file(model.get(), saved.c_str()); + } + const double error = nmse(logits, load_logits(saved, model_params)); + LOG("FP8 model save: %s nmse=%.2e\n", path.c_str(), error); + if (!(error < 1e-10)) { + throw std::runtime_error("saved FP8 test model differs from original model"); + } + std::filesystem::rename(saved, path); +} + +static void check_fp8_scales(const std::string & input, const std::string & reference, const llama_model_params & model_params) { + const auto reference_logits = load_logits(reference, model_params); + const double source_error = nmse(reference_logits, load_logits(input, model_params)); + LOG("FP8 sidecar scales: nmse=%.2e\n", source_error); + if (!(source_error < 1e-5)) { + throw std::runtime_error("FP8 sidecar scales differ from folded weights"); + } +} + +static void check_split_model(const std::string & path, const std::filesystem::path & dir, const llama_model_params & model_params) { + ggml_context * ctx_raw = nullptr; + gguf_context_ptr metadata(gguf_init_from_file(path.c_str(), {false, &ctx_raw})); + ggml_context_ptr ctx(ctx_raw); + if (!metadata || !ctx) { + throw std::runtime_error("failed to read FP8 test model for splitting"); + } + const auto split_first = (dir / "split-00001-of-00002.gguf").string(); + const auto split_second = (dir / "split-00002-of-00002.gguf").string(); + for (int split = 0; split < 2; ++split) { + gguf_context_ptr shard(gguf_init_empty()); + gguf_set_kv(shard.get(), metadata.get()); + gguf_set_val_u16(shard.get(), "split.no", split); + gguf_set_val_u16(shard.get(), "split.count", 2); + gguf_set_val_i32(shard.get(), "split.tensors.count", gguf_get_n_tensors(metadata.get())); + for (int64_t i = 0; i < gguf_get_n_tensors(metadata.get()); ++i) { + auto * tensor = ggml_get_tensor(ctx.get(), gguf_get_tensor_name(metadata.get(), i)); + // Put output scales in the second shard to test lookup across splits. + const bool output_scale = strcmp(tensor->name, "output.scale") == 0 || strcmp(tensor->name, "output.input_scale") == 0; + if (output_scale == (split == 1)) { + gguf_add_tensor(shard.get(), tensor); + } + } + if (!gguf_write_to_file(shard.get(), (split == 0 ? split_first : split_second).c_str(), false)) { + throw std::runtime_error("failed to write split FP8 test model"); + } + } + const double split_error = nmse(load_logits(path, model_params), load_logits(split_first, model_params, true)); + LOG("FP8 split model: nmse=%.2e\n", split_error); + if (!(split_error < 1e-10)) { + throw std::runtime_error("split FP8 model differs from unsplit model"); + } +} + +static void check_fp8_requantize( + const std::string & input, const std::string & reference, + const std::string & output, const std::string & reference_output, + int64_t n_embd, const llama_model_params & model_params) { + for (llama_ftype ftype : {LLAMA_FTYPE_ALL_F32, LLAMA_FTYPE_MOSTLY_F16, LLAMA_FTYPE_MOSTLY_BF16, LLAMA_FTYPE_MOSTLY_Q4_K_M}) { + for (int nthread : {1, 3}) { + std::vector importance(n_embd, 1.0f); + llama_model_imatrix_data imatrix[] = { + {"blk.0.attn_q.weight", importance.data(), importance.size()}, + {nullptr, nullptr, 0}, + }; + auto params = llama_model_quantize_default_params(); + params.ftype = ftype; + params.nthread = nthread; + if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M && nthread == 3) { + params.imatrix = imatrix; + } + if (llama_model_quantize(input.c_str(), output.c_str(), ¶ms) != 0 || + llama_model_quantize(reference.c_str(), reference_output.c_str(), ¶ms) != 0) { + throw std::runtime_error("FP8 model requantization failed"); + } + const double error = nmse(load_logits(reference_output, model_params), load_logits(output, model_params)); + LOG("FP8 requantization: ftype=%d threads=%d nmse=%.2e\n", ftype, nthread, error); + if (!(error < 1e-4)) { + throw std::runtime_error("FP8 requantization differs from folded weights"); + } + } + } +} + +static void test_fp8(const std::string & model_path) { + const auto dir = std::filesystem::temp_directory_path() / ("llama-fp8-" + std::to_string(std::random_device{}())); + if (!std::filesystem::create_directory(dir)) { + throw std::runtime_error("failed to create FP8 test directory"); + } + struct test_files { + std::filesystem::path dir; + ~test_files() { + std::error_code ec; + std::filesystem::remove_all(dir, ec); + } + } files{dir}; + + const auto input = (files.dir / "fp8.gguf").string(); + const auto reference = (files.dir / "reference.gguf").string(); + const auto output = (files.dir / "output.gguf").string(); + const auto reference_output = (files.dir / "reference-output.gguf").string(); + + // Use the same CPU kernels for scaled weights and the folded reference. + const llama_model_tensor_buft_override overrides[] = { + {".*", ggml_backend_cpu_buffer_type()}, + {nullptr, nullptr}, + }; + ggml_backend_dev_t devices[] = {nullptr}; + auto model_params = llama_model_default_params(); + model_params.devices = devices; + model_params.tensor_buft_overrides = overrides; + model_params.progress_callback = silent_model_load_progress; + + const int64_t n_embd = prepare_models(model_path, input, reference); + check_model_save(input, model_params); + check_model_save(reference, model_params); + check_fp8_scales(input, reference, model_params); + check_split_model(reference, files.dir, model_params); + check_split_model(input, files.dir, model_params); + check_fp8_requantize(input, reference, output, reference_output, n_embd, model_params); +} + +int main(int argc, char ** argv) { + if (argc == 2 && (strcmp(argv[1], "-h") == 0 || strcmp(argv[1], "--help") == 0)) { + printf("Usage: %s -m model.gguf | --models DIR\n", argv[0]); + return 0; + } + if (argc != 3 || (strcmp(argv[1], "-m") != 0 && strcmp(argv[1], "--model") != 0 && strcmp(argv[1], "--models") != 0)) { + fprintf(stderr, "Usage: %s -m model.gguf | --models DIR\n", argv[0]); + return 1; + } + common_init(); + llama_backend_init(); + int result = 0; + try { + std::vector models; + if (strcmp(argv[1], "--models") == 0) { + if (!std::filesystem::is_directory(argv[2])) { + throw std::runtime_error("models directory does not exist: " + std::string(argv[2])); + } + for (const auto & entry : std::filesystem::directory_iterator(argv[2])) { + if (entry.is_regular_file() && entry.path().extension() == ".gguf") { + models.push_back(entry.path().string()); + } + } + std::sort(models.begin(), models.end()); + if (models.empty()) { + throw std::runtime_error("no .gguf models found in: " + std::string(argv[2])); + } + } else { + models.push_back(argv[2]); + } + + std::vector passed; + for (const auto & model : models) { + try { + test_fp8(model); + passed.push_back(true); + } catch (const std::exception & err) { + fprintf(stderr, "FP8 test failed for '%s': %s\n", model.c_str(), err.what()); + passed.push_back(false); + result = 1; + } + } + printf("\n| Model | Result |\n|-------|--------|\n"); + for (size_t i = 0; i < models.size(); ++i) { + printf("| %s | %s |\n", std::filesystem::path(models[i]).filename().string().c_str(), passed[i] ? "PASS" : "FAIL"); + } + } catch (const std::exception & err) { + fprintf(stderr, "FP8 test failed: %s\n", err.what()); + result = 1; + } + llama_backend_free(); + return result; +} diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp index 0479e1020..d82d6db79 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index 1930a09b2..ab88a237f 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -7,6 +7,7 @@ #include #include #include +#include #include #include #include @@ -155,6 +156,152 @@ static int test_vec_dot_f32(bool verbose) { return num_failed; } +// Positive finite E4M3 values in encoding order, from OCP FP8 rev. 1.1, section 5.1. +static const float F8_E4M3_VALUES[] = { + 0.0f, 0.001953125f, 0.00390625f, 0.005859375f, 0.0078125f, 0.009765625f, 0.01171875f, 0.013671875f, + 0.015625f, 0.017578125f, 0.01953125f, 0.021484375f, 0.0234375f, 0.025390625f, 0.02734375f, 0.029296875f, + 0.03125f, 0.03515625f, 0.0390625f, 0.04296875f, 0.046875f, 0.05078125f, 0.0546875f, 0.05859375f, + 0.0625f, 0.0703125f, 0.078125f, 0.0859375f, 0.09375f, 0.1015625f, 0.109375f, 0.1171875f, + 0.125f, 0.140625f, 0.15625f, 0.171875f, 0.1875f, 0.203125f, 0.21875f, 0.234375f, + 0.25f, 0.28125f, 0.3125f, 0.34375f, 0.375f, 0.40625f, 0.4375f, 0.46875f, + 0.5f, 0.5625f, 0.625f, 0.6875f, 0.75f, 0.8125f, 0.875f, 0.9375f, + 1.0f, 1.125f, 1.25f, 1.375f, 1.5f, 1.625f, 1.75f, 1.875f, + 2.0f, 2.25f, 2.5f, 2.75f, 3.0f, 3.25f, 3.5f, 3.75f, + 4.0f, 4.5f, 5.0f, 5.5f, 6.0f, 6.5f, 7.0f, 7.5f, + 8.0f, 9.0f, 10.0f, 11.0f, 12.0f, 13.0f, 14.0f, 15.0f, + 16.0f, 18.0f, 20.0f, 22.0f, 24.0f, 26.0f, 28.0f, 30.0f, + 32.0f, 36.0f, 40.0f, 44.0f, 48.0f, 52.0f, 56.0f, 60.0f, + 64.0f, 72.0f, 80.0f, 88.0f, 96.0f, 104.0f, 112.0f, 120.0f, + 128.0f, 144.0f, 160.0f, 176.0f, 192.0f, 208.0f, 224.0f, 240.0f, + 256.0f, 288.0f, 320.0f, 352.0f, 384.0f, 416.0f, 448.0f, +}; + +static uint8_t fp32_to_f8_e4m3_ref(float value) { + if (isnan(value)) { + return 0x7F; + } + const uint8_t sign = signbit(value) ? 0x80 : 0; + const double magnitude = fabs((double) value); + if (magnitude >= F8_E4M3_VALUES[126]) { + return sign | 126; + } + int nearest = 0; + double distance = magnitude; + for (int i = 1; i < 127; ++i) { + const double candidate = fabs(magnitude - F8_E4M3_VALUES[i]); + if (candidate < distance || (candidate == distance && (i & 1) == 0)) { + nearest = i; + distance = candidate; + } + } + return sign | (uint8_t) nearest; +} + +static int test_f8_e4m3(bool verbose) { + static_assert(sizeof(ggml_fp8_e4m3_t) == 1); + + const auto * traits = ggml_get_type_traits(GGML_TYPE_F8_E4M3); + const auto * traits_cpu = ggml_get_type_traits_cpu(GGML_TYPE_F8_E4M3); + std::vector encoded(256); + std::vector roundtrip(256); + std::vector decoded(256); + int num_failed = 0; + + const bool traits_failed = traits->blck_size != 1 || traits->type_size != sizeof(ggml_fp8_e4m3_t) || traits->is_quantized; + num_failed += traits_failed; + if (verbose || traits_failed) { + printf("f8_e4m3 scalar type traits: %s\n", RESULT_STR[traits_failed]); + } + + for (int i = 0; i < 256; ++i) { + encoded[i] = (uint8_t) i; + } + traits->to_float(encoded.data(), decoded.data(), decoded.size()); + traits_cpu->from_float(decoded.data(), roundtrip.data(), decoded.size()); + + for (int i = 0; i < 256; ++i) { + const bool is_nan = (i & 0x7F) == 0x7F; + const float expected = is_nan ? NAN : copysignf(F8_E4M3_VALUES[i & 0x7F], i & 0x80 ? -1.0f : 1.0f); + const bool decode_failed = is_nan ? !isnan(decoded[i]) + : decoded[i] != expected || (expected == 0.0f && signbit(decoded[i]) != signbit(expected)); + const bool encode_failed = !is_nan && roundtrip[i] != encoded[i]; + if (decode_failed || encode_failed) { + num_failed++; + if (verbose) { + printf("f8_e4m3 code 0x%02x failed: decoded=%f expected=%f roundtrip=0x%02x\n", + i, decoded[i], expected, roundtrip[i]); + } + } + } + + std::vector inputs; + for (float value : F8_E4M3_VALUES) { + inputs.push_back(value); + inputs.push_back(-value); + } + for (int i = 0; i < 126; ++i) { + const float midpoint = (F8_E4M3_VALUES[i] + F8_E4M3_VALUES[i + 1]) / 2.0f; + for (float value : {nextafterf(midpoint, -INFINITY), midpoint, nextafterf(midpoint, INFINITY)}) { + inputs.push_back(value); + inputs.push_back(-value); + } + } + for (float value : {std::numeric_limits::denorm_min(), std::numeric_limits::max(), + nextafterf(448.0f, INFINITY), nextafterf(464.0f, -INFINITY), 464.0f, + nextafterf(464.0f, INFINITY), INFINITY, NAN}) { + inputs.push_back(value); + inputs.push_back(-value); + } + + std::vector quantized(inputs.size()); + std::vector quantized_ref(inputs.size()); + traits_cpu->from_float(inputs.data(), quantized.data(), inputs.size()); + traits->from_float_ref(inputs.data(), quantized_ref.data(), inputs.size()); + for (size_t i = 0; i < inputs.size(); ++i) { + const uint8_t expected = fp32_to_f8_e4m3_ref(inputs[i]); + if (quantized[i] != expected || quantized_ref[i] != expected) { + num_failed++; + if (verbose) { + printf("f8_e4m3 input %a failed: encoded=0x%02x reference=0x%02x expected=0x%02x\n", + inputs[i], quantized[i], quantized_ref[i], expected); + } + } + } + + for (int n : {0, 1, 7, 8, 15, 16, 31, 32, 63, 64, 65, 129, 257}) { + std::vector x(n + 1); + std::vector y(n + 1); + double expected = 0.0; + for (int i = 0; i < n; ++i) { + const uint8_t code = (uint8_t) (i % 254 + (i % 254 >= 127)); + x[i + 1] = code; + y[i + 1] = (i % 7 - 3) * 0.125f; + expected += copysignf(F8_E4M3_VALUES[code & 0x7F], code & 0x80 ? -1.0f : 1.0f) * y[i + 1]; + } + float actual; + traits_cpu->vec_dot(n, &actual, 0, x.data() + 1, 0, y.data() + 1, 0, 1); + const bool failed = fabs(actual - expected) > 1e-5 * std::max(1.0, fabs(expected)); + num_failed += failed; + if (verbose || failed) { + printf("f8_e4m3 dot n=%d: %s (ref=%f got=%f)\n", n, RESULT_STR[failed], expected, actual); + } + } + + for (uint8_t code : {0x7F, 0xFF}) { + std::vector x(65, 0); + std::vector y(65, 1.0f); + x[0] = x.back() = code; + float actual; + traits_cpu->vec_dot(x.size(), &actual, 0, x.data(), 0, y.data(), 0, 1); + num_failed += !isnan(actual); + } + + if (verbose || num_failed) { + printf("f8_e4m3 exhaustive conversion: %s (%d failures)\n", RESULT_STR[num_failed != 0], num_failed); + } + return num_failed; +} + static int test_vec_dot_q(bool verbose) { int num_failed = 0; @@ -299,6 +446,7 @@ int main(int argc, char * argv[]) { int num_failed = 0; num_failed += test_vec_dot_f32(verbose); + num_failed += test_f8_e4m3(verbose); num_failed += test_vec_dot_q(verbose); num_failed += test_quantize_imatrix_degenerate(verbose); diff --git a/tests/test_conversion_fp8.py b/tests/test_conversion_fp8.py new file mode 100644 index 000000000..be9c3228c --- /dev/null +++ b/tests/test_conversion_fp8.py @@ -0,0 +1,370 @@ +"""Test FP8 conversion with small synthetic checkpoints; no GPU or model downloads required. + +Install dependencies and run from the repository root: + python -m pip install -r requirements/requirements-convert_hf_to_gguf.txt + python -m unittest discover -s tests -p 'test_conversion_fp8.py' -v + +Use unittest discovery so the repository root is available for imports. +""" + +from __future__ import annotations + +import json +from pathlib import Path +import tempfile +import unittest + +import numpy as np +import torch +from safetensors.torch import save_file + +from conversion.base import LazyTorchTensor, gguf +from conversion.deepseek import DeepseekV2Model +from conversion.llama import LlamaModel +from conversion.nemotron import NemotronHModel +from conversion.phi import PhiMoeModel +from conversion.qwen import Qwen2MoeModel, Qwen3NextModel + + +class ForcedF16Model(Qwen2MoeModel): + model_arch = gguf.MODEL_ARCH.QWEN2MOE + + def tensor_force_quant(self, name, new_name, bid, n_dims): + if name.endswith("q_proj.weight"): + return gguf.GGMLQuantizationType.F16 + return super().tensor_force_quant(name, new_name, bid, n_dims) + + +class OffsetNormModel(Qwen2MoeModel): + model_arch = gguf.MODEL_ARCH.QWEN2MOE + + def modify_tensors(self, data_torch, name, bid): + if name.endswith("input_layernorm.weight"): + data_torch = data_torch + 1.0 + yield from super().modify_tensors(data_torch, name, bid) + + +class TestFP8Conversion(unittest.TestCase): + def convert(self, model_class, tensors, eager=False, check_lazy=False, scan_nvfp4=False, ftype=gguf.LlamaFileType.GUESSED, hparams_override=None, **kwargs): + with tempfile.TemporaryDirectory() as directory: + return self._convert(Path(directory), model_class, tensors, eager, check_lazy, scan_nvfp4, ftype, hparams_override, **kwargs) + + def _convert(self, path, model_class, tensors, eager, check_lazy, scan_nvfp4, ftype, hparams_override, **kwargs): + hparams = { + "model_type": "qwen2_moe", + "architectures": ["Qwen2MoeForCausalLM"], + "num_hidden_layers": 1, + "hidden_size": 4, + "num_attention_heads": 1, + "num_key_value_heads": 1, + "num_experts": 2, + "qk_nope_head_dim": 2, + "v_head_dim": 2, + "quantization_config": {"quant_method": "modelopt"}, + } + hparams.update(hparams_override or {}) + (path / "config.json").write_text(json.dumps(hparams)) + save_file(tensors, str(path / "model.safetensors")) + model = model_class(path, ftype, path / "out.gguf", hparams=hparams, eager=eager, **kwargs) + reads: list[str] = [] + if check_lazy: + self.assertFalse(eager) + + def track(name, gen): + tensor = gen() + + def load(source): + reads.append(name) + return LazyTorchTensor.to_eager(source()) + + return LazyTorchTensor(meta=torch.empty(tuple(tensor.shape), dtype=tensor.dtype, device="meta"), args=(gen,), func=load) + + model.model_tensors = { + name: lambda name=name, gen=gen: track(name, gen) + for name, gen in model.model_tensors.items() + } + if scan_nvfp4: + model._generate_nvfp4_tensors() + self.assertEqual(reads, []) + model.prepare_tensors() + if check_lazy: + deferred = {name for name, tensor in tensors.items() if tensor.numel() != 1 or name.endswith((".weight", ".bias"))} + self.assertFalse(deferred.intersection(reads), f"Materialized during preparation: {reads}") + for name, info in model.gguf_writer.tensors[0].items(): + if name.endswith((".weight", ".bias")): + self.assertIsInstance(info.tensor, gguf.LazyNumpyTensor) + result = { + name: (info.dtype, gguf.LazyNumpyTensor.to_eager(info.tensor).copy()) + for name, info in model.gguf_writer.tensors[0].items() + } + if check_lazy: + self.assertTrue(deferred.intersection(reads)) + return result + + def test_lazy_fp8_export(self): + name = "model.layers.0.self_attn.q_proj.weight" + weight = torch.tensor([[1., -1.], [2., -2.]]).to(torch.float8_e4m3fn) + for channel_scale in (False, True): + for scan_nvfp4 in (False, True): + with self.subTest(channel_scale=channel_scale, scan_nvfp4=scan_nvfp4): + tensors = { + name: weight, + name + "_scale": torch.full((2, 1) if channel_scale else (1,), 0.25), + } + result = self.convert(Qwen2MoeModel, tensors, check_lazy=True, scan_nvfp4=scan_nvfp4) + dtype, data = result["blk.0.attn_q.weight"] + if channel_scale: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + np.testing.assert_array_equal(data, weight.float().numpy() * 0.25) + else: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F8_E4M3) + np.testing.assert_array_equal(gguf.dequantize(data, dtype), weight.float().numpy()) + np.testing.assert_array_equal(result["blk.0.attn_q.scale"][1], [0.25]) + + def test_mixed_compressed_fp8_block_scale(self): + name = "model.layers.0.self_attn.q_proj.weight" + scale = torch.tensor([[0.25, 0.5], [1.0, 2.0]]) + tensors = { + name: torch.ones((3, 32)).to(torch.float8_e4m3fn), + name + "_scale": scale, + } + config = { + "quant_method": "compressed-tensors", + "format": "mixed-precision", + "config_groups": { + "nvfp4": {"format": "nvfp4-pack-quantized"}, + "fp8": {"format": "float-quantized", "weights": { + "type": "float", "num_bits": 8, "strategy": "block", "block_structure": [2, 16], + }}, + }, + } + result = self.convert(Qwen2MoeModel, tensors, check_lazy=True, hparams_override={"quantization_config": config}) + self.assertEqual(set(result), {"blk.0.attn_q.weight"}) + dtype, data = result["blk.0.attn_q.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + expected = scale.repeat_interleave(2, 0).repeat_interleave(16, 1)[:3, :32] + np.testing.assert_array_equal(data, expected.numpy()) + + def test_mixed_compressed_fp8_block_layout_errors(self): + name = "model.layers.0.self_attn.q_proj.weight" + tensors = { + name: torch.ones((3, 32)).to(torch.float8_e4m3fn), + name + "_scale": torch.ones((2, 2)), + } + for configs, message in ( + ([], "requires a block layout"), + ([{"strategy": "block"}], "requires a block layout"), + ([{"strategy": "block", "block_structure": [2, 16], "group_size": 16}], "requires a block layout"), + ([{"strategy": "block", "block_structure": [2, 16]}, {"strategy": "block", "block_structure": [2, 32]}], "ambiguous block layouts"), + ): + with self.subTest(configs=configs): + groups: dict[str, dict[str, object]] = {"nvfp4": {"format": "nvfp4-pack-quantized"}} + for i, config in enumerate(configs): + groups[str(i)] = {"format": "float-quantized", "weights": {"type": "float", "num_bits": 8, **config}} + config = {"quant_method": "compressed-tensors", "format": "mixed-precision", "config_groups": groups} + with self.assertRaisesRegex(ValueError, message): + self.convert(Qwen2MoeModel, tensors, hparams_override={"quantization_config": config}) + + def test_mla_scale_fallback(self): + weight = torch.arange(16, dtype=torch.float32).reshape(4, 4).to(torch.float8_e4m3fn) + prefix = "model.layers.0.self_attn." + tensors = {} + for proj in ("q_a_proj", "q_b_proj", "kv_a_proj_with_mqa", "kv_b_proj"): + tensors[prefix + proj + ".weight"] = weight.clone() + tensors[prefix + proj + ".weight_scale"] = torch.tensor([0.25]) + for eager in (False, True): + with self.subTest(eager=eager): + result = self.convert(DeepseekV2Model, tensors, eager=eager) + self.assertFalse(any(name.endswith(".scale") for name in result)) + np.testing.assert_array_equal(result["blk.0.attn_k_b.weight"][1], (weight.float()[:2].T * 0.25).numpy()[None]) + np.testing.assert_array_equal(result["blk.0.attn_v_b.weight"][1], (weight.float()[2:] * 0.25).numpy()[None]) + + def test_output_type_override(self): + name = "model.layers.0.self_attn.q_proj.weight" + weight = torch.tensor([[1., -1.], [2., -2.]]).to(torch.float8_e4m3fn) + for eager in (False, True): + for scaled in (False, True): + with self.subTest(eager=eager, scaled=scaled): + tensors = {name: weight} + if scaled: + tensors[name + "_scale"] = torch.tensor([0.25]) + result = self.convert(ForcedF16Model, tensors, eager=eager, check_lazy=not eager) + dtype, data = result["blk.0.attn_q.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + self.assertNotIn("blk.0.attn_q.scale", result) + np.testing.assert_array_equal(data, weight.float().numpy() * (0.25 if scaled else 1.)) + + def test_fp8_bias_is_numeric_f32(self): + bias = torch.tensor([1., -1.]).to(torch.float8_e4m3fn) + result = self.convert(Qwen2MoeModel, {"model.layers.0.self_attn.q_proj.bias": bias}, check_lazy=True) + dtype, data = result["blk.0.attn_q.bias"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F32) + np.testing.assert_array_equal(data, [1., -1.]) + + def test_fp8_mandatory_f32(self): + for name, target, shape in ( + ("model.layers.0.linear_attn.conv1d.weight", "blk.0.ssm_conv1d.weight", (4, 1, 4)), + ("model.layers.0.mlp.gate.weight", "blk.0.ffn_gate_inp.weight", (2, 4)), + ): + for eager in (False, True): + for scaled in (False, True): + with self.subTest(name=name, eager=eager, scaled=scaled): + weight = torch.arange(int(np.prod(shape)), dtype=torch.float32).reshape(shape).to(torch.float8_e4m3fn) + tensors = {name: weight} + if scaled: + tensors[name + "_scale"] = torch.tensor([0.25]) + result = self.convert(Qwen3NextModel, tensors, eager=eager, check_lazy=not eager, hparams_override={"mtp_num_hidden_layers": 1}) + dtype, data = result[target] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F32) + np.testing.assert_array_equal(data, weight.float().squeeze().numpy() * (0.25 if scaled else 1.)) + self.assertNotIn(target.removesuffix("weight") + "scale", result) + + def test_fp8_norm_transform(self): + for shape in ((4,), (2, 2)): + with self.subTest(shape=shape): + weight = torch.ones(shape, dtype=torch.float32).to(torch.float8_e4m3fn) + result = self.convert(OffsetNormModel, {"model.layers.0.input_layernorm.weight": weight}, check_lazy=True) + dtype, data = result["blk.0.attn_norm.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F32) + np.testing.assert_array_equal(data, np.full(shape, 2.)) + + def test_explicit_outtypes(self): + name = "model.layers.0.self_attn.q_proj.weight" + weight = torch.ones((4, 256), dtype=torch.float32).to(torch.float8_e4m3fn) + for ftype, qtype in ( + (gguf.LlamaFileType.ALL_F32, gguf.GGMLQuantizationType.F32), + (gguf.LlamaFileType.MOSTLY_F16, gguf.GGMLQuantizationType.F16), + (gguf.LlamaFileType.MOSTLY_BF16, gguf.GGMLQuantizationType.BF16), + (gguf.LlamaFileType.MOSTLY_Q8_0, gguf.GGMLQuantizationType.Q8_0), + (gguf.LlamaFileType.MOSTLY_TQ1_0, gguf.GGMLQuantizationType.TQ1_0), + (gguf.LlamaFileType.MOSTLY_TQ2_0, gguf.GGMLQuantizationType.TQ2_0), + ): + for eager in (False, True): + for scaled in (False, True): + with self.subTest(ftype=ftype, eager=eager, scaled=scaled): + tensors = {name: weight} + if scaled: + tensors[name + "_scale"] = torch.tensor([0.25]) + tensors[name.removesuffix("weight") + "input_scale"] = torch.tensor([0.5]) + result = self.convert(Qwen2MoeModel, tensors, ftype=ftype, eager=eager, check_lazy=not eager) + self.assertEqual(set(result), {"blk.0.attn_q.weight"}) + dtype, data = result["blk.0.attn_q.weight"] + self.assertEqual(dtype, qtype) + np.testing.assert_allclose(gguf.dequantize(data, dtype), np.full((4, 256), 0.25 if scaled else 1.), atol=1e-4) + + def test_embedding_scale_fallback(self): + name = "model.embed_tokens.weight" + weight = torch.ones((4, 4), dtype=torch.float32).to(torch.float8_e4m3fn) + for eager in (False, True): + with self.subTest(eager=eager): + result = self.convert(Qwen2MoeModel, {name: weight, name + "_scale": torch.tensor([0.25])}, eager=eager, check_lazy=not eager) + self.assertEqual(set(result), {"token_embd.weight"}) + dtype, data = result["token_embd.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + np.testing.assert_array_equal(data, np.full((4, 4), 0.25)) + + def test_fp8_scale_formats(self): + name = "model.layers.0.self_attn.q_proj.weight" + weight = torch.ones((4, 4), dtype=torch.float32).to(torch.float8_e4m3fn) + for suffix, config in ( + ("_scale_inv", {"quant_method": "fp8"}), + ("_scale", {"quant_method": "compressed-tensors", "format": "float-quantized", "config_groups": {"group_0": {"weights": {"type": "float", "num_bits": 8, "strategy": "tensor"}}}}), + ): + for ftype in (gguf.LlamaFileType.GUESSED, gguf.LlamaFileType.MOSTLY_F16): + with self.subTest(suffix=suffix, ftype=ftype): + result = self.convert(Qwen2MoeModel, {name: weight, name + suffix: torch.tensor([0.25])}, ftype=ftype, hparams_override={"quantization_config": config}, check_lazy=True) + dtype, data = result["blk.0.attn_q.weight"] + if ftype == gguf.LlamaFileType.GUESSED: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F8_E4M3) + np.testing.assert_array_equal(result["blk.0.attn_q.scale"][1], [0.25]) + else: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + self.assertEqual(set(result), {"blk.0.attn_q.weight"}) + np.testing.assert_array_equal(data, np.full((4, 4), 0.25)) + + def test_expert_naming(self): + for model_class, prefix, projections, hparams in ( + (LlamaModel, "model.layers.0.block_sparse_moe", (("w1", "gate"), ("w3", "up"), ("w2", "down")), {"num_local_experts": 2}), + (PhiMoeModel, "model.layers.0.block_sparse_moe", (("w1", "gate"), ("w3", "up"), ("w2", "down")), {"max_position_embeddings": 16, "rope_scaling": {"original_max_position_embeddings": 16}}), + (NemotronHModel, "backbone.layers.0.mixer", (("up_proj", "up"), ("down_proj", "down")), {"num_experts_per_tok": 1, "n_routed_experts": 2, "head_dim": 4, "num_heads": 1, "n_groups": 1, "hybrid_override_pattern": "E"}), + ): + for eager in (False, True): + for mixed in (False, True): + with self.subTest(model=model_class.__name__, eager=eager, mixed=mixed): + tensors = {} + for expert in (1, 0): + for proj, _ in projections: + name = f"{prefix}.experts.{expert}.{proj}.weight" + dtype = torch.float16 if mixed and expert == 1 else torch.float8_e4m3fn + tensors[name] = torch.ones((4, 4), dtype=torch.float32).to(dtype) + tensors[name + "_scale"] = torch.tensor([0.25 * (expert + 1)]) + result = self.convert(model_class, tensors, hparams_override=hparams, eager=eager, check_lazy=not eager) + for _, proj in projections: + dtype, data = result[f"blk.0.ffn_{proj}_exps.weight"] + if mixed: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + np.testing.assert_array_equal(data, np.array([np.full((4, 4), 0.25), np.full((4, 4), 0.5)])) + self.assertNotIn(f"blk.0.ffn_{proj}_exps.scale", result) + else: + self.assertEqual(dtype, gguf.GGMLQuantizationType.F8_E4M3) + np.testing.assert_array_equal(gguf.dequantize(data, dtype), np.ones((2, 4, 4))) + np.testing.assert_array_equal(result[f"blk.0.ffn_{proj}_exps.scale"][1], [0.25, 0.5]) + + def test_mixed_experts(self): + for eager in (False, True): + for other in ("f16", "channel", "unscaled"): + with self.subTest(eager=eager, other=other): + tensors = {} + for expert in range(2): + for proj in ("gate", "up", "down"): + name = f"model.layers.0.mlp.experts.{expert}.{proj}_proj.weight" + dtype = torch.float16 if expert == 1 and other == "f16" else torch.float8_e4m3fn + tensors[name] = torch.ones((4, 4), dtype=torch.float32).to(dtype) + if expert == 0: + tensors[name + "_scale"] = torch.tensor([0.25]) + elif other == "channel": + tensors[name + "_scale"] = torch.full((4, 1), 0.5) + result = self.convert(Qwen2MoeModel, tensors, eager=eager, check_lazy=not eager) + self.assertEqual(len(result), 3) + for proj in ("gate", "up", "down"): + dtype, data = result[f"blk.0.ffn_{proj}_exps.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F16) + np.testing.assert_array_equal(data[0], np.full((4, 4), 0.25)) + np.testing.assert_array_equal(data[1], np.full((4, 4), 0.5 if other == "channel" else 1.)) + + def test_uniform_experts_stay_fp8(self): + tensors = {} + for expert in range(2): + for proj in ("gate", "up", "down"): + name = f"model.layers.0.mlp.experts.{expert}.{proj}_proj.weight" + tensors[name] = torch.ones((4, 4), dtype=torch.float32).to(torch.float8_e4m3fn) + tensors[name + "_scale"] = torch.tensor([0.25 * (expert + 1)]) + result = self.convert(Qwen2MoeModel, tensors, fuse_gate_up_exps=True, check_lazy=True) + for proj in ("gate", "up", "down"): + dtype, data = result[f"blk.0.ffn_{proj}_exps.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.F8_E4M3) + np.testing.assert_array_equal(gguf.dequantize(data, dtype), np.ones((2, 4, 4))) + np.testing.assert_array_equal(result[f"blk.0.ffn_{proj}_exps.scale"][1], [0.25, 0.5]) + + unscaled = {name: tensor for name, tensor in tensors.items() if name.endswith(".weight")} + result = self.convert(Qwen2MoeModel, unscaled, check_lazy=True) + self.assertEqual(len(result), 3) + for dtype, data in result.values(): + self.assertEqual(dtype, gguf.GGMLQuantizationType.F8_E4M3) + np.testing.assert_array_equal(gguf.dequantize(data, dtype), np.ones((2, 4, 4))) + + def test_fp8_as_q8(self): + name = "model.layers.0.self_attn.q_proj.weight" + tensors = { + name: torch.ones((4, 32), dtype=torch.float32).to(torch.float8_e4m3fn), + name + "_scale": torch.tensor([0.25]), + } + result = self.convert(Qwen2MoeModel, tensors, fp8_as_q8=True) + self.assertNotIn("blk.0.attn_q.scale", result) + dtype, data = result["blk.0.attn_q.weight"] + self.assertEqual(dtype, gguf.GGMLQuantizationType.Q8_0) + np.testing.assert_allclose(gguf.dequantize(data, dtype), np.full((4, 32), 0.25), atol=1e-4) + + +if __name__ == "__main__": + unittest.main()