From b012d20c71054189e948e2122b350a7c5a71382e Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Mon, 19 Oct 2026 17:44:36 +0300 Subject: [PATCH 01/15] Default the Granite 4.2 probe test to the Hub instead of skipping test_granite_4_2_chatml was gated on GRANITE_4_2_BASE pointing at an existing local path, on the stated grounds that "the 4.2 base ships on a local path, not the Hub". That is no longer true: ibm-granite/granite-4.2-3b is public and ungated, and this test needs only its tokenizer -- no weights and no GPU -- so there was nothing left for the gate to protect and it skipped on every machine that had not set the variable by hand. It now defaults to the Hub id with GRANITE_4_2_BASE still honoured as an override for a local or pre-release 4.2, and _tok keeps skipping on a fetch failure so an offline run degrades to a skip rather than a red test. While the assertions were being touched, add the one the docstring always promised and never checked: that the probe returns id 100256, not merely something that stringifies to <|im_start|>. The file goes to 9 passed / 0 skipped on CPU. Found while auditing every skip the GPU suite reports; that inventory is kept as a local working note rather than committed, since it is a point-in-time snapshot that goes stale as soon as a gate changes. Signed-off-by: itzikvaisman --- .gitignore | 3 +++ tests/composer/test_lora_substitute_probe.py | 18 ++++++++++++------ 2 files changed, 15 insertions(+), 6 deletions(-) diff --git a/.gitignore b/.gitignore index 819657dd..98944bd9 100644 --- a/.gitignore +++ b/.gitignore @@ -64,3 +64,6 @@ docs/KV_CACHE_OVERHEAD_REMOVAL.md docs/KV_CACHE_OVERHEAD_REMOVAL.html docs/KV_CACHE_OVERHEAD_REMOVAL*.html docs/UNIFIED_SWITCH_FUSED_CANONICAL.md +# Skip audit: a point-in-time snapshot of which tests were skipped and why. It +# goes stale the moment a gate changes, so it stays a local working note. +docs/SKIPPED_TESTS.html diff --git a/tests/composer/test_lora_substitute_probe.py b/tests/composer/test_lora_substitute_probe.py index 33aae30a..741fd2ea 100644 --- a/tests/composer/test_lora_substitute_probe.py +++ b/tests/composer/test_lora_substitute_probe.py @@ -53,14 +53,20 @@ def test_granite_4_0_micro(self): def test_granite_4_2_chatml(self): """Granite 4.2 (ChatML): the probe returns <|im_start|> (id 100256). - The 4.2 base ships on a local path (env GRANITE_4_2_BASE), not the Hub; - skips when unavailable so the suite still runs locally. + The ChatML counterpart of the two tests above: 4.0/4.1 use the + <|start_of_role|> role-marker template and 4.2 uses ChatML, so the probe + has to derive the substitute token from the template instead of + hardcoding either. ``GRANITE_4_2_BASE`` still overrides the source so a + local/pre-release 4.2 can be pointed at, but it is no longer required: + ibm-granite/granite-4.2-3b is public and ungated, and only its tokenizer + is needed (no weights, no GPU). ``_tok`` still skips on a fetch failure, + so an offline run degrades to a skip rather than a failure. """ - path = os.environ.get("GRANITE_4_2_BASE") - if not path or not os.path.exists(path): - pytest.skip("GRANITE_4_2_BASE not set / path missing") - tok = self._tok(path) + tok = self._tok( + os.environ.get("GRANITE_4_2_BASE") or "ibm-granite/granite-4.2-3b" + ) sub_id = _probe_lora_substitute_token_id(tok) + assert sub_id == 100256 assert tok.convert_ids_to_tokens([sub_id])[0] == "<|im_start|>" From 9de8bb90b7fb0dc42812bb2c3a440011a26a4c13 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Wed, 21 Oct 2026 11:03:17 +0300 Subject: [PATCH 02/15] Replace three opt-in skip flags with gates that tell the truth The skipped-test audit found that GRANITE_SWITCH_E2E_MODELS was doing two unrelated jobs, and doing one of them wrong. Where it gates cost, keep it -- composing a real ~3B checkpoint is genuinely expensive and should stay opt-in. But the knowledge of which variable to export lived only in shell history, so add `make test-e2e`. It runs tests/hf, tests/vllm and tests/integration, not just integration: the eleven gated modules split 4/5/1 across those directories, so an integration-only target would leave six still skipping and quietly under-deliver. Separate pytest invocations per directory so one leaking CUDA context cannot poison the next. Where the same flag gated tests/hf/test_quantization.py it was wrong -- what constrains those sixteen tests is capability, not cost. The flag skipped them before importorskip ever ran, so the quantization path had zero coverage on every CI run. Split it into gates that each report honestly: CUDA at module level, bitsandbytes/optimum-quanto via importorskip in the owning fixture, and a real probe for Quanto float8. Lifting the flag alone would only have swapped one silent skip for another, since neither backend was installed, so add a `quant` dependency group to dev, dev-vllm27 and dev-vllm30. Pinned, unlike most of this file: these JIT-compile CUDA extensions, so a floating version could turn the suite red from a background upgrade rather than from a commit. The four float8 tests needed a probe rather than an import check because they ERROR instead of skipping. The visible exception is a transformers "issues during automatic conversion of the weights" naming base_layer.weight, which reads like a granite-switch LoRA fault and is not one: ptxas aborts with "Feature 'cp.async' requires .target sm_80 or higher" while building marlin_cuda.o -- an INT4 kernel float8 never uses, pulled in because quanto_cuda builds all-or-nothing. Reproduced on an A100 with a bare nn.Linear and no granite-switch code, which is also what the probe runs, in a subprocess so a hard abort cannot take the session down. A skip now carries the real ptxas error instead of a misattributed LoRA one. Signed-off-by: itzikvaisman --- Makefile | 42 ++++++++++++++- pyproject.toml | 24 +++++++-- tests/hf/test_quantization.py | 96 ++++++++++++++++++++++++++++++++--- uv.lock | 53 ++++++++++++++----- 4 files changed, 191 insertions(+), 24 deletions(-) diff --git a/Makefile b/Makefile index 1dbbb328..b309f617 100644 --- a/Makefile +++ b/Makefile @@ -1,4 +1,4 @@ -.PHONY: test test-unit test-composer test-hf test-vllm test-integration test-all test-cpu test-gpu test-gpu-full test-tp test-regression-fast test-regression-real test-regression-hf test-peft-equiv test-peft-equiv-hf lint fmt pre-commit help +.PHONY: test test-unit test-composer test-hf test-vllm test-integration test-all test-cpu test-gpu test-gpu-full test-tp test-e2e test-regression-fast test-regression-real test-regression-hf test-peft-equiv test-peft-equiv-hf lint fmt pre-commit help # Default Python (can override: make test PYTHON=python3.11) PYTHON ?= python @@ -6,6 +6,26 @@ PYTHON ?= python # Pytest flags per CLAUDE.md guidelines PYTEST_FLAGS = -v -s --tb=short +# ── The real-checkpoint (e2e) gate ─────────────────────────────────────────── +# Eleven modules across tests/hf, tests/integration and tests/vllm compose a real +# ~3B Granite plus adapters, and skip themselves unless GRANITE_SWITCH_E2E_MODELS=1. +# That flag is an opt-in for COST, not a capability check: nothing is wrong with the +# machine when they skip, they are just expensive (a multi-GB download on a cold +# cache). `make test-e2e` is the supported way to turn them on, so the knowledge of +# which variable that is lives here instead of in everyone's shell history. +# +# GRANITE_SWITCH_E2E_DIR is where composed checkpoints are cached between runs; the +# modules warm-reuse any dir that already has a config.json, so a second `make +# test-e2e` pays the compose cost once. Override it to keep the cache off /tmp: +# make test-e2e GRANITE_SWITCH_E2E_DIR=/scratch/gs-e2e +# +# `?=` on both, so an already-exported value wins and CI can set them in the +# environment instead. +GRANITE_SWITCH_E2E_MODELS ?= 1 +GRANITE_SWITCH_E2E_DIR ?= /tmp/granite_switch_e2e +E2E_ENV = GRANITE_SWITCH_E2E_MODELS=$(GRANITE_SWITCH_E2E_MODELS) \ + GRANITE_SWITCH_E2E_DIR=$(GRANITE_SWITCH_E2E_DIR) + # Individual test suites test-unit: $(PYTHON) -m pytest tests/unit/ $(PYTEST_FLAGS) @@ -34,6 +54,22 @@ test-gpu-full: test-gpu test-regression-fast # GPU tests + regression suite test-tp: $(PYTHON) -m pytest tests/vllm/test_tp_integration.py tests/vllm/test_tp_lora.py $(PYTEST_FLAGS) +# Everything the integration suite covers, WITH the real-checkpoint modules enabled. +# +# Three directories, not just tests/integration/: the e2e-gated modules are spread +# across tests/hf (4), tests/integration (5) and tests/vllm (1), so running only +# tests/integration/ would leave six of them still skipping and the target would +# quietly under-deliver. +# +# Separate pytest invocations on purpose. These modules each build a vLLM engine or +# serve a model, and a per-directory process keeps one leaking CUDA context from +# poisoning the next. No leading `-`: unlike the regression targets below, a failure +# here should stop the run — it means a real checkpoint behaved wrongly. +test-e2e: + $(E2E_ENV) $(PYTHON) -m pytest tests/hf/ $(PYTEST_FLAGS) + $(E2E_ENV) $(PYTHON) -m pytest tests/vllm/ $(PYTEST_FLAGS) + $(E2E_ENV) $(PYTHON) -m pytest tests/integration/ $(PYTEST_FLAGS) + test-all: test-cpu test-gpu # Everything # Regression test suites (no -x: run all tests even if some fail) @@ -82,6 +118,10 @@ help: @echo " test-gpu - Run all GPU tests" @echo " test-gpu-full - Run GPU tests + regression-fast" @echo " test-tp - Run TP integration + TP LoRA tests (requires 2+ GPUs)" + @echo " test-e2e - Run hf + vllm + integration WITH the real-checkpoint" + @echo " (GRANITE_SWITCH_E2E_MODELS) modules enabled. Composes a" + @echo " real ~3B checkpoint; cached in GRANITE_SWITCH_E2E_DIR" + @echo " (default /tmp/granite_switch_e2e). Requires GPU." @echo " test-all - Run all tests" @echo " test-regression-fast - Run fast synthetic regression tests" @echo " test-regression-real - Run real model regression tests (requires_model)" diff --git a/pyproject.toml b/pyproject.toml index ee201ae3..8de14871 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -82,15 +82,31 @@ vllm27 = ["vllm>=0.27.0,<0.28.0"] # `audio` is included so the audio tests can actually run: the ASR path needs # vLLM's audio deps (av/soundfile/resampy) at runtime, and no group pulled them # in before (integration tests failed with ModuleNotFoundError on a synced pod). -dev = ["pytest", "pytest-cov", { include-group = "vllm26" }, "granite-switch[hf,compose,audio]"] -dev-vllm27 = ["pytest", "pytest-cov", { include-group = "vllm27" }, "granite-switch[hf,compose,audio]"] -dev-vllm30 = ["pytest", "pytest-cov", { include-group = "vllm30" }, "granite-switch[hf,compose,audio]"] +# +# `quant` carries the two quantization backends tests/hf/test_quantization.py drives. +# They are in every dev group because that module no longer opts in through an env +# flag -- it runs wherever it can, and without these it would skip on importorskip +# instead, which is the same silent nothing the flag used to produce. +# +# PINNED, unlike most of this file. These are JIT-compiling CUDA extensions, so a +# new release can change whether a kernel builds at all on a given image -- the +# Quanto float8 `quanto_cuda`/ptxas break is exactly that failure mode, and it is +# inside one of these two. Floating them would mean a red suite traceable to a +# background upgrade rather than to a commit. Both versions below are measured green +# on the Vela A100 image. Raise them deliberately, with a GPU run. +quant = ["bitsandbytes==0.50.2", "optimum-quanto==0.2.7"] +dev = ["pytest", "pytest-cov", { include-group = "vllm26" }, { include-group = "quant" }, "granite-switch[hf,compose,audio]"] +dev-vllm27 = ["pytest", "pytest-cov", { include-group = "vllm27" }, { include-group = "quant" }, "granite-switch[hf,compose,audio]"] +dev-vllm30 = ["pytest", "pytest-cov", { include-group = "vllm30" }, { include-group = "quant" }, "granite-switch[hf,compose,audio]"] # The newest vLLM line, as an alias. The GPU workflow (derive_legs.py) and # CLAUDE.md ask for `dev-latest` rather than the concrete `dev-vllmNN`, so moving # the ceiling is repointing this one line instead of chasing the version through # every caller. Currently the 0.30 line. dev-latest = [{ include-group = "dev-vllm30" }] -test = ["pytest", "pytest-cov", "bitsandbytes", "optimum-quanto", { include-group = "dev" }] +# The two quantization backends come in through `dev` -> `quant` now, PINNED. They +# used to be listed here bare (unversioned), which would have silently overridden +# the pin with "whatever is newest". +test = ["pytest", "pytest-cov", { include-group = "dev" }] [tool.uv] default-groups = ["vllm26"] diff --git a/tests/hf/test_quantization.py b/tests/hf/test_quantization.py index e4a21c98..f4cccc84 100644 --- a/tests/hf/test_quantization.py +++ b/tests/hf/test_quantization.py @@ -10,8 +10,18 @@ - BitsAndBytes: NF4 (INT4), FP4 - Quanto: INT4, FP8 -No hardware gating — all methods dequantize to BF16 at compute time on HF backend. -Requires: CUDA GPU, bitsandbytes, optimum-quanto. +GATING. This module runs wherever it *can* run — there is no opt-in env flag. It +used to sit behind ``GRANITE_SWITCH_E2E_MODELS=1`` like the e2e suites, which meant +all 16 tests skipped on every CI run and the quantization path was never exercised +at all. What actually constrains them is capability, not cost policy, so each +requirement is now checked where it applies and reported honestly: + +- CUDA: module-level skip (every backend here needs a GPU). +- ``bitsandbytes`` / ``optimum-quanto``: ``importorskip`` in the owning fixture, so + a missing backend skips only its own tests. +- Quanto **float8**: a real probe (``_probe_quanto_float8``) — see + ``require_quanto_float8`` for why a capability check is needed and what it + catches. These four used to ERROR in fixture setup on an A100. Model: a real MultiSwitch checkpoint composed on demand from granitelib-rag, which carries both an aLoRA intrinsic (``answerability``) and a LoRA one @@ -33,8 +43,8 @@ pytest.mark.requires_model, pytest.mark.gpu, pytest.mark.skipif( - os.environ.get("GRANITE_SWITCH_E2E_MODELS") != "1", - reason="composes a real ~3B mixed checkpoint; set GRANITE_SWITCH_E2E_MODELS=1", + not torch.cuda.is_available(), + reason="every quantization backend here needs a CUDA GPU", ), ] @@ -51,7 +61,13 @@ @pytest.fixture(scope="module") def model_path(): - """Compose (or warm-reuse) a rag MultiSwitch checkpoint; return its dir.""" + """Compose (or warm-reuse) a rag MultiSwitch checkpoint; return its dir. + + The compose is the one genuinely expensive thing in this module (a ~3B base + model plus adapters). It is NOT behind an opt-in flag: `make test-e2e` is the + way this module is meant to be run, and the warm-reuse below makes a second run + free. ``GRANITE_SWITCH_E2E_DIR`` relocates the cache. + """ out_dir = _E2E_ROOT / "quant-rag" if (out_dir / "config.json").exists(): print(f"warm-reuse {out_dir}", file=sys.stderr) @@ -326,8 +342,76 @@ def test_base_layers_are_quantized(self, quanto_model): # =========================================================================== +# The probe body. Deliberately contains NO granite-switch code: it quantizes a bare +# nn.Linear, so a failure can only come from optimum-quanto/CUDA, never from us. +# ``freeze()`` is the operative call — that is what materializes the quantized +# weight and what JIT-builds the ``quanto_cuda`` extension. +_FLOAT8_PROBE = """ +import torch +from optimum.quanto import freeze, qfloat8, quantize + +m = torch.nn.Linear(64, 64).cuda() +quantize(m, weights=qfloat8) +freeze(m) +m(torch.randn(2, 64, device="cuda")) +print("FLOAT8_OK") +""" + + +def _probe_quanto_float8(): + """Ask a subprocess whether Quanto can actually quantize to float8 here. + + A subprocess for two reasons: the extension build and a CUDA context must not + be inherited into the test session, and a hard abort (``ptxas`` has been seen + to abort rather than raise) would otherwise take the whole run down. + + Returns ``(ok, reason)``. + """ + res = subprocess.run( + [sys.executable, "-c", _FLOAT8_PROBE], + capture_output=True, + text=True, + timeout=900, + ) + if "FLOAT8_OK" in res.stdout: + return True, "" + tail = (res.stderr or res.stdout or "").strip().splitlines() + return False, " | ".join(tail[-3:])[:400] or f"exit {res.returncode}" + + +@pytest.fixture(scope="module") +def require_quanto_float8(): + """Skip the FP8 tests unless Quanto float8 works on this machine. + + These four tests used to ERROR rather than skip. The visible exception was a + transformers ``RuntimeError: issues during automatic conversion of the weights`` + naming ``base_layer.weight``, which reads like a granite-switch LoRA fault and is + not one: the real failure is inside ``quanto.convert → freeze() → + quantize_weight``, where optimum-quanto JIT-builds its ``quanto_cuda`` extension + and ``ptxas`` aborts with ``Feature 'cp.async' requires .target sm_80 or higher`` + while compiling ``compute_75.ptx``. Reproduced on an A100 (sm_80) with a bare + ``nn.Linear`` and no granite-switch code at all. + + Note the failing object is ``marlin_cuda.o`` — an INT4 kernel — because + ``quanto_cuda`` bundles marlin+awq+unpack and builds all-or-nothing, so float8's + ``freeze()`` pays for compiling a kernel it never uses. That is why INT4 can be + green while float8 is not. + + The probe does not care which of those is the root cause: it asks the only + question that matters to these tests — can this machine freeze a float8 linear — + and skips with the real error when it cannot. A green result here therefore means + "float8 verified", and a skip means "float8 unavailable", never a misattributed + LoRA failure. + """ + pytest.importorskip("optimum.quanto") + ok, reason = _probe_quanto_float8() + print(f"\nQuanto float8 skip-guard: ok={ok} reason={reason!r}") + if not ok: + pytest.skip(f"Quanto cannot quantize to float8 on this machine: {reason}") + + @pytest.fixture(scope="module") -def fp8_model(model_path): +def fp8_model(model_path, require_quanto_float8): """Load granite-switch from HF with Quanto FP8 quantization.""" pytest.importorskip("optimum.quanto") from transformers import AutoModelForCausalLM, AutoTokenizer, QuantoConfig diff --git a/uv.lock b/uv.lock index 8e5f16b9..84e0906a 100644 --- a/uv.lock +++ b/uv.lock @@ -435,19 +435,21 @@ wheels = [ [[package]] name = "bitsandbytes" -version = "0.49.2" +version = "0.50.2" source = { registry = "https://pypi.org/simple" } dependencies = [ - { name = "numpy", version = "2.3.5", source = { registry = "https://pypi.org/simple" }, marker = "(python_full_version < '3.13' and extra == 'group-14-granite-switch-test') or (python_full_version >= '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (python_full_version >= '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, - { name = "numpy", version = "2.4.6", source = { registry = "https://pypi.org/simple" }, marker = "(python_full_version >= '3.13' and extra == 'group-14-granite-switch-test') or (python_full_version < '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (python_full_version < '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "numpy", version = "2.3.5", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version < '3.13' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "numpy", version = "2.4.6", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version >= '3.13' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, { name = "packaging" }, - { name = "torch", version = "2.11.0", source = { registry = "https://pypi.org/simple" } }, + { name = "torch", version = "2.11.0", source = { registry = "https://pypi.org/simple" }, marker = "extra == 'group-14-granite-switch-dev' or extra == 'group-14-granite-switch-test' or extra == 'group-14-granite-switch-vllm26' or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "torch", version = "2.13.0", source = { registry = "https://pypi.org/simple" }, marker = "extra == 'group-14-granite-switch-dev-latest' or extra == 'group-14-granite-switch-dev-vllm27' or extra == 'group-14-granite-switch-dev-vllm30' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra != 'group-14-granite-switch-vllm26')" }, ] wheels = [ - { url = "https://files.pythonhosted.org/packages/d8/7d/f1fe0992334b18cd8494f89aeec1dcc674635584fcd9f115784fea3a1d05/bitsandbytes-0.49.2-py3-none-macosx_14_0_arm64.whl", hash = "sha256:87be5975edeac5396d699ecbc39dfc47cf2c026daaf2d5852a94368611a6823f", size = 131940, upload-time = "2026-02-16T21:26:04.572Z" }, - { url = "https://files.pythonhosted.org/packages/29/71/acff7af06c818664aa87ff73e17a52c7788ad746b72aea09d3cb8e424348/bitsandbytes-0.49.2-py3-none-manylinux_2_24_aarch64.whl", hash = "sha256:2fc0830c5f7169be36e60e11f2be067c8f812dfcb829801a8703735842450750", size = 31442815, upload-time = "2026-02-16T21:26:06.783Z" }, - { url = "https://files.pythonhosted.org/packages/19/57/3443d6f183436fbdaf5000aac332c4d5ddb056665d459244a5608e98ae92/bitsandbytes-0.49.2-py3-none-manylinux_2_24_x86_64.whl", hash = "sha256:54b771f06e1a3c73af5c7f16ccf0fc23a846052813d4b008d10cb6e017dd1c8c", size = 60651714, upload-time = "2026-02-16T21:26:11.579Z" }, - { url = "https://files.pythonhosted.org/packages/b6/d4/501655842ad6771fb077f576d78cbedb5445d15b1c3c91343ed58ca46f0e/bitsandbytes-0.49.2-py3-none-win_amd64.whl", hash = "sha256:2e0ddd09cd778155388023cbe81f00afbb7c000c214caef3ce83386e7144df7d", size = 55372289, upload-time = "2026-02-16T21:26:16.267Z" }, + { url = "https://files.pythonhosted.org/packages/55/bf/5290208ce1ecf0f2e6a916fc72a75f6e68021ecfd69e7014fc95998532eb/bitsandbytes-0.50.2-py3-none-macosx_14_0_arm64.whl", hash = "sha256:4311f52a880b341bada639e4edd1a3c8d786830c9c93cdde29eaa1f062c8f8e5", size = 123541, upload-time = "2026-08-27T00:10:48.726Z" }, + { url = "https://files.pythonhosted.org/packages/88/d5/b2cb5b5a9daf7349a02b1af2c49b6a044fda2702c9cc5dc296f648358327/bitsandbytes-0.50.2-py3-none-manylinux_2_24_aarch64.whl", hash = "sha256:d5772560dd94c4d9c57f50c9b017450a1707f7687bfd4b3dc86f7342aafe721e", size = 23777467, upload-time = "2026-08-27T00:10:50.92Z" }, + { url = "https://files.pythonhosted.org/packages/a5/6e/e4e8b75716dbe5e50964f070266e06f4e6806ce051bfb97f52ee162b9310/bitsandbytes-0.50.2-py3-none-manylinux_2_24_x86_64.whl", hash = "sha256:55348a9a4a21bfd99cf8c7b32fe67b4030ae5c2a05738e03c1747f65fa6ec283", size = 43139553, upload-time = "2026-08-27T00:10:54.751Z" }, + { url = "https://files.pythonhosted.org/packages/72/82/742dc27a1feab90c8f87f2ed14e6d72d05f9e1cf764b4d2ba30aa9b4a2cb/bitsandbytes-0.50.2-py3-none-win_amd64.whl", hash = "sha256:c697963c8fda3dcd0d7ebd9b5211ae4067feef7cd06e0350d4e816a434fe683d", size = 39096375, upload-time = "2026-08-27T00:10:58.297Z" }, + { url = "https://files.pythonhosted.org/packages/a2/57/61636c5b11b0a32e505127a6dce6fa8fcbf73978babe8fa37082ab547f1c/bitsandbytes-0.50.2-py3-none-win_arm64.whl", hash = "sha256:8437ab68a04ea56daf1d6ecb54230fb1d88be4b89fe2d79bc399bc0203b487cf", size = 1058684, upload-time = "2026-08-27T00:11:00.664Z" }, ] [[package]] @@ -1631,29 +1633,41 @@ vllm = [ [package.dev-dependencies] dev = [ + { name = "bitsandbytes" }, { name = "granite-switch", extra = ["audio", "compose", "hf"], marker = "extra == 'group-14-granite-switch-dev' or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "optimum-quanto" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", version = "0.26.0", source = { registry = "https://pypi.org/simple" } }, ] dev-latest = [ + { name = "bitsandbytes" }, { name = "granite-switch", extra = ["audio", "compose", "hf"], marker = "extra == 'group-14-granite-switch-dev-latest' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "optimum-quanto" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", version = "0.30.0", source = { registry = "https://pypi.org/simple" } }, ] dev-vllm27 = [ + { name = "bitsandbytes" }, { name = "granite-switch", extra = ["audio", "compose", "hf"], marker = "extra == 'group-14-granite-switch-dev-vllm27' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "optimum-quanto" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", version = "0.27.1", source = { registry = "https://pypi.org/simple" } }, ] dev-vllm30 = [ + { name = "bitsandbytes" }, { name = "granite-switch", extra = ["audio", "compose", "hf"], marker = "extra == 'group-14-granite-switch-dev-vllm30' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra != 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra != 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra != 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27')" }, + { name = "optimum-quanto" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", version = "0.30.0", source = { registry = "https://pypi.org/simple" } }, ] +quant = [ + { name = "bitsandbytes" }, + { name = "optimum-quanto" }, +] test = [ { name = "bitsandbytes" }, { name = "granite-switch", extra = ["audio", "compose", "hf"], marker = "extra == 'group-14-granite-switch-test' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30')" }, @@ -1702,33 +1716,45 @@ provides-extras = ["hf", "vllm", "compose", "build", "audio", "tutorials"] [package.metadata.requires-dev] dev = [ + { name = "bitsandbytes", specifier = "==0.50.2" }, { name = "granite-switch", extras = ["hf", "compose", "audio"] }, + { name = "optimum-quanto", specifier = "==0.2.7" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", specifier = ">=0.26.0,<0.27.0" }, ] dev-latest = [ + { name = "bitsandbytes", specifier = "==0.50.2" }, { name = "granite-switch", extras = ["hf", "compose", "audio"] }, + { name = "optimum-quanto", specifier = "==0.2.7" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", specifier = ">=0.30.0,<0.31.0" }, ] dev-vllm27 = [ + { name = "bitsandbytes", specifier = "==0.50.2" }, { name = "granite-switch", extras = ["hf", "compose", "audio"] }, + { name = "optimum-quanto", specifier = "==0.2.7" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", specifier = ">=0.27.0,<0.28.0" }, ] dev-vllm30 = [ + { name = "bitsandbytes", specifier = "==0.50.2" }, { name = "granite-switch", extras = ["hf", "compose", "audio"] }, + { name = "optimum-quanto", specifier = "==0.2.7" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", specifier = ">=0.30.0,<0.31.0" }, ] +quant = [ + { name = "bitsandbytes", specifier = "==0.50.2" }, + { name = "optimum-quanto", specifier = "==0.2.7" }, +] test = [ - { name = "bitsandbytes" }, + { name = "bitsandbytes", specifier = "==0.50.2" }, { name = "granite-switch", extras = ["hf", "compose", "audio"] }, - { name = "optimum-quanto" }, + { name = "optimum-quanto", specifier = "==0.2.7" }, { name = "pytest" }, { name = "pytest-cov" }, { name = "vllm", specifier = ">=0.26.0,<0.27.0" }, @@ -4012,10 +4038,11 @@ source = { registry = "https://pypi.org/simple" } dependencies = [ { name = "huggingface-hub" }, { name = "ninja" }, - { name = "numpy", version = "2.3.5", source = { registry = "https://pypi.org/simple" }, marker = "(python_full_version < '3.13' and extra == 'group-14-granite-switch-test') or (python_full_version >= '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (python_full_version >= '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, - { name = "numpy", version = "2.4.6", source = { registry = "https://pypi.org/simple" }, marker = "(python_full_version >= '3.13' and extra == 'group-14-granite-switch-test') or (python_full_version < '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (python_full_version < '3.13' and extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "numpy", version = "2.3.5", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version < '3.13' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "numpy", version = "2.4.6", source = { registry = "https://pypi.org/simple" }, marker = "python_full_version >= '3.13' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-latest') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-test') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm26') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, { name = "safetensors" }, - { name = "torch", version = "2.11.0", source = { registry = "https://pypi.org/simple" } }, + { name = "torch", version = "2.11.0", source = { registry = "https://pypi.org/simple" }, marker = "extra == 'group-14-granite-switch-dev' or extra == 'group-14-granite-switch-test' or extra == 'group-14-granite-switch-vllm26' or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-dev-vllm27') or (extra == 'group-14-granite-switch-dev-latest' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-dev-vllm30') or (extra == 'group-14-granite-switch-dev-vllm27' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-dev-vllm30' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm27' and extra == 'group-14-granite-switch-vllm30')" }, + { name = "torch", version = "2.13.0", source = { registry = "https://pypi.org/simple" }, marker = "extra == 'group-14-granite-switch-dev-latest' or extra == 'group-14-granite-switch-dev-vllm27' or extra == 'group-14-granite-switch-dev-vllm30' or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-dev' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-test' and extra == 'group-14-granite-switch-vllm30') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm27') or (extra == 'group-14-granite-switch-vllm26' and extra == 'group-14-granite-switch-vllm30') or (extra != 'group-14-granite-switch-dev' and extra != 'group-14-granite-switch-test' and extra != 'group-14-granite-switch-vllm26')" }, ] sdist = { url = "https://files.pythonhosted.org/packages/3e/df/03ae85090b33d81f06b4dd7a43b16cef6ee5f1d36d8fdcce864964895c70/optimum_quanto-0.2.7.tar.gz", hash = "sha256:91b5c2dc8a9100297dc7924a93747fb77ab010784b5e1f6d0208976ba054dade", size = 361601, upload-time = "2025-03-06T08:07:51.578Z" } wheels = [ From d9516df7ae228e2a48e2090b5703ae77e2df6472 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Wed, 21 Oct 2026 12:25:07 +0300 Subject: [PATCH 03/15] Stop generating the two BLOCK_M=1 bitmask cases that cannot exist test_two_adapters_split and test_partial_tile_padding each carried a pytest.skip for block_m=1, which reported two tests as skipped on every run. The skips were correct but mislabelled what they were doing, and the two were not the same thing: - split[1] genuinely fails without its skip. M // 2 is 0 at M=1, so adapter_indices is [2] and the reference returns bitmask 2 where the test asserts 3. One token cannot be split between two adapters. - padding[1] passes. pad_m = (-2) % 1 == 0, so no padding exists and the reference already satisfies both assertions. Neither is a case that merely went unchecked, so a skip overstates it: a tile of one token cannot hold a split, and M % 1 is always 0 so there is never a partial tile. Narrow the parametrization instead -- a second fixture supplies the same sweep without 1 -- and the configurations are simply not generated. Nothing is reported as skipped and nothing asserts the degenerate answer, which would have been worse than either: a green cell pinning arithmetic no caller relies on, at a tiling production never uses (BLOCK_M is 32). BLOCK_M=1 keeps its coverage from the other five tests, including test_different_patterns_per_tile[1], which already exercises the tile boundary padding[1] was nominally about. 36 collected / 2 skipped becomes 34 / 0. Signed-off-by: itzikvaisman --- tests/unit/test_bitmask_computation.py | 50 +++++++++++++++++++------- 1 file changed, 37 insertions(+), 13 deletions(-) diff --git a/tests/unit/test_bitmask_computation.py b/tests/unit/test_bitmask_computation.py index 4f28f891..6857e243 100644 --- a/tests/unit/test_bitmask_computation.py +++ b/tests/unit/test_bitmask_computation.py @@ -11,11 +11,41 @@ import granite_switch.vllm.core.lora_kernel_meta as meta_mod from granite_switch.vllm.core.lora_kernel_meta import _compute_bitmasks_reference +# Production BLOCK_M is 32 (kernels/switch_lora_kernel.py); the others are synthetic +# probes of the tile arithmetic, so these tests keep holding if the kernel retiles. +_BLOCK_M_VALUES = [1, 8, 16, 32, 64] -@pytest.fixture(params=[1, 8, 16, 32, 64]) + +def _use_block_m(monkeypatch, value): + monkeypatch.setattr(meta_mod, "BLOCK_M", value) + return value + + +@pytest.fixture(params=_BLOCK_M_VALUES) def block_m(request, monkeypatch): - monkeypatch.setattr(meta_mod, "BLOCK_M", request.param) - return request.param + return _use_block_m(monkeypatch, request.param) + + +@pytest.fixture(params=[b for b in _BLOCK_M_VALUES if b > 1]) +def block_m_multi_token(request, monkeypatch): + """The same sweep minus BLOCK_M=1: tiles that hold more than one token. + + Two tests need a tile wide enough to hold two *different* things at once, and + at BLOCK_M=1 that is arithmetically impossible rather than merely untested: + + - Splitting a tile between two adapters needs >= 2 tokens. ``M // 2`` is 0 at + M=1, so the "split" collapses to a single adapter and the expected bitmask + is unreachable. + - A partially-filled tile needs ``M % BLOCK_M != 0``, and every M is divisible + by 1, so there is no padding for the test to find spurious bits in. + + Narrowing the sweep rather than skipping inside the test body, because the two + read differently in a report: a skip says "might work, nobody checked", when + the truth is "this configuration does not exist". Asserting the degenerate + outcome instead would be worse than either -- a green cell pinning arithmetic + no caller relies on, for a tiling production never uses. + """ + return _use_block_m(monkeypatch, request.param) def _make_remap_tables_t(remap_tables): @@ -43,12 +73,9 @@ def test_single_adapter_all_positions(self, block_m): class TestMultipleAdaptersOneTile: - def test_two_adapters_split(self, block_m): - if block_m < 2: - pytest.skip("need at least 2 tokens to split between adapters") - + def test_two_adapters_split(self, block_m_multi_token): # One tile: first half adapter 1, second half adapter 2 - M = block_m + M = block_m_multi_token half = M // 2 adapter_indices = torch.cat( [ @@ -160,16 +187,13 @@ def test_more_than_64_adapters_raises(self): class TestPaddingNoSpuriousBits: - def test_partial_tile_padding(self, block_m): - if block_m == 1: - pytest.skip("no padding when block_m=1") - + def test_partial_tile_padding(self, block_m_multi_token): # M = block_m + 1 → two tiles, second tile has 1 real token + padding # All tokens use adapter 1 except the single token in tile 2 # which uses base (0). Padding should not introduce bits. adapter_indices = torch.cat( [ - torch.ones(block_m, dtype=torch.long), # tile 0: adapter 1 + torch.ones(block_m_multi_token, dtype=torch.long), # tile 0: adapter 1 torch.zeros(1, dtype=torch.long), # tile 1: base only ] ) From 40ac57cae9c9f7872cc8addf8d537c788734102e Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Wed, 21 Oct 2026 12:26:08 +0300 Subject: [PATCH 04/15] Run the HF skinning check against a noise bound instead of skipping it test_skinning_equivalence_hf carried an unconditional @pytest.mark.skip whose stated reason was sound -- the skin's fused QKV/gate-up projections sum in a different order than upstream HuggingFace's separate ones, so the worker's `max_diff == 0.0` was unreachable -- but the conclusion drawn was too strong. The sibling vLLM worker had already faced the same arithmetic and answered it with a tolerance (LOGPROB_ATOL) rather than a skip. Do the same here, because both vLLM skinning tests need a GPU and this is the only end-to-end check that the composer's weight transfer reproduces the original model on CPU. The worker now applies three checks, ordered most to least meaningful: 1. argmax agreement, exact -- what greedy decoding actually emits. 2. top-100 token set agreement, exact, as sets so benign tie ordering does not count as a regression. 3. magnitude agreement within ULP_SLACK ULP. Ranking leads because magnitude is the weaker signal: log_softmax is shift-invariant, so adding a constant to every logit leaves the distribution identical while moving max|diff| arbitrarily. Verified -- a uniform +1.5 shift changes log-probabilities by 3.8e-06, i.e. nothing. Check 3 still earns its place by catching the order-preserving faults ranking cannot see, such as a wrong logits_scaling. The bound is derived from the dtype rather than fixed. A flat atol does not transfer from the vLLM worker, which compares logprobs: one ULP of a raw logit measures ~5e-06 in fp32 but ~0.34 in bf16, so any constant tight enough to mean something in fp32 would reject plain arithmetic in bf16. ULP_SLACK is 8 against a measured noise floor near zero -- matmul accumulates in fp32 and rounds once, so two reduction orders give the same value or adjacent ones. Exercised on synthetic logits: identical, fp32 ULP noise, bf16 round-trip and a uniform shift all pass; per-token shifts, swapped positions, a 1.05x scale and a reorder inside the top-100 all fail. The real download-compose-forward path is unverified until a GPU run; if granite-4.0-micro's true noise exceeds 8 ULP the printed max|diff| and bound make that a one-look adjustment. Signed-off-by: itzikvaisman --- CLAUDE.md | 12 +- .../composer/_skinning_equivalence_worker.py | 155 +++++++++++++++--- tests/composer/test_skinning_equivalence.py | 30 ++-- 3 files changed, 162 insertions(+), 35 deletions(-) diff --git a/CLAUDE.md b/CLAUDE.md index 0e3481ca..1ca567a5 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -404,8 +404,16 @@ architectures (4.0, 4.1) use `attention_bias=False` and `mlp_bias=False`. The GraniteSwitch HF backend uses fused QKV and gate-up projections, symmetric with the vLLM backend architecture. Upstream HuggingFace `GraniteMoeHybridForCausalLM` uses separate projections. Fused projections change the floating-point reduction order, so bit-exact skinning equivalence -with the upstream HF model is not achievable. The HF skinning tests in -`tests/composer/test_skinning_equivalence.py` are skipped for this reason. +with the upstream HF model is not achievable. The HF skinning test in +`tests/composer/test_skinning_equivalence.py` used to be skipped outright for this reason; it +now runs and compares the way every other equivalence check here does — exact agreement on +argmax and the top-100 token set, plus a magnitude bound derived from the dtype's epsilon +(`ULP_SLACK` in `_skinning_equivalence_worker.py`) rather than a fixed atol, since one ULP of a +logit is ~5e-06 in fp32 but ~0.34 in bf16. Ranking is the load-bearing part: adding a constant +to every logit leaves the distribution unchanged (`log_softmax` is shift-invariant), so a +magnitude check alone can neither convict nor acquit. Worth keeping rather than deleting because +both vLLM skinning tests need a GPU, making this the only CPU-runnable check that the composer's +weight transfer reproduces the original model. The vLLM check (`tests/vllm/test_generation_equivalence.py`) is the authoritative equivalence test, but it is **not** token-exact either. The skinned model runs its projections through the diff --git a/tests/composer/_skinning_equivalence_worker.py b/tests/composer/_skinning_equivalence_worker.py index 3ad9bc59..0eaf7ce0 100644 --- a/tests/composer/_skinning_equivalence_worker.py +++ b/tests/composer/_skinning_equivalence_worker.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -"""Subprocess worker: verify a skinned GraniteSwitch model is logit-exact vs the original. +"""Subprocess worker: verify a skinned GraniteSwitch model matches the original's logits. Invoked as:: @@ -7,13 +7,24 @@ Loads the original model in its **native dtype** (the precision it was published in), computes reference logits, frees it, then builds a GraniteSwitch skin -(num_adapters=0) via the builder's weight-transfer pipeline, and asserts -bit-exact equivalence. +(num_adapters=0) via the builder's weight-transfer pipeline, and asserts the two +agree. Native dtype is critical: skinned models must be equivalent in the precision users will actually deploy. Float32 may mask real differences caused by attention scaling order (pre-multiply Q vs post-divide scores). +AGREEMENT IS CHECKED WITHIN A TOLERANCE, NOT BIT-FOR-BIT. This worker asserted +``max_diff == 0.0`` until it was found to be asking for something unreachable: +the skin uses fused QKV / gate-up projections while upstream HuggingFace uses +separate ones, and fusing changes the order of a floating-point sum. ``a+b+c`` +and ``(a+b)+c`` are not bit-identical, so a few ULP of disagreement is arithmetic, +not a defect, and the test could only ever fail. It sat behind an unconditional +``@pytest.mark.skip`` for that reason, which cost the composer its only +weight-transfer check that runs without a GPU. + +See ``_compare_logits`` for what replaced it and why that still catches real bugs. + Runs in its own process so CUDA context is fully torn down on exit. """ @@ -28,6 +39,31 @@ SEQ_LEN = 12 # Short fixed sequence — enough to exercise all layer types. +# How many ULP of disagreement to allow, as a multiple of the dtype's epsilon scaled +# by the logit magnitude. DERIVED, not a fixed atol: the sibling vLLM worker can use a +# flat 1e-5 because it compares logprobs, but this one compares raw logits in the +# model's native dtype, and one ULP there spans four orders of magnitude depending on +# dtype -- measured, at a logit of ~43 it is 5.2e-06 in fp32 but 3.4e-01 in bf16. A +# flat tolerance tight enough to mean anything in fp32 would reject plain arithmetic +# in bf16, and one loose enough for bf16 would be vacuous in fp32. +# +# 8 ULP, because the measured noise is far below that: a matmul accumulates in fp32 +# and rounds once at the end, so two reduction orders usually produce the SAME output +# value and at worst adjacent ones. Simulating fused-vs-blockwise projections at +# H=2048 gave a max difference of 0 ULP in both dtypes, and an fp32 chunked-vs-whole +# dot product differed by ~1.4 ULP. 8 leaves room for a wider model without being the +# kind of bound that passes anything. +ULP_SLACK = 8 + +# Positions must also agree on WHICH tokens are most likely, not just on magnitudes. +# This is the load-bearing check, and the magnitude bound above is the weaker one: a +# constant added to every logit shifts max|diff| arbitrarily far while being invisible +# after softmax (log_softmax is shift-invariant), so magnitude alone can neither +# convict nor acquit. Ranking cannot be fooled that way -- a wrong weight permutation +# or a swapped K/V reorders the distribution, which is what a caller actually sees. +# Mirrors the vLLM worker, which compares top-K sets for the same reason. +TOP_K = 100 + def _native_dtype(config): """Determine the model's native dtype from its HuggingFace config. @@ -45,6 +81,99 @@ def _native_dtype(config): return torch.float32 +def _compare_logits(reference, switch, label, dtype): + """Check skinned logits against the original's. Returns 0 on agreement, 1 on diff. + + Three checks, ordered most to least meaningful. They fail for different reasons, + and reporting them separately is what makes a failure diagnosable rather than + just "the numbers moved": + + 1. **argmax agreement, exact.** Every position must rank the same token first. + No tolerance -- this is what greedy decoding actually emits, and a skin that + flips the top token is wrong however small the gap that flipped it. + 2. **top-K set agreement, exact.** The ``TOP_K`` most likely tokens must be the + same set at every position. Catches a distribution that is subtly reordered + below the argmax, which the vLLM worker checks for the same reason. + 3. **magnitude agreement, within ``ULP_SLACK`` ULP.** The weakest of the three + and deliberately last: a uniform shift of every logit leaves the + distribution identical (log_softmax is shift-invariant) yet moves max|diff| + arbitrarily, so this check alone can neither convict nor acquit. It earns + its place by catching the case the ranking checks cannot -- a scale change + that preserves order, e.g. a wrong ``logits_scaling`` or + ``attention_multiplier``. + + All three are printed either way, so a run drifting toward the bound is visible + before it crosses it. + """ + # Compare in fp32: a bf16 subtraction would itself round the difference being + # measured, and the two tensors are already materialized at native precision. + ref = reference.float() + sw = switch.float() + + diff = (ref - sw).abs() + max_diff = diff.max().item() + mean_diff = diff.mean().item() + num_nonzero = int((diff > 0).sum().item()) + total = diff.numel() + + # Tolerance from the dtype's resolution at this model's logit magnitude. + scale = max(ref.abs().max().item(), 1.0) + tol = ULP_SLACK * torch.finfo(dtype).eps * scale + + ref_top = ref.argmax(dim=-1) + sw_top = sw.argmax(dim=-1) + argmax_mismatches = int((ref_top != sw_top).sum().item()) + + # Top-K as a set per position: ties can order two equal logits either way, which + # is not a regression, so compare membership rather than the sorted sequence. + k = min(TOP_K, ref.shape[-1]) + ref_k = ref.topk(k, dim=-1).indices + sw_k = sw.topk(k, dim=-1).indices + flat_ref = ref_k.reshape(-1, k) + flat_sw = sw_k.reshape(-1, k) + topk_mismatches = sum( + 1 + for i in range(flat_ref.shape[0]) + if set(flat_ref[i].tolist()) != set(flat_sw[i].tolist()) + ) + + print( + f" dtype = {dtype} (1 ULP at |logit|<={scale:.3g} is {tol / ULP_SLACK:.3e})" + ) + print(f" max |diff| = {max_diff:.6e} tolerance = {tol:.6e} ({ULP_SLACK} ULP)") + print(f" mean |diff| = {mean_diff:.6e}") + print(f" differing elements: {num_nonzero} / {total}") + print(f" argmax mismatches: {argmax_mismatches} / {ref_top.numel()}") + print(f" top-{k} set mismatches: {topk_mismatches} / {flat_ref.shape[0]}") + + if argmax_mismatches: + print( + f"\nFAIL: {label} — {argmax_mismatches} position(s) predict a different " + f"top token; the skin is not equivalent (dtype={dtype})" + ) + return 1 + if topk_mismatches: + print( + f"\nFAIL: {label} — {topk_mismatches} position(s) have a different top-{k} " + f"token set; the distribution is reordered (dtype={dtype})" + ) + return 1 + if max_diff > tol: + print( + f"\nFAIL: {label} — logits differ by more than reduction-order noise " + f"(max |diff| = {max_diff:.6e} > {tol:.6e}, dtype={dtype})" + ) + return 1 + if max_diff == 0.0: + print(f"\nPASS: {label} — bit-exact equivalence ({dtype})") + return 0 + print( + f"\nPASS: {label} — equivalent within {ULP_SLACK} ULP " + f"(max |diff| = {max_diff:.6e} <= {tol:.6e}, dtype={dtype})" + ) + return 0 + + def main(): parser = argparse.ArgumentParser(description=__doc__) parser.add_argument("model", help="HuggingFace model name or local path") @@ -93,25 +222,7 @@ def main(): # ── Phase 3: compare ────────────────────────────────────────── print("\nPhase 3: comparing logits...") - - diff = (switch_logits - reference_logits).abs() - max_diff = diff.max().item() - mean_diff = diff.mean().item() - num_nonzero = (diff > 0).sum().item() - total = diff.numel() - - print(f" max |diff| = {max_diff:.6e}") - print(f" mean |diff| = {mean_diff:.6e}") - print(f" nonzero elements: {num_nonzero} / {total}") - - if max_diff == 0.0: - print(f"\nPASS: {model_name} — bit-exact equivalence ({dtype})") - return 0 - - print( - f"\nFAIL: {model_name} — logits differ (max |diff| = {max_diff:.6e}, dtype={dtype})" - ) - return 1 + return _compare_logits(reference_logits, switch_logits, model_name, dtype) if __name__ == "__main__": diff --git a/tests/composer/test_skinning_equivalence.py b/tests/composer/test_skinning_equivalence.py index 9be11867..6d8fc617 100644 --- a/tests/composer/test_skinning_equivalence.py +++ b/tests/composer/test_skinning_equivalence.py @@ -1,5 +1,10 @@ # SPDX-License-Identifier: Apache-2.0 -"""Verify skinned GraniteSwitch models produce bit-exact logits vs the originals. +"""Verify skinned GraniteSwitch models reproduce the originals' logits. + +Neither backend checks bit-exactness: the skin's fused projections sum in a +different order than the unfused reference, so both workers compare within a +noise bound (an ULP multiple for HF logits, LOGPROB_ATOL for vLLM logprobs) and +pair it with an exact top-token check. See CLAUDE.md Gotcha #9. Each model runs in its own subprocess so CUDA context is fully torn down between tests (avoids OOM and stale-state issues that plague in-process model reloading). @@ -174,22 +179,25 @@ def _run_vllm_test(model_name, timeout, fast=False): # ── HF tests (slow) ────────────────────────────────────────────── # -# SKIPPED: The GraniteSwitch HF backend uses fused QKV/gate-up projections -# (symmetric with vLLM) rather than the separate projections used by upstream -# HuggingFace GraniteMoeHybrid. Fused projections change the floating-point -# reduction order, so bit-exact equivalence with the upstream HF model is not -# achievable. The vLLM skinning tests below are the authoritative equivalence -# check — both sides use the same fused-projection architecture. +# The GraniteSwitch HF backend uses fused QKV/gate-up projections (symmetric with +# vLLM) rather than the separate projections upstream HuggingFace uses, and fusing +# changes the order of a floating-point sum. This test therefore used to carry an +# unconditional @pytest.mark.skip, on the correct observation that bit-exact +# equivalence is unreachable -- but the conclusion drawn from it was too strong. +# Its worker now compares within a dtype-derived ULP bound plus an exact argmax +# check (see _skinning_equivalence_worker._compare_logits), which is the same move +# the vLLM worker already made with LOGPROB_ATOL, so the test can run instead. +# +# Worth having rather than deleting: both vLLM skinning tests below need a GPU, so +# this is the only end-to-end check that the composer's weight transfer produces a +# model matching the original, and it runs on CPU. -@pytest.mark.skip( - reason="HF backend uses fused projections (not bit-exact with upstream HF)" -) @pytest.mark.slow @pytest.mark.requires_model @pytest.mark.parametrize("model_name", ALL_MODELS, ids=_short_name) def test_skinning_equivalence_hf(model_name): - """Skin *model_name* to GraniteSwitch and assert bit-exact logits via HF.""" + """Skin *model_name* to GraniteSwitch and compare logits via HF (CPU).""" _run_worker(HF_WORKER, model_name, HF_TIMEOUT) From 98b8617ba92dcbb74901995959232aa2ec632930 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Wed, 21 Oct 2026 12:44:56 +0300 Subject: [PATCH 05/15] Add a `skips` GPU scope for the tests the audit un-gated The three files changed by the skip audit cannot be verified on a CPU box, and `full` is not currently a usable way to reach them: every full-suite dispatch since 2026-10-18 has been cancelled from outside the pod at ~95-120 min, which stops a 3h46m suite around 40% and reports no summary. The limit is none of the three this workflow configures (360/390/420 min). So give them a narrow scope that finishes well inside that window. What each path is here to prove: tests/hf/test_quantization.py 16 tests that used to skip on an env flag before importorskip could run; 4 now gate on a Quanto float8 probe whose answer only exists on a GPU. tests/composer/test_skinning_... was @pytest.mark.skip outright; now compares within a dtype-derived ULP bound chosen against simulated noise, not the real model's. tests/unit/test_bitmask_computation cheap, and not collectable off-GPU at all: it imports lora_kernel_meta, which imports vllm.distributed. Deliberately not wired into the /gpu-test comment path's defaults, and not a substitute for `full`. Remove this arm once the ~95 min cancellation is diagnosed and `full` can be trusted again. Signed-off-by: itzikvaisman --- .github/workflows/gpu-tests.yaml | 34 ++++++++++++++++++++++++++++++-- 1 file changed, 32 insertions(+), 2 deletions(-) diff --git a/.github/workflows/gpu-tests.yaml b/.github/workflows/gpu-tests.yaml index b5d0c5cf..22dd8462 100644 --- a/.github/workflows/gpu-tests.yaml +++ b/.github/workflows/gpu-tests.yaml @@ -60,11 +60,11 @@ on: # family is added HERE and nowhere else -- this entry plus its arm in the # Resolve test scope step below. suite: - description: "Which tests to run: full (all five suites), short (vllm + integration), dev (one fast GPU file), audio (the audio/ASR marker), multi (all five on 2 GPUs)." + description: "Which tests to run: full (all five suites), short (vllm + integration), dev (one fast GPU file), audio (the audio/ASR marker), multi (all five on 2 GPUs), skips (the formerly-skipped tests being un-gated)." required: false default: full type: choice - options: [full, short, dev, audio, multi] + options: [full, short, dev, audio, multi, skips] # RESULT CACHE OPT-OUT. Default off, i.e. cache on: the expensive thing here is # a GPU run, and re-running one for a commit+scope that was already tested buys # nothing in the common case (a reviewer re-reading a PR, a second maintainer @@ -238,6 +238,36 @@ jobs: SUFFIX="audio" LEG="$LABEL-audio" ;; + # The tests that the skip audit un-gated, as one narrow scope. Exists + # because `full` is not currently a usable way to check them: every + # full-suite run since 2026-10-18 has been cancelled from outside the pod + # at ~95-120 min (see the three-timeout note above -- it is none of + # ours), so a 3h46m suite is stopped around 40% and reports no summary. + # These three paths finish far inside that window. + # + # What each one is here to prove, since all three were changed by edits + # that no CPU box can check: + # tests/hf/test_quantization.py 16 tests that used to skip on an + # env flag; 4 of them now gate on a + # real Quanto float8 probe that only + # means anything on a GPU. + # tests/composer/test_skinning_...py was @pytest.mark.skip outright; + # now compares within a dtype-derived + # ULP bound whose real noise floor is + # unmeasured until it runs. + # tests/unit/test_bitmask_computation 34 tests, no longer 2 skips. Cheap, + # and it cannot even be COLLECTED off + # GPU: it imports lora_kernel_meta, + # which imports vllm.distributed. + # + # NOT a substitute for `full`, and deliberately not wired into the + # /gpu-test comment path's defaults. Delete this arm once the ~95 min + # cancellation is fixed and `full` is trustworthy again. + skips) + PATHS="tests/hf/test_quantization.py tests/composer/test_skinning_equivalence.py tests/unit/test_bitmask_computation.py" + SUFFIX="skips" + LEG="$LABEL-skips" + ;; # Everything `full` runs, but on TWO GPUs. Not a different test # selection -- a different SHAPE of pod. # From 42730b79d55763f3c617dd2f8699dc801d23ff35 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Wed, 21 Oct 2026 14:51:12 +0300 Subject: [PATCH 06/15] Gate the HF skinning test on distribution agreement, not exact top-K The previous commit un-skipped this test by comparing logit magnitudes within an 8-ULP bound plus an exact top-100 token-set match. The GPU leg showed the set check was unreachable: max |diff| = 7.5e-01 tolerance = 9.1e-01 (8 ULP) <- passed argmax mismatches: 0 / 12 <- passed top-100 set mismatches: 11 / 12 <- FAILED bf16 resolution is ~0.11 at these logits, so the rank-100 boundary is a pile of exact ties -- measured 4, 7 and 7 tokens sharing one bf16 value at the cut on the first three positions. Which tie-mate lands inside the cut is arbitrary, not a regression. Comparing sets rather than sorted order was an attempt at tie-robustness, but membership *at the boundary* is exactly as tie-sensitive as order. Replace both with the gate tests/vllm/_generation_equivalence_worker.py already calibrated -- same question, so a second set of thresholds would just be a second thing to keep calibrated: top-k JSD + Jaccard swept over k=1..100, plus the exact argmax check (measured a strict maximum at all 12 positions, so safe to demand exactly). The magnitude bound was the weaker problem -- not wrong, just near-vacuous, since a uniform shift moves max|diff| arbitrarily while leaving the distribution identical. Those three are all post-log_softmax, which leaves them blind to a uniform logit SCALING: a 5% error measures mean JSD 7.2e-04, inside the benign floor, while being the signature of a mistransferred logits_scaling (typically 8.0, so one step is ~12%) or attention_multiplier. Add a least-squares slope fit with free intercept, which separates it because a scale error is systematic across all 1.2M elements where reduction-order noise is random: 3.2e-05 for the real skin versus 1.0e-02 for a 1% scaling. Verified against 11 cases through the real comparator: the actual skin, an identical copy, a +1.5 uniform shift and a bf16 round-trip pass; a per-token shift, 1%/5% scaling, logits_scaling 8.0->8.1, two swapped positions, one corrupted position, a 100-token weight permutation and a subtle +0.9 all fail. Real skin passes with 27x margin (max JSD 1.9e-03 against a 0.05 gate). Signed-off-by: itzikvaisman --- .github/workflows/gpu-tests.yaml | 10 +- CLAUDE.md | 31 ++- .../composer/_skinning_equivalence_worker.py | 255 +++++++++++------- tests/composer/test_skinning_equivalence.py | 12 +- 4 files changed, 197 insertions(+), 111 deletions(-) diff --git a/.github/workflows/gpu-tests.yaml b/.github/workflows/gpu-tests.yaml index 22dd8462..ff8924ff 100644 --- a/.github/workflows/gpu-tests.yaml +++ b/.github/workflows/gpu-tests.yaml @@ -251,10 +251,12 @@ jobs: # env flag; 4 of them now gate on a # real Quanto float8 probe that only # means anything on a GPU. - # tests/composer/test_skinning_...py was @pytest.mark.skip outright; - # now compares within a dtype-derived - # ULP bound whose real noise floor is - # unmeasured until it runs. + # tests/composer/test_skinning_...py was @pytest.mark.skip outright; now + # gates top-k JSD/Jaccard + argmax + + # a slope fit. This scope is what + # measured the bf16 noise floor those + # are calibrated against -- a CPU box + # runs it in fp32 and sees none of it. # tests/unit/test_bitmask_computation 34 tests, no longer 2 skips. Cheap, # and it cannot even be COLLECTED off # GPU: it imports lora_kernel_meta, diff --git a/CLAUDE.md b/CLAUDE.md index 1ca567a5..8e9b97f9 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -406,14 +406,29 @@ backend architecture. Upstream HuggingFace `GraniteMoeHybridForCausalLM` uses se Fused projections change the floating-point reduction order, so bit-exact skinning equivalence with the upstream HF model is not achievable. The HF skinning test in `tests/composer/test_skinning_equivalence.py` used to be skipped outright for this reason; it -now runs and compares the way every other equivalence check here does — exact agreement on -argmax and the top-100 token set, plus a magnitude bound derived from the dtype's epsilon -(`ULP_SLACK` in `_skinning_equivalence_worker.py`) rather than a fixed atol, since one ULP of a -logit is ~5e-06 in fp32 but ~0.34 in bf16. Ranking is the load-bearing part: adding a constant -to every logit leaves the distribution unchanged (`log_softmax` is shift-invariant), so a -magnitude check alone can neither convict nor acquit. Worth keeping rather than deleting because -both vLLM skinning tests need a GPU, making this the only CPU-runnable check that the composer's -weight transfer reproduces the original model. +now runs and gates **distribution** equivalence instead, with the same metric and thresholds +`tests/vllm/_generation_equivalence_worker.py` already calibrated: top-k JSD + Jaccard swept +over k, plus an exact argmax check and a least-squares slope fit. Four checks because the first +three live after `log_softmax` and the fourth covers what normalizing hides: + +- **argmax, exact** — what greedy decoding emits; measured to be a strict maximum at every + position, so no tolerance is needed even in bf16. +- **mean + max top-k JSD** — mean catches "wrong everywhere a little", max catches one + catastrophically wrong position that a mean would dilute. +- **mean top-k Jaccard distance**, deliberately loose. An exact top-K *set* match is + unreachable in bf16 and was the first attempt: bf16 resolution is ~0.11 at these logits, so + the rank-100 boundary is a pile of exact ties (measured 4, 7 and 7 tokens sharing one bf16 + value at the cut) and which tie-mate lands inside the cut is arbitrary. Jaccard measures how + much the top-k overlaps rather than demanding it match. +- **slope of skinned-vs-reference logits** (`SLOPE_TOL`) — the only non-scale-invariant check, + which is its whole job. A uniform logit *scaling* preserves ranking and barely dents the + normalized shape (a 5% error measures inside the benign JSD floor), yet it is the signature of + a mistransferred `logits_scaling` / `attention_multiplier`. A scale error is systematic where + reduction-order noise is random, so the fit separates them by ~300x. Intercept left free, so a + benign uniform shift still reads as slope 1. + +Worth keeping rather than deleting because both vLLM skinning tests need a GPU, making this the +only CPU-runnable check that the composer's weight transfer reproduces the original model. The vLLM check (`tests/vllm/test_generation_equivalence.py`) is the authoritative equivalence test, but it is **not** token-exact either. The skinned model runs its projections through the diff --git a/tests/composer/_skinning_equivalence_worker.py b/tests/composer/_skinning_equivalence_worker.py index 0eaf7ce0..5c895f49 100644 --- a/tests/composer/_skinning_equivalence_worker.py +++ b/tests/composer/_skinning_equivalence_worker.py @@ -14,22 +14,26 @@ users will actually deploy. Float32 may mask real differences caused by attention scaling order (pre-multiply Q vs post-divide scores). -AGREEMENT IS CHECKED WITHIN A TOLERANCE, NOT BIT-FOR-BIT. This worker asserted +DISTRIBUTION AGREEMENT IS CHECKED, NOT BIT-FOR-BIT EQUALITY. This worker asserted ``max_diff == 0.0`` until it was found to be asking for something unreachable: the skin uses fused QKV / gate-up projections while upstream HuggingFace uses separate ones, and fusing changes the order of a floating-point sum. ``a+b+c`` -and ``(a+b)+c`` are not bit-identical, so a few ULP of disagreement is arithmetic, +and ``(a+b)+c`` are not bit-identical, so a small disagreement is arithmetic, not a defect, and the test could only ever fail. It sat behind an unconditional ``@pytest.mark.skip`` for that reason, which cost the composer its only weight-transfer check that runs without a GPU. -See ``_compare_logits`` for what replaced it and why that still catches real bugs. +What replaced it is the same gate the vLLM generation-equivalence worker uses +(top-k JSD + Jaccard over a swept k), plus an exact argmax check and a slope fit. +See the constants below for the measured noise floor those are calibrated against, +and ``_compare_logits`` for why each of the four is there. Runs in its own process so CUDA context is fully torn down on exit. """ import argparse import gc +import os import sys import torch @@ -37,32 +41,75 @@ from granite_switch.composer import GraniteSwitchComposer +# Make tests.shared importable when run as a bare subprocess (cwd-independent), +# the same way the vLLM generation-equivalence worker does it. +_REPO = os.path.dirname(os.path.dirname(os.path.dirname(os.path.abspath(__file__)))) +if _REPO not in sys.path: + sys.path.insert(0, _REPO) + +from tests.shared.logit_metrics import jaccard, jsd_bits + SEQ_LEN = 12 # Short fixed sequence — enough to exercise all layer types. -# How many ULP of disagreement to allow, as a multiple of the dtype's epsilon scaled -# by the logit magnitude. DERIVED, not a fixed atol: the sibling vLLM worker can use a -# flat 1e-5 because it compares logprobs, but this one compares raw logits in the -# model's native dtype, and one ULP there spans four orders of magnitude depending on -# dtype -- measured, at a logit of ~43 it is 5.2e-06 in fp32 but 3.4e-01 in bf16. A -# flat tolerance tight enough to mean anything in fp32 would reject plain arithmetic -# in bf16, and one loose enough for bf16 would be vacuous in fp32. +# Distribution-agreement gates, and the same three the vLLM generation-equivalence +# worker uses -- deliberately, since both are asking the identical question (does a +# fused-projection skin match the unfused original) and a second set of numbers here +# would just be a second thing to keep calibrated. +# +# WHAT THESE REPLACED, AND WHY. This file first compared raw logit MAGNITUDES within +# an 8-ULP bound plus an exact top-100 SET match. Both were wrong for bf16 and the +# GPU run that caught it is worth recording: +# +# max |diff| = 7.5e-01 tolerance = 9.1e-01 (8 ULP) <- passed +# argmax mismatches: 0 / 12 <- passed +# top-100 set mismatches: 11 / 12 <- FAILED +# +# The top-K set check fails on a correct skin because bf16 has ~0.11 resolution at +# these logits, so the rank-100 boundary is a pile of EXACT TIES -- measured 4, 7 and +# 7 tokens sharing the identical bf16 value at the cut on the first three positions. +# Which of those tie-mates lands inside the cut is arbitrary, not a regression. +# Comparing sets rather than sorted order was an attempt to be tie-robust, but +# membership *at the boundary* is exactly as tie-sensitive as order is. Jaccard over +# a swept k replaces it: it measures HOW MUCH the top-k overlaps instead of demanding +# it match exactly, so a few tie-mates trading places costs a few percent instead of +# failing outright. # -# 8 ULP, because the measured noise is far below that: a matmul accumulates in fp32 -# and rounds once at the end, so two reduction orders usually produce the SAME output -# value and at worst adjacent ones. Simulating fused-vs-blockwise projections at -# H=2048 gave a max difference of 0 ULP in both dtypes, and an fp32 chunked-vs-whole -# dot product differed by ~1.4 ULP. 8 leaves room for a wider model without being the -# kind of bound that passes anything. -ULP_SLACK = 8 - -# Positions must also agree on WHICH tokens are most likely, not just on magnitudes. -# This is the load-bearing check, and the magnitude bound above is the weaker one: a -# constant added to every logit shifts max|diff| arbitrarily far while being invisible -# after softmax (log_softmax is shift-invariant), so magnitude alone can neither -# convict nor acquit. Ranking cannot be fooled that way -- a wrong weight permutation -# or a swapped K/V reorders the distribution, which is what a caller actually sees. -# Mirrors the vLLM worker, which compares top-K sets for the same reason. -TOP_K = 100 +# The magnitude bound was the weaker problem: it is not wrong, just nearly vacuous. +# A uniform shift of every logit moves max|diff| arbitrarily while leaving the +# distribution identical (log_softmax is shift-invariant), so it can neither convict +# nor acquit. JSD is measured on log_softmax output and is therefore immune to that +# by construction. +# +# MEASURED on granite-4.0-micro (bf16, 12 positions, k swept over 1..100): mean JSD +# 6.5e-04, max JSD 1.9e-03, mean(1-Jaccard) 0.094. The gates sit ~2-25x above that +# floor. Verified to still convict: a swapped K/V (max JSD 0.52), a 100-token weight +# permutation (0.77), one corrupted position (0.046), a +0.9 per-token shift (mean +# JSD 0.030) and a 30% logit scaling (0.024) all fail. +MEAN_JSD_THRESH = 0.003 +MAX_JSD_THRESH = 0.05 +JACC_THRESH = 0.20 + +# Swept rather than fixed, because the two failure modes live at opposite ends: a +# flipped near-tie only shows up at small k, and a reordered tail only at large k. +K_SWEEP = (1, 5, 10, 20, 50, 100) + +# Max allowed |a-1| where `a` is the least-squares slope of skinned-vs-reference +# logits. +# +# This check exists because the three gates above are all computed AFTER +# log_softmax, which makes them blind to one real bug: a uniform SCALING of every +# logit. Scaling preserves ranking (so argmax and Jaccard see nothing) and barely +# dents the normalized shape at small factors -- a 5% error measured mean JSD +# 7.2e-04, inside the benign floor. That is not a hypothetical: it is the signature +# of a wrong ``logits_scaling`` or ``attention_multiplier``, the two Granite config +# multipliers a skin could plausibly mistransfer, and `logits_scaling` is typically +# 8.0 so a one-step error is a ~12% scale change. +# +# A slope fit separates it cleanly because a scale error is SYSTEMATIC across all +# 1.2M elements while reduction-order noise is random: measured |a-1| = 3.2e-05 for +# the real skin versus 1.0e-02 for a 1% scaling -- a 300x gap. The intercept is left +# free, so a benign uniform shift still reads as slope 1 exactly. +SLOPE_TOL = 1e-3 def _native_dtype(config): @@ -81,70 +128,92 @@ def _native_dtype(config): return torch.float32 +def _slope(ref, sw): + """Least-squares slope of ``sw`` against ``ref``, intercept free. + + Intercept free on purpose: a uniform additive shift is invisible after softmax, + so it must not register here either. Returns 1.0 exactly for a shifted copy. + """ + x = ref.reshape(-1) + y = sw.reshape(-1) + xc = x - x.mean() + return float((xc * (y - y.mean())).sum() / (xc * xc).sum()) + + def _compare_logits(reference, switch, label, dtype): """Check skinned logits against the original's. Returns 0 on agreement, 1 on diff. - Three checks, ordered most to least meaningful. They fail for different reasons, - and reporting them separately is what makes a failure diagnosable rather than - just "the numbers moved": + Four checks. The first three are computed on ``log_softmax`` output -- the + distribution a caller actually observes -- and the fourth covers the one real + bug that normalizing hides. They fail for different reasons, and reporting them + separately is what makes a failure diagnosable rather than just "the numbers + moved": 1. **argmax agreement, exact.** Every position must rank the same token first. - No tolerance -- this is what greedy decoding actually emits, and a skin that - flips the top token is wrong however small the gap that flipped it. - 2. **top-K set agreement, exact.** The ``TOP_K`` most likely tokens must be the - same set at every position. Catches a distribution that is subtly reordered - below the argmax, which the vLLM worker checks for the same reason. - 3. **magnitude agreement, within ``ULP_SLACK`` ULP.** The weakest of the three - and deliberately last: a uniform shift of every logit leaves the - distribution identical (log_softmax is shift-invariant) yet moves max|diff| - arbitrarily, so this check alone can neither convict nor acquit. It earns - its place by catching the case the ranking checks cannot -- a scale change - that preserves order, e.g. a wrong ``logits_scaling`` or - ``attention_multiplier``. - - All three are printed either way, so a run drifting toward the bound is visible + No tolerance: this is what greedy decoding emits, and a skin that flips the + top token is wrong however small the gap that flipped it. Safe to demand + exactly even in bf16 -- measured, the argmax is a strict maximum at all 12 + positions (no tie-mates), unlike the rank-100 boundary. + 2. **mean and max JSD over a swept k.** The primary signal. Mean catches a + distribution that is wrong everywhere a little; max catches one position + that is catastrophically wrong, which a mean over 12 positions would dilute. + 3. **mean top-k Jaccard distance.** A loose set-overlap guard, same role as in + the vLLM worker. Deliberately loose: at these bf16 logits several tokens tie + exactly at the k-th rank, so a few percent of churn is arithmetic. + 4. **slope of skinned-vs-reference logits.** The only check here that is NOT + scale-invariant, which is precisely its job -- see ``SLOPE_TOL``. + + All four are printed either way, so a run drifting toward a bound is visible before it crosses it. """ # Compare in fp32: a bf16 subtraction would itself round the difference being # measured, and the two tensors are already materialized at native precision. - ref = reference.float() - sw = switch.float() - - diff = (ref - sw).abs() - max_diff = diff.max().item() - mean_diff = diff.mean().item() - num_nonzero = int((diff > 0).sum().item()) - total = diff.numel() - - # Tolerance from the dtype's resolution at this model's logit magnitude. - scale = max(ref.abs().max().item(), 1.0) - tol = ULP_SLACK * torch.finfo(dtype).eps * scale - - ref_top = ref.argmax(dim=-1) - sw_top = sw.argmax(dim=-1) - argmax_mismatches = int((ref_top != sw_top).sum().item()) - - # Top-K as a set per position: ties can order two equal logits either way, which - # is not a regression, so compare membership rather than the sorted sequence. - k = min(TOP_K, ref.shape[-1]) - ref_k = ref.topk(k, dim=-1).indices - sw_k = sw.topk(k, dim=-1).indices - flat_ref = ref_k.reshape(-1, k) - flat_sw = sw_k.reshape(-1, k) - topk_mismatches = sum( - 1 - for i in range(flat_ref.shape[0]) - if set(flat_ref[i].tolist()) != set(flat_sw[i].tolist()) - ) - - print( - f" dtype = {dtype} (1 ULP at |logit|<={scale:.3g} is {tol / ULP_SLACK:.3e})" - ) - print(f" max |diff| = {max_diff:.6e} tolerance = {tol:.6e} ({ULP_SLACK} ULP)") - print(f" mean |diff| = {mean_diff:.6e}") - print(f" differing elements: {num_nonzero} / {total}") - print(f" argmax mismatches: {argmax_mismatches} / {ref_top.numel()}") - print(f" top-{k} set mismatches: {topk_mismatches} / {flat_ref.shape[0]}") + ref = reference.float()[0] + sw = switch.float()[0] + + ref_lp = torch.log_softmax(ref, dim=-1) + sw_lp = torch.log_softmax(sw, dim=-1) + + argmax_mismatches = int((ref.argmax(dim=-1) != sw.argmax(dim=-1)).sum().item()) + slope = _slope(ref, sw) + num_pos = ref.shape[0] + + print(f" dtype = {dtype} positions = {num_pos} vocab = {ref.shape[-1]}") + print(f" argmax mismatches: {argmax_mismatches} / {num_pos}") + print(f" slope |a-1| = {abs(slope - 1):.3e} tolerance = {SLOPE_TOL:.3e}") + print(f" {'k':>5}{'mean(1-Jacc)':>15}{'mean JSD':>12}{'max JSD':>12}{'@pos':>7}") + + failures = [] + for k in K_SWEEP: + if k > ref.shape[-1]: + continue + jacc_dists = [] + jsds = [] + for pos in range(num_pos): + a = ref[pos].topk(k).indices.tolist() + b = sw[pos].topk(k).indices.tolist() + ids = sorted(set(a) | set(b)) + # Dicts only over the union, not the whole 100k vocab: jsd_bits + # renormalizes over `ids` anyway, so the tail would be discarded. + p_lp = {t: ref_lp[pos, t].item() for t in ids} + q_lp = {t: sw_lp[pos, t].item() for t in ids} + jacc_dists.append(1.0 - jaccard(a, b)) + jsds.append(jsd_bits(p_lp, q_lp, ids)) + + mean_jd = sum(jacc_dists) / len(jacc_dists) + mean_js = sum(jsds) / len(jsds) + max_js = max(jsds) + max_pos = jsds.index(max_js) + print(f" {k:5d}{mean_jd:15.5f}{mean_js:12.6f}{max_js:12.6f}{max_pos:7d}") + + if mean_jd > JACC_THRESH: + failures.append(f"k={k}: mean(1-Jaccard)={mean_jd:.4f} > {JACC_THRESH}") + if mean_js > MEAN_JSD_THRESH: + failures.append(f"k={k}: mean JSD={mean_js:.6f} > {MEAN_JSD_THRESH}") + if max_js > MAX_JSD_THRESH: + failures.append( + f"k={k}: max JSD={max_js:.6f} @pos {max_pos} > {MAX_JSD_THRESH}" + ) if argmax_mismatches: print( @@ -152,24 +221,22 @@ def _compare_logits(reference, switch, label, dtype): f"top token; the skin is not equivalent (dtype={dtype})" ) return 1 - if topk_mismatches: + if abs(slope - 1) > SLOPE_TOL: print( - f"\nFAIL: {label} — {topk_mismatches} position(s) have a different top-{k} " - f"token set; the distribution is reordered (dtype={dtype})" + f"\nFAIL: {label} — skinned logits are scaled by {slope:.6f} relative to " + f"the original (|a-1| = {abs(slope - 1):.3e} > {SLOPE_TOL:.3e}); suspect a " + f"mistransferred logits_scaling / attention_multiplier (dtype={dtype})" ) return 1 - if max_diff > tol: - print( - f"\nFAIL: {label} — logits differ by more than reduction-order noise " - f"(max |diff| = {max_diff:.6e} > {tol:.6e}, dtype={dtype})" - ) + if failures: + print(f"\nFAIL: {label} — distributions disagree beyond the noise floor:") + for line in failures: + print(f" {line}") return 1 - if max_diff == 0.0: - print(f"\nPASS: {label} — bit-exact equivalence ({dtype})") - return 0 print( - f"\nPASS: {label} — equivalent within {ULP_SLACK} ULP " - f"(max |diff| = {max_diff:.6e} <= {tol:.6e}, dtype={dtype})" + f"\nPASS: {label} — distribution-equivalent (mean JSD <= {MEAN_JSD_THRESH}, " + f"max JSD <= {MAX_JSD_THRESH}, mean 1-Jaccard <= {JACC_THRESH}, " + f"slope within {SLOPE_TOL:.0e}, dtype={dtype})" ) return 0 diff --git a/tests/composer/test_skinning_equivalence.py b/tests/composer/test_skinning_equivalence.py index 6d8fc617..bb0adb0b 100644 --- a/tests/composer/test_skinning_equivalence.py +++ b/tests/composer/test_skinning_equivalence.py @@ -3,8 +3,8 @@ Neither backend checks bit-exactness: the skin's fused projections sum in a different order than the unfused reference, so both workers compare within a -noise bound (an ULP multiple for HF logits, LOGPROB_ATOL for vLLM logprobs) and -pair it with an exact top-token check. See CLAUDE.md Gotcha #9. +measured noise bound and pair it with an exact top-token check. See CLAUDE.md +Gotcha #9. Each model runs in its own subprocess so CUDA context is fully torn down between tests (avoids OOM and stale-state issues that plague in-process model reloading). @@ -184,9 +184,11 @@ def _run_vllm_test(model_name, timeout, fast=False): # changes the order of a floating-point sum. This test therefore used to carry an # unconditional @pytest.mark.skip, on the correct observation that bit-exact # equivalence is unreachable -- but the conclusion drawn from it was too strong. -# Its worker now compares within a dtype-derived ULP bound plus an exact argmax -# check (see _skinning_equivalence_worker._compare_logits), which is the same move -# the vLLM worker already made with LOGPROB_ATOL, so the test can run instead. +# Its worker now gates DISTRIBUTION equivalence (top-k JSD + Jaccard swept over k, +# plus an exact argmax check and a slope fit; see +# _skinning_equivalence_worker._compare_logits), reusing the metric and thresholds +# tests/vllm/_generation_equivalence_worker.py already calibrated, so the test can +# run instead. # # Worth having rather than deleting: both vLLM skinning tests below need a GPU, so # this is the only end-to-end check that the composer's weight transfer produces a From e2cc4390062344353531d969dd657efec7d17fe7 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Thu, 22 Oct 2026 10:19:25 +0300 Subject: [PATCH 07/15] Measure skin logit scale by spread ratio, not an OLS slope The slope check added in the previous commit read 3.2e-05 on arm64 but 4.0e-03 on CI's x86 leg, failing its 1e-3 gate on a skin that was otherwise clean: argmax 0/12, JSD at the benign floor, nothing else moved. The platform was not giving the wrong answer -- OLS is the wrong estimator. cov(x,y)/var(x) assumes a noise-free predictor. Here x = t + e_x and y = t + e_y (t the exact logit, e each backend's reduction error), so cov = var(t) while var(x) = var(t) + var(e_x), and the fit converges to 1 - var(e_x)/var(t). The reference is itself bf16, so there is always some e_x, and the bias goes as the SQUARE of the residual -- a platform with 2x the reduction noise reports 4x the scale error for an equally correct skin. Confirmed by injecting noise into both sides of the measured logit pair: extra noise residual OLS |a-1| std ratio |a-1| none 2.21e-02 3.2e-05 2.8e-04 1x 3.13e-02 2.7e-04 2.2e-04 2x 4.94e-02 9.0e-04 3.2e-04 4x 9.11e-02 3.9e-03 2.6e-04 <- the CI reading 6x 1.34e-01 8.7e-03 3.2e-04 Use std(sw)/std(ref) instead: symmetric in the two noises so they cancel rather than bias it, and still exactly `a` under a true uniform scaling. Mean-centred via std, so a benign uniform shift still reads exactly 1. SLOPE_TOL goes to 3e-3, ~10x over the measured 2.8e-04 floor and still ~35x under the 1.03e-02 a 1% scaling produces. Also print the reduction residual, ungated. It is the quantity that explains a scale reading, so a future platform difference is one run to diagnose instead of another round of bisecting CI. Re-verified through the real comparator: real skin, identical, uniform shifts (+1.5, +0.9) and a bf16 round-trip pass; 1%/5% scalings, logits_scaling 8.0->8.1, a per-token shift, swapped positions, a corrupted position and a 100-token weight permutation all fail. Signed-off-by: itzikvaisman --- CLAUDE.md | 16 ++-- .../composer/_skinning_equivalence_worker.py | 93 ++++++++++++++----- 2 files changed, 78 insertions(+), 31 deletions(-) diff --git a/CLAUDE.md b/CLAUDE.md index 8e9b97f9..8d1c8660 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -420,12 +420,16 @@ three live after `log_softmax` and the fourth covers what normalizing hides: the rank-100 boundary is a pile of exact ties (measured 4, 7 and 7 tokens sharing one bf16 value at the cut) and which tie-mate lands inside the cut is arbitrary. Jaccard measures how much the top-k overlaps rather than demanding it match. -- **slope of skinned-vs-reference logits** (`SLOPE_TOL`) — the only non-scale-invariant check, - which is its whole job. A uniform logit *scaling* preserves ranking and barely dents the - normalized shape (a 5% error measures inside the benign JSD floor), yet it is the signature of - a mistransferred `logits_scaling` / `attention_multiplier`. A scale error is systematic where - reduction-order noise is random, so the fit separates them by ~300x. Intercept left free, so a - benign uniform shift still reads as slope 1. +- **relative logit scale** (`SLOPE_TOL`) — the only non-scale-invariant check, which is its + whole job. A uniform logit *scaling* preserves ranking and barely dents the normalized shape + (a 5% error measures inside the benign JSD floor), yet it is the signature of a mistransferred + `logits_scaling` / `attention_multiplier`. Measured as `std(sw)/std(ref)`, **not** an OLS + slope: OLS is `cov(x,y)/var(x)`, and since the reference is itself bf16 the predictor carries + noise, so OLS converges to `1 - var(e_x)/var(t)` — biased down by the *square* of the + reduction residual. That made it platform-dependent (3.2e-05 on arm64, 4.0e-03 on CI's x86, + failing a 1e-3 gate on a correct skin). The ratio is symmetric in the two noises so they + cancel, and it still convicts: 2.8e-04 for the real skin against 1.03e-02 for a 1% scaling. + Mean-centred, so a benign uniform shift still reads exactly 1. Worth keeping rather than deleting because both vLLM skinning tests need a GPU, making this the only CPU-runnable check that the composer's weight transfer reproduces the original model. diff --git a/tests/composer/_skinning_equivalence_worker.py b/tests/composer/_skinning_equivalence_worker.py index 5c895f49..2af4e8bf 100644 --- a/tests/composer/_skinning_equivalence_worker.py +++ b/tests/composer/_skinning_equivalence_worker.py @@ -24,7 +24,7 @@ weight-transfer check that runs without a GPU. What replaced it is the same gate the vLLM generation-equivalence worker uses -(top-k JSD + Jaccard over a swept k), plus an exact argmax check and a slope fit. +(top-k JSD + Jaccard over a swept k), plus an exact argmax check and a scale check. See the constants below for the measured noise floor those are calibrated against, and ``_compare_logits`` for why each of the four is there. @@ -93,8 +93,7 @@ # flipped near-tie only shows up at small k, and a reordered tail only at large k. K_SWEEP = (1, 5, 10, 20, 50, 100) -# Max allowed |a-1| where `a` is the least-squares slope of skinned-vs-reference -# logits. +# Max allowed |a-1| where `a` is the ratio of skinned to reference logit spread. # # This check exists because the three gates above are all computed AFTER # log_softmax, which makes them blind to one real bug: a uniform SCALING of every @@ -105,11 +104,39 @@ # multipliers a skin could plausibly mistransfer, and `logits_scaling` is typically # 8.0 so a one-step error is a ~12% scale change. # -# A slope fit separates it cleanly because a scale error is SYSTEMATIC across all -# 1.2M elements while reduction-order noise is random: measured |a-1| = 3.2e-05 for -# the real skin versus 1.0e-02 for a 1% scaling -- a 300x gap. The intercept is left -# free, so a benign uniform shift still reads as slope 1 exactly. -SLOPE_TOL = 1e-3 +# A scale estimate separates that from reduction noise because a scale error is +# SYSTEMATIC across all 1.2M elements while reduction-order noise is random. +# +# NOT AN ORDINARY LEAST-SQUARES SLOPE, and the reason is a CI failure worth +# recording. This started as OLS -- cov(x,y)/var(x) -- which measured 3.2e-05 on +# the dev box but 4.0e-03 on CI's x86 leg, failing a 1e-3 gate on a skin that was +# otherwise clean (argmax 0/12, JSD at the benign floor). OLS is the wrong +# estimator here, not the platform the wrong answer: with x = t + e_x and +# y = t + e_y (t the exact logit, e each backend's reduction error), +# cov(x,y) = var(t) but var(x) = var(t) + var(e_x), so OLS converges to +# 1 - var(e_x)/var(t). It is biased DOWN by noise in the PREDICTOR -- classic +# regression dilution -- and the reference is itself bf16, so there is always some. +# Worse, the bias goes as the SQUARE of the residual, so a platform with 2x the +# reduction noise reads 4x the error for an equally correct skin. +# +# std(y)/std(x) is used instead: symmetric in the two noises, so they cancel +# rather than bias it, while still reading exactly `a` under a true uniform +# scaling. CONFIRMED by injecting extra noise into both sides of the measured +# logit pair -- OLS crosses the old 1e-3 gate on a correct skin while the ratio +# does not move at all: +# +# extra noise residual OLS |a-1| scale |a-1| +# none 2.21e-02 3.2e-05 2.8e-04 +# 1x 3.13e-02 2.7e-04 2.2e-04 +# 2x 4.94e-02 9.0e-04 3.2e-04 +# 4x 9.11e-02 3.9e-03 2.6e-04 <- the CI reading +# 6x 1.34e-01 8.7e-03 3.2e-04 +# +# Still convicts, which is the only reason the swap is worth making: 1.03e-02 for +# a 1% scaling and 1.21e-02 for logits_scaling 8.0->8.1, against 2.8e-04 for the +# real skin -- a ~40x gap, and one that does not move when the platform's +# arithmetic does. +SLOPE_TOL = 3e-3 def _native_dtype(config): @@ -128,16 +155,30 @@ def _native_dtype(config): return torch.float32 -def _slope(ref, sw): - """Least-squares slope of ``sw`` against ``ref``, intercept free. +def _scale(ref, sw): + """Ratio of ``sw``'s logit spread to ``ref``'s: 1.0 iff unscaled. + + Mean-centred via ``std``, so a uniform additive shift -- invisible after + softmax -- does not register here either. Returns exactly 1.0 for a shifted + copy and exactly ``a`` for an ``a``-scaled one. + + Symmetric in the two sides on purpose: this replaced an OLS slope that was + biased by noise in the reference and so read differently per platform. See + ``SLOPE_TOL``. + """ + return float(sw.reshape(-1).std() / ref.reshape(-1).std()) + + +def _residual(ref, sw): + """``std(sw - ref) / std(ref)`` -- the reduction-noise floor, printed only. - Intercept free on purpose: a uniform additive shift is invisible after softmax, - so it must not register here either. Returns 1.0 exactly for a shifted copy. + Not gated: it is a raw magnitude ratio, so a legitimately different float + reduction order moves it, which is what made the original magnitude bound + near-vacuous. Reported because it is the quantity that *explains* a scale + reading, which makes a future platform difference diagnosable in one run. + ARM measures 2.2e-02. """ - x = ref.reshape(-1) - y = sw.reshape(-1) - xc = x - x.mean() - return float((xc * (y - y.mean())).sum() / (xc * xc).sum()) + return float((sw - ref).reshape(-1).std() / ref.reshape(-1).std()) def _compare_logits(reference, switch, label, dtype): @@ -160,11 +201,12 @@ def _compare_logits(reference, switch, label, dtype): 3. **mean top-k Jaccard distance.** A loose set-overlap guard, same role as in the vLLM worker. Deliberately loose: at these bf16 logits several tokens tie exactly at the k-th rank, so a few percent of churn is arithmetic. - 4. **slope of skinned-vs-reference logits.** The only check here that is NOT - scale-invariant, which is precisely its job -- see ``SLOPE_TOL``. + 4. **relative logit scale.** The only check here that is NOT scale-invariant, + which is precisely its job -- see ``SLOPE_TOL``. All four are printed either way, so a run drifting toward a bound is visible - before it crosses it. + before it crosses it, along with the ungated reduction residual that explains + where the scale estimate's own noise comes from. """ # Compare in fp32: a bf16 subtraction would itself round the difference being # measured, and the two tensors are already materialized at native precision. @@ -175,12 +217,13 @@ def _compare_logits(reference, switch, label, dtype): sw_lp = torch.log_softmax(sw, dim=-1) argmax_mismatches = int((ref.argmax(dim=-1) != sw.argmax(dim=-1)).sum().item()) - slope = _slope(ref, sw) + scale = _scale(ref, sw) num_pos = ref.shape[0] print(f" dtype = {dtype} positions = {num_pos} vocab = {ref.shape[-1]}") print(f" argmax mismatches: {argmax_mismatches} / {num_pos}") - print(f" slope |a-1| = {abs(slope - 1):.3e} tolerance = {SLOPE_TOL:.3e}") + print(f" scale |a-1| = {abs(scale - 1):.3e} tolerance = {SLOPE_TOL:.3e}") + print(f" reduction residual = {_residual(ref, sw):.3e} (informational)") print(f" {'k':>5}{'mean(1-Jacc)':>15}{'mean JSD':>12}{'max JSD':>12}{'@pos':>7}") failures = [] @@ -221,10 +264,10 @@ def _compare_logits(reference, switch, label, dtype): f"top token; the skin is not equivalent (dtype={dtype})" ) return 1 - if abs(slope - 1) > SLOPE_TOL: + if abs(scale - 1) > SLOPE_TOL: print( - f"\nFAIL: {label} — skinned logits are scaled by {slope:.6f} relative to " - f"the original (|a-1| = {abs(slope - 1):.3e} > {SLOPE_TOL:.3e}); suspect a " + f"\nFAIL: {label} — skinned logits are scaled by {scale:.6f} relative to " + f"the original (|a-1| = {abs(scale - 1):.3e} > {SLOPE_TOL:.3e}); suspect a " f"mistransferred logits_scaling / attention_multiplier (dtype={dtype})" ) return 1 @@ -236,7 +279,7 @@ def _compare_logits(reference, switch, label, dtype): print( f"\nPASS: {label} — distribution-equivalent (mean JSD <= {MEAN_JSD_THRESH}, " f"max JSD <= {MAX_JSD_THRESH}, mean 1-Jaccard <= {JACC_THRESH}, " - f"slope within {SLOPE_TOL:.0e}, dtype={dtype})" + f"scale within {SLOPE_TOL:.0e}, dtype={dtype})" ) return 0 From 85e2a6ec5f808df620ab825c16633cf12f06c428 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Thu, 22 Oct 2026 12:09:26 +0300 Subject: [PATCH 08/15] Launch tutorial vLLM servers with the caller's interpreter launch_vllm spawned the API server as `python -m vllm.entrypoints.openai.api_server`, letting PATH decide which interpreter ran it. In a container with both a system Python and a project venv those are different interpreters, and the one first on PATH need not have vLLM installed. That is not hypothetical: five tests/integration serving modules failed this way under `make test-e2e PYTHON=.../.venv/bin/python` on a GPU pod. The venv ran the tests, but the child got /usr/local/bin/python and died on "No module named 'vllm'" before binding its port, so every test in those modules timed out waiting for a server that was already dead. Use sys.executable so the server always runs in the same environment as its caller. Verified on the pod: test_conversation_prefix_cache now reaches SERVE_OK where it previously could not start the server. Signed-off-by: itzikvaisman --- src/granite_switch/tutorials/vllm_server.py | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/src/granite_switch/tutorials/vllm_server.py b/src/granite_switch/tutorials/vllm_server.py index 59b7e10b..b80eb557 100644 --- a/src/granite_switch/tutorials/vllm_server.py +++ b/src/granite_switch/tutorials/vllm_server.py @@ -6,6 +6,7 @@ import os import signal import subprocess +import sys import time from collections.abc import Sequence @@ -24,8 +25,16 @@ def launch_vllm( extra_args: Sequence[str] = (), max_model_len: int = DEFAULT_MAX_MODEL_LEN, ) -> subprocess.Popen: + # sys.executable, not "python": the server has to run in the SAME environment + # as the caller, and the bare name resolves through PATH to whatever comes + # first. On a container with a system Python plus a project venv those are + # different interpreters, and the one on PATH need not have vLLM installed -- + # five tests/integration serving modules failed exactly that way under + # `make test-e2e PYTHON=.../.venv/bin/python`, where the venv ran the tests + # but the child got /usr/local/bin/python and died on + # "No module named 'vllm'" before binding its port. cmd = [ - "python", + sys.executable, "-m", "vllm.entrypoints.openai.api_server", "--model", From fd66895a655b77ebb9106b3e0b484f364ae4e18b Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Thu, 22 Oct 2026 12:11:09 +0300 Subject: [PATCH 09/15] Gate skinned multipliers on the config, not on the logit scale The HF skinning check gated a numeric estimate of the skin's logit scale to catch a mistransferred Granite multiplier, which the three distribution checks cannot see (they run after log_softmax, and a uniform scaling barely dents a normalized shape -- a 5% error measures inside the benign JSD floor). That gate could not be made portable. As an OLS slope it read 3.2e-05 on arm64 and 4.0e-03 on CI's x86, and the previous commit swapped in std(sw)/std(ref) on the theory that OLS was diluted by noise in the predictor. The theory was wrong: the ratio still read 3.7e-03 on x86 against 2.8e-04 on arm64. Both platforms report the same reduction residual (2.318e-02 vs 2.213e-02, within 5%), so noise is not the difference. Independent per-side noise bounds |a-1| by residual^2/2; arm64 sits at 1.1x that ceiling while x86 sits at 14x it, so x86 really does contract these logits ~0.37% on a skin that is otherwise clean (argmax 0/12, JSD at the floor). No estimator fixes a real platform difference, and a gate loose enough for x86 would sit 2.7x under what a 1% scaling reads. Check the config instead, where the bug actually lives and needs no tolerance. compose_utils.py propagates these with getattr(base_config, field, 1.0), so a field upstream renames or drops silently becomes 1.0 -- confirmed for all four, including logits_scaling, which is 10 on granite-4.0-micro. Comparing against the base catches that exactly, on every platform, and covers residual_multiplier and embedding_multiplier too, which no logit-scale reading could isolate. test_save_load_compose only checks these survive a round-trip (loaded == built), never that built matches the base. Scale and residual are still printed, ungated, since they are what made the platform gap diagnosable in one run. Validated by injection: all four silent 1.0 defaults convict and a clean config passes; the logit gates still convict a vocab tilt, a top-100 permutation, a corrupted position and swapped positions, while passing the real skin, a bf16 round-trip and uniform/per-position shifts. Signed-off-by: itzikvaisman --- .github/workflows/gpu-tests.yaml | 14 +- CLAUDE.md | 34 ++-- .../composer/_skinning_equivalence_worker.py | 173 ++++++++++-------- tests/composer/test_skinning_equivalence.py | 10 +- 4 files changed, 136 insertions(+), 95 deletions(-) diff --git a/.github/workflows/gpu-tests.yaml b/.github/workflows/gpu-tests.yaml index ff8924ff..fdfcb287 100644 --- a/.github/workflows/gpu-tests.yaml +++ b/.github/workflows/gpu-tests.yaml @@ -252,11 +252,15 @@ jobs: # real Quanto float8 probe that only # means anything on a GPU. # tests/composer/test_skinning_...py was @pytest.mark.skip outright; now - # gates top-k JSD/Jaccard + argmax + - # a slope fit. This scope is what - # measured the bf16 noise floor those - # are calibrated against -- a CPU box - # runs it in fp32 and sees none of it. + # gates top-k JSD/Jaccard + argmax, + # plus exact config multipliers. This + # scope is what measured the bf16 noise + # floor the thresholds are calibrated + # against -- a CPU box runs it in fp32 + # and sees none of it. It is also what + # caught x86 contracting the logits + # 0.37%, which retired a numeric + # logit-scale gate as unportable. # tests/unit/test_bitmask_computation 34 tests, no longer 2 skips. Cheap, # and it cannot even be COLLECTED off # GPU: it imports lora_kernel_meta, diff --git a/CLAUDE.md b/CLAUDE.md index 8d1c8660..0ac09cbc 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -408,8 +408,7 @@ with the upstream HF model is not achievable. The HF skinning test in `tests/composer/test_skinning_equivalence.py` used to be skipped outright for this reason; it now runs and gates **distribution** equivalence instead, with the same metric and thresholds `tests/vllm/_generation_equivalence_worker.py` already calibrated: top-k JSD + Jaccard swept -over k, plus an exact argmax check and a least-squares slope fit. Four checks because the first -three live after `log_softmax` and the fourth covers what normalizing hides: +over k, plus an exact argmax check. Three logit checks, all after `log_softmax`: - **argmax, exact** — what greedy decoding emits; measured to be a strict maximum at every position, so no tolerance is needed even in bf16. @@ -420,16 +419,27 @@ three live after `log_softmax` and the fourth covers what normalizing hides: the rank-100 boundary is a pile of exact ties (measured 4, 7 and 7 tokens sharing one bf16 value at the cut) and which tie-mate lands inside the cut is arbitrary. Jaccard measures how much the top-k overlaps rather than demanding it match. -- **relative logit scale** (`SLOPE_TOL`) — the only non-scale-invariant check, which is its - whole job. A uniform logit *scaling* preserves ranking and barely dents the normalized shape - (a 5% error measures inside the benign JSD floor), yet it is the signature of a mistransferred - `logits_scaling` / `attention_multiplier`. Measured as `std(sw)/std(ref)`, **not** an OLS - slope: OLS is `cov(x,y)/var(x)`, and since the reference is itself bf16 the predictor carries - noise, so OLS converges to `1 - var(e_x)/var(t)` — biased down by the *square* of the - reduction residual. That made it platform-dependent (3.2e-05 on arm64, 4.0e-03 on CI's x86, - failing a 1e-3 gate on a correct skin). The ratio is symmetric in the two noises so they - cancel, and it still convicts: 2.8e-04 for the real skin against 1.03e-02 for a 1% scaling. - Mean-centred, so a benign uniform shift still reads exactly 1. +Being post-`log_softmax`, those three are blind to a uniform logit *scaling* — the signature of +a mistransferred Granite multiplier. That is gated separately, and **on the config rather than +the logits**: `_compare_multipliers` asserts all four of `logits_scaling`, +`attention_multiplier`, `residual_multiplier` and `embedding_multiplier` match the base config +exactly. The bug is reachable — `compose_utils.py:189` propagates them with +`getattr(base_config, field, default)` where the default is `1.0`, so a field upstream renames +or drops silently becomes 1.0 (verified for all four; absent `logits_scaling`, which is 10 on +granite-4.0-micro, would compose as 1.0). `test_save_load_compose.py` only checks these survive +a save/load round-trip (`loaded == built`), never that `built` matches the *base*, so this is +the only coverage of that gap. + +A numeric logit-scale gate was tried first and **could not be made platform-independent**. As an +OLS slope it read 3.2e-05 on arm64 and 4.0e-03 on CI's x86; switching to `std(sw)/std(ref)` on a +regression-dilution theory did not help (3.7e-03 on x86 vs 2.8e-04 on arm64). The data refutes +the noise explanation: both platforms report the same reduction residual (2.318e-02 vs +2.213e-02, within 5%), and since independent per-side noise bounds `|a-1|` by `residual²/2`, +arm64 sits at 1.1x that ceiling (pure noise) while x86 sits at **14x** it. x86 genuinely +contracts the logits ~0.37% on a skin that is otherwise clean. No estimator fixes a real +platform difference, and a gate loose enough for x86 (~4e-3) would sit only 2.7x below what a 1% +scaling reads. The scale and residual are still *printed* (ungated) because they make a future +platform difference diagnosable in one run. Worth keeping rather than deleting because both vLLM skinning tests need a GPU, making this the only CPU-runnable check that the composer's weight transfer reproduces the original model. diff --git a/tests/composer/_skinning_equivalence_worker.py b/tests/composer/_skinning_equivalence_worker.py index 2af4e8bf..d70a3de8 100644 --- a/tests/composer/_skinning_equivalence_worker.py +++ b/tests/composer/_skinning_equivalence_worker.py @@ -24,9 +24,16 @@ weight-transfer check that runs without a GPU. What replaced it is the same gate the vLLM generation-equivalence worker uses -(top-k JSD + Jaccard over a swept k), plus an exact argmax check and a scale check. -See the constants below for the measured noise floor those are calibrated against, -and ``_compare_logits`` for why each of the four is there. +(top-k JSD + Jaccard over a swept k) plus an exact argmax check -- see the constants +below for the measured noise floor those are calibrated against, and +``_compare_logits`` for why each of the three is there. + +Those three all live after ``log_softmax`` and are therefore blind to a uniform +logit SCALING, the signature of a mistransferred Granite multiplier. That bug is +caught on the config instead, by exact comparison against the base +(``_compare_multipliers``), rather than inferred from the logits -- a numeric +estimate of the scale turned out to differ 13x between arm64 and x86 on the same +correct skin. See ``_MULTIPLIERS``. Runs in its own process so CUDA context is fully torn down on exit. """ @@ -93,50 +100,47 @@ # flipped near-tie only shows up at small k, and a reordered tail only at large k. K_SWEEP = (1, 5, 10, 20, 50, 100) -# Max allowed |a-1| where `a` is the ratio of skinned to reference logit spread. -# -# This check exists because the three gates above are all computed AFTER -# log_softmax, which makes them blind to one real bug: a uniform SCALING of every -# logit. Scaling preserves ranking (so argmax and Jaccard see nothing) and barely -# dents the normalized shape at small factors -- a 5% error measured mean JSD -# 7.2e-04, inside the benign floor. That is not a hypothetical: it is the signature -# of a wrong ``logits_scaling`` or ``attention_multiplier``, the two Granite config -# multipliers a skin could plausibly mistransfer, and `logits_scaling` is typically -# 8.0 so a one-step error is a ~12% scale change. -# -# A scale estimate separates that from reduction noise because a scale error is -# SYSTEMATIC across all 1.2M elements while reduction-order noise is random. +# The four Granite scalar multipliers, checked for EXACT equality between the base +# config and the skin's. # -# NOT AN ORDINARY LEAST-SQUARES SLOPE, and the reason is a CI failure worth -# recording. This started as OLS -- cov(x,y)/var(x) -- which measured 3.2e-05 on -# the dev box but 4.0e-03 on CI's x86 leg, failing a 1e-3 gate on a skin that was -# otherwise clean (argmax 0/12, JSD at the benign floor). OLS is the wrong -# estimator here, not the platform the wrong answer: with x = t + e_x and -# y = t + e_y (t the exact logit, e each backend's reduction error), -# cov(x,y) = var(t) but var(x) = var(t) + var(e_x), so OLS converges to -# 1 - var(e_x)/var(t). It is biased DOWN by noise in the PREDICTOR -- classic -# regression dilution -- and the reference is itself bf16, so there is always some. -# Worse, the bias goes as the SQUARE of the residual, so a platform with 2x the -# reduction noise reads 4x the error for an equally correct skin. +# WHY A CONFIG CHECK AND NOT A NUMERIC ONE. The three distribution gates above are +# all computed after ``log_softmax``, which makes them blind to one real bug: a +# uniform SCALING of every logit. Scaling preserves ranking (argmax and Jaccard see +# nothing) and barely dents the normalized shape -- a 5% error measures mean JSD +# 7.2e-04, inside the benign floor. That is not hypothetical: it is the signature of +# a mistransferred multiplier, and the transfer really can drop one silently. +# ``compose_utils.py`` propagates these with ``getattr(base_config, field, default)`` +# where the default is ``1.0``, so a field that upstream renames or removes does not +# raise -- it quietly becomes 1.0. Verified for all four: absent ``logits_scaling`` +# (10 on granite-4.0-micro) would compose as 1.0. # -# std(y)/std(x) is used instead: symmetric in the two noises, so they cancel -# rather than bias it, while still reading exactly `a` under a true uniform -# scaling. CONFIRMED by injecting extra noise into both sides of the measured -# logit pair -- OLS crosses the old 1e-3 gate on a correct skin while the ratio -# does not move at all: +# So the bug is reachable, but it is a CONFIG bug, and comparing the configs catches +# it exactly, on every platform, for all four multipliers rather than just the two +# that happen to move the logit scale. ``test_save_load_compose.py`` only checks +# these survive a save/load round-trip (``loaded == built``), never that ``built`` +# matches the BASE, so this is the only place that gap is covered. # -# extra noise residual OLS |a-1| scale |a-1| -# none 2.21e-02 3.2e-05 2.8e-04 -# 1x 3.13e-02 2.7e-04 2.2e-04 -# 2x 4.94e-02 9.0e-04 3.2e-04 -# 4x 9.11e-02 3.9e-03 2.6e-04 <- the CI reading -# 6x 1.34e-01 8.7e-03 3.2e-04 +# WHAT THIS REPLACED. A gated numeric estimate of the logit scale, which could not be +# made platform-independent. It began as an OLS slope (1e-3 gate), read 3.2e-05 on +# arm64 and 4.0e-03 on CI's x86 leg, and was changed to ``std(sw)/std(ref)`` on the +# theory that OLS was diluted by noise in the predictor. THE DATA REFUTED THAT: the +# ratio still read 3.74e-03 on x86 against 2.77e-04 on arm64. The two platforms +# report the SAME reduction residual (2.318e-02 vs 2.213e-02, within 5%), so noise is +# not the difference. Independent per-side noise bounds |a-1| at r^2/2, and arm64 sits +# at 1.13x that ceiling -- pure noise -- while x86 sits at 13.9x it. x86 genuinely +# contracts the logits ~0.37%, systematically, on a skin that is otherwise clean +# (argmax 0/12, JSD at the benign floor). No choice of estimator fixes a real +# platform difference, and a gate loose enough to pass x86 (~4e-3) would sit only +# 2.7x under what a 1% scaling reads (1.03e-02) -- too thin to trust either way. # -# Still convicts, which is the only reason the swap is worth making: 1.03e-02 for -# a 1% scaling and 1.21e-02 for logits_scaling 8.0->8.1, against 2.8e-04 for the -# real skin -- a ~40x gap, and one that does not move when the platform's -# arithmetic does. -SLOPE_TOL = 3e-3 +# ``_scale`` is still computed and printed, just not gated: it is the quantity that +# makes a future platform difference diagnosable in one run. +_MULTIPLIERS = ( + "logits_scaling", + "attention_multiplier", + "residual_multiplier", + "embedding_multiplier", +) def _native_dtype(config): @@ -159,12 +163,11 @@ def _scale(ref, sw): """Ratio of ``sw``'s logit spread to ``ref``'s: 1.0 iff unscaled. Mean-centred via ``std``, so a uniform additive shift -- invisible after - softmax -- does not register here either. Returns exactly 1.0 for a shifted - copy and exactly ``a`` for an ``a``-scaled one. + softmax -- does not register here either. - Symmetric in the two sides on purpose: this replaced an OLS slope that was - biased by noise in the reference and so read differently per platform. See - ``SLOPE_TOL``. + PRINTED, NOT GATED: it is platform-dependent (2.8e-04 on arm64, 3.7e-03 on + x86 for the same correct skin) because x86 really does contract these logits. + The bug it was meant to catch is gated exactly via ``_MULTIPLIERS`` instead. """ return float(sw.reshape(-1).std() / ref.reshape(-1).std()) @@ -172,21 +175,36 @@ def _scale(ref, sw): def _residual(ref, sw): """``std(sw - ref) / std(ref)`` -- the reduction-noise floor, printed only. - Not gated: it is a raw magnitude ratio, so a legitimately different float - reduction order moves it, which is what made the original magnitude bound - near-vacuous. Reported because it is the quantity that *explains* a scale - reading, which makes a future platform difference diagnosable in one run. - ARM measures 2.2e-02. + Reported next to ``_scale`` because it bounds it: with independent per-side + noise, ``|scale - 1| <= residual^2 / 2``. A scale reading above that ceiling + is systematic rather than noise, which is how the x86/arm64 gap was diagnosed + (both report residual ~2.2e-02; arm64 lands at 1.1x the ceiling, x86 at 14x). """ return float((sw - ref).reshape(-1).std() / ref.reshape(-1).std()) +def _compare_multipliers(base_config, switch_config): + """Check the Granite scalar multipliers transferred exactly. Returns a failure list. + + Exact equality, no tolerance: these are config floats copied verbatim, not + computed, so any difference is a transfer defect. See ``_MULTIPLIERS``. + """ + failures = [] + for name in _MULTIPLIERS: + want = getattr(base_config, name, None) + got = getattr(switch_config, name, None) + status = "OK" if want == got else "MISMATCH" + print(f" {name:24} base={want!r:12} switch={got!r:12} {status}") + if want != got: + failures.append(f"{name}: base={want!r} but skin={got!r}") + return failures + + def _compare_logits(reference, switch, label, dtype): """Check skinned logits against the original's. Returns 0 on agreement, 1 on diff. - Four checks. The first three are computed on ``log_softmax`` output -- the - distribution a caller actually observes -- and the fourth covers the one real - bug that normalizing hides. They fail for different reasons, and reporting them + Three gated checks, all computed on ``log_softmax`` output -- the distribution a + caller actually observes. They fail for different reasons, and reporting them separately is what makes a failure diagnosable rather than just "the numbers moved": @@ -201,12 +219,13 @@ def _compare_logits(reference, switch, label, dtype): 3. **mean top-k Jaccard distance.** A loose set-overlap guard, same role as in the vLLM worker. Deliberately loose: at these bf16 logits several tokens tie exactly at the k-th rank, so a few percent of churn is arithmetic. - 4. **relative logit scale.** The only check here that is NOT scale-invariant, - which is precisely its job -- see ``SLOPE_TOL``. - All four are printed either way, so a run drifting toward a bound is visible - before it crosses it, along with the ungated reduction residual that explains - where the scale estimate's own noise comes from. + All three are printed either way, so a run drifting toward a bound is visible + before it crosses it. The scale and reduction residual are printed alongside but + NOT gated -- being post-``log_softmax``, the three above are blind to a uniform + logit scaling, but that bug is caught exactly on the config instead + (``_compare_multipliers``) rather than inferred from platform-dependent + arithmetic. See ``_MULTIPLIERS`` for the full reasoning. """ # Compare in fp32: a bf16 subtraction would itself round the difference being # measured, and the two tensors are already materialized at native precision. @@ -222,8 +241,12 @@ def _compare_logits(reference, switch, label, dtype): print(f" dtype = {dtype} positions = {num_pos} vocab = {ref.shape[-1]}") print(f" argmax mismatches: {argmax_mismatches} / {num_pos}") - print(f" scale |a-1| = {abs(scale - 1):.3e} tolerance = {SLOPE_TOL:.3e}") - print(f" reduction residual = {_residual(ref, sw):.3e} (informational)") + resid = _residual(ref, sw) + print( + f" scale |a-1| = {abs(scale - 1):.3e} noise ceiling = {resid**2 / 2:.3e}" + f" (informational, not gated)" + ) + print(f" reduction residual = {resid:.3e} (informational)") print(f" {'k':>5}{'mean(1-Jacc)':>15}{'mean JSD':>12}{'max JSD':>12}{'@pos':>7}") failures = [] @@ -264,13 +287,6 @@ def _compare_logits(reference, switch, label, dtype): f"top token; the skin is not equivalent (dtype={dtype})" ) return 1 - if abs(scale - 1) > SLOPE_TOL: - print( - f"\nFAIL: {label} — skinned logits are scaled by {scale:.6f} relative to " - f"the original (|a-1| = {abs(scale - 1):.3e} > {SLOPE_TOL:.3e}); suspect a " - f"mistransferred logits_scaling / attention_multiplier (dtype={dtype})" - ) - return 1 if failures: print(f"\nFAIL: {label} — distributions disagree beyond the noise floor:") for line in failures: @@ -278,8 +294,7 @@ def _compare_logits(reference, switch, label, dtype): return 1 print( f"\nPASS: {label} — distribution-equivalent (mean JSD <= {MEAN_JSD_THRESH}, " - f"max JSD <= {MAX_JSD_THRESH}, mean 1-Jaccard <= {JACC_THRESH}, " - f"scale within {SLOPE_TOL:.0e}, dtype={dtype})" + f"max JSD <= {MAX_JSD_THRESH}, mean 1-Jaccard <= {JACC_THRESH}, dtype={dtype})" ) return 0 @@ -331,8 +346,18 @@ def main(): print(f" switch_logits shape: {switch_logits.shape}") # ── Phase 3: compare ────────────────────────────────────────── - print("\nPhase 3: comparing logits...") - return _compare_logits(reference_logits, switch_logits, model_name, dtype) + print("\nPhase 3: comparing Granite scalar multipliers...") + mult_failures = _compare_multipliers(base_config, switch_model.config) + + print("\nPhase 4: comparing logits...") + rc = _compare_logits(reference_logits, switch_logits, model_name, dtype) + + if mult_failures: + print(f"\nFAIL: {model_name} — Granite multipliers did not transfer:") + for line in mult_failures: + print(f" {line}") + return 1 + return rc if __name__ == "__main__": diff --git a/tests/composer/test_skinning_equivalence.py b/tests/composer/test_skinning_equivalence.py index bb0adb0b..37429a86 100644 --- a/tests/composer/test_skinning_equivalence.py +++ b/tests/composer/test_skinning_equivalence.py @@ -185,10 +185,12 @@ def _run_vllm_test(model_name, timeout, fast=False): # unconditional @pytest.mark.skip, on the correct observation that bit-exact # equivalence is unreachable -- but the conclusion drawn from it was too strong. # Its worker now gates DISTRIBUTION equivalence (top-k JSD + Jaccard swept over k, -# plus an exact argmax check and a slope fit; see -# _skinning_equivalence_worker._compare_logits), reusing the metric and thresholds -# tests/vllm/_generation_equivalence_worker.py already calibrated, so the test can -# run instead. +# plus an exact argmax check; see _skinning_equivalence_worker._compare_logits), +# reusing the metric and thresholds tests/vllm/_generation_equivalence_worker.py +# already calibrated, so the test can run instead. Those three checks live after +# log_softmax and so cannot see a uniform logit scaling; that bug is gated exactly +# on the config instead (_compare_multipliers), because a numeric estimate of the +# scale differs 13x between arm64 and x86 on the same correct skin. # # Worth having rather than deleting: both vLLM skinning tests below need a GPU, so # this is the only end-to-end check that the composer's weight transfer produces a From 5b8194565e3d59763b02aa04178dd0a97d002c42 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Thu, 22 Oct 2026 13:36:40 +0300 Subject: [PATCH 10/15] Parse the adapter map by structure, not by a fixed key sequence _ADAPTER_MAP_RE spelled out the exact key order of an adapter_map entry and required the closing brace right after 'invocation_text'. The composer's aLoRA branch now emits a fourth key as well (invocation_drop_chars, tokenizer_setup.py:986-992), so every aLoRA entry stopped matching while LoRA entries -- still three keys -- kept parsing. The failure that produced was badly misleading. On an all-aLoRA checkpoint the map came back EMPTY, and both workers that share this parser reported "need >=2 alora adapters, got []" against a sound 3-adapter model, pointing at the adapters instead of at the reader. Match an entry by its braces and read keys out of the body instead, so the map stays the composer's to extend. Verified against all four entry shapes the composer emits (4-key aLoRA, 3-key LoRA, SR, and the legacy 3-key aLoRA form), and it ignores unrelated Jinja dicts. Confirmed on a GPU pod: test_conversation_prefix_cache now resolves all three aLoRA adapters and passes, measuring the routing the empty map had been hiding (PRESERVE reuses 80 cached positions against RE_PREFILL's 64, with distinct decode write addresses). Signed-off-by: itzikvaisman --- .../_conversation_prefix_cache_worker.py | 41 +++++++++++++++---- 1 file changed, 33 insertions(+), 8 deletions(-) diff --git a/tests/integration/_conversation_prefix_cache_worker.py b/tests/integration/_conversation_prefix_cache_worker.py index ba10ffa5..339fddfa 100644 --- a/tests/integration/_conversation_prefix_cache_worker.py +++ b/tests/integration/_conversation_prefix_cache_worker.py @@ -45,10 +45,21 @@ DEFAULT_PORT = 8137 MAX_TOKENS = 24 -_ADAPTER_MAP_RE = re.compile( - r"'(?P[^']+)':\s*\{'token':\s*'(?P[^']+)',\s*" - r"'type':\s*'(?P[^']+)'(?:,\s*'invocation_text':\s*'(?P[^']*)')?\}" -) +# Match an entry by its BRACES, then pull keys out of the body, rather than +# pinning the exact key sequence. +# +# The previous pattern spelled out `{'token': .., 'type': ..[, 'invocation_text': +# ..]}` and required the `}` right after. That silently stopped matching every +# aLoRA entry once the composer began emitting a fourth key +# ('invocation_drop_chars', tokenizer_setup.py:986-992): LoRA entries still parsed +# (3 keys, unchanged) but aLoRA ones did not, so an all-aLoRA checkpoint produced +# an EMPTY map and the workers reported "need >=2 alora adapters, got []" against a +# perfectly good 3-adapter model. Nothing was wrong with the checkpoint, and the +# error pointed at the adapters rather than at the parser. +# +# `[^{}]*` for the body is safe because these entries never nest. +_ADAPTER_ENTRY_RE = re.compile(r"'(?P[^']+)':\s*\{(?P[^{}]*)\}") +_ADAPTER_KEY_RE = re.compile(r"'(?P\w+)':\s*'(?P[^']*)'") def _parse_adapter_map(chat_template): @@ -57,11 +68,25 @@ def _parse_adapter_map(chat_template): The composer bakes this map into the template, so it is the checkpoint's own account of where each adapter's control token belongs -- no second source of truth to drift. + + Tolerant of extra keys by design: the map is the composer's to extend, and a + reader that must be edited in lockstep turns a new key into a confusing + failure somewhere else. ``invocation_text`` is None when absent (LoRA, SR). """ - return { - m.group("name"): (m.group("token"), m.group("type"), m.group("inv")) - for m in _ADAPTER_MAP_RE.finditer(chat_template) - } + out = {} + for m in _ADAPTER_ENTRY_RE.finditer(chat_template): + keys = { + k.group("key"): k.group("val") + for k in _ADAPTER_KEY_RE.finditer(m.group("body")) + } + if "token" not in keys or "type" not in keys: + continue # not an adapter entry + out[m.group("name")] = ( + keys["token"], + keys["type"], + keys.get("invocation_text"), + ) + return out def cmd_build(args): From cd7102100d34cc1deabd658b8db75f6b9d356fca Mon Sep 17 00:00:00 2001 From: aviv ron Date: Thu, 22 Oct 2026 16:50:33 +0300 Subject: [PATCH 11/15] Cap vLLM memory in the skinning equivalence worker The skinning vLLM tests failed on every `skips` run at engine init: vLLM's default gpu_memory_utilization (0.92 in newer releases) is a fraction of total device memory, and earlier in-process HF tests leave ~9 GiB held by the parent pytest, so only 70.19/79.25 GiB is free against the 72.91 GiB requested. Pass gpu_memory_utilization=0.8 and max_model_len=2048 (inputs top out at 256 tokens; this stops vLLM sizing KV cache for the full 131k context). The original and the skin both run through cmd_run, so they keep identical settings. Signed-off-by: aviv ron --- tests/composer/_skinning_equivalence_worker_vllm.py | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/tests/composer/_skinning_equivalence_worker_vllm.py b/tests/composer/_skinning_equivalence_worker_vllm.py index b6f3df72..ad5550c8 100644 --- a/tests/composer/_skinning_equivalence_worker_vllm.py +++ b/tests/composer/_skinning_equivalence_worker_vllm.py @@ -136,7 +136,12 @@ def cmd_run(args): dtype_s = _dtype_str(dtype) print(f" native_dtype={dtype}") - # Create vLLM instance + # Create vLLM instance. Both the original and the skin go through this path, so + # they share these memory settings and the comparison stays like-for-like. + # gpu_memory_utilization: vLLM's default (0.92 in newer releases) is a fraction of + # TOTAL memory and exceeds what is free once earlier in-process HF tests in the + # parent pytest hold ~9 GiB. max_model_len: the longest input is 256 tokens, so + # don't let vLLM size the KV cache for the model's full 131k context. print(f"Creating vLLM LLM for {model_path}...") llm_kwargs = dict( model=model_path, @@ -145,6 +150,8 @@ def cmd_run(args): enforce_eager=True, enable_prefix_caching=False, skip_tokenizer_init=True, + gpu_memory_utilization=0.8, + max_model_len=2048, ) llm = LLM(**llm_kwargs) From 4e4ff82083966f80a73813dba3e72b0e7247b8c2 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Sun, 25 Oct 2026 09:34:31 +0200 Subject: [PATCH 12/15] Strip trailing whitespace from the SWITCH kernels write-up Four lines in docs/SWITCH_KERNELS.md carry trailing whitespace, which the `trailing-whitespace` pre-commit hook strips. The file arrived in 7e72456 (#148) and is already on main, so the hook fails for every branch cut from current main rather than for anything the branch changed -- this PR included, where it also held back the CPU test job that gates on pre-commit passing. Whitespace only: `git diff -w` on the file is empty, so no prose is altered and the rendered Markdown is unchanged. Signed-off-by: itzikvaisman --- docs/SWITCH_KERNELS.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/docs/SWITCH_KERNELS.md b/docs/SWITCH_KERNELS.md index 7843034b..d7092553 100644 --- a/docs/SWITCH_KERNELS.md +++ b/docs/SWITCH_KERNELS.md @@ -10,8 +10,8 @@ rank-32 ones, and **up to 5.9× faster** with one rank-512 adapter. At idle (no ## Motivation and Availability As part of our mission to create a modular model architecture where adapter libraries are used to enhance -a model, we encountered the problem that existing infrastructure for serving adapters introduced -unnecessary performance overheads. These overheads are noticeable even when a single adapter, regardless of +a model, we encountered the problem that existing infrastructure for serving adapters introduced +unnecessary performance overheads. These overheads are noticeable even when a single adapter, regardless of its size, is introduced, but become very pronounced when the adapters have diverse sizes, which we anticipate will be a common scenario when a developer is integrating library components from different contributors into a single checkpoint. @@ -36,7 +36,7 @@ expensive in our setting. ## What SWITCH does differently SWITCH is named after its design: **S**witched delta, **W**-ext fusion, **I**n-place, -**T**iled, **C**onditional skip, **H**eterogeneous rank. +**T**iled, **C**onditional skip, **H**eterogeneous rank. Existing Punica technology, as implemented in vLLM, exhibits the following properties: @@ -108,7 +108,7 @@ adapter. The figure below compares each system **as it currently implemented**: ![Idle vs adapter-active throughput at 12 adapters loaded, against the base model with LoRA off](figures/fig3_dispatch_3b_b32.png) -Of note is that the overhead at idle that the SWITCH implementation in Granite Switch and vLLM exhibit is unrelated to the SWITCH kernels themselves, and rather +Of note is that the overhead at idle that the SWITCH implementation in Granite Switch and vLLM exhibit is unrelated to the SWITCH kernels themselves, and rather related to the fact that SWITCH is operating as a single computational graph, whereas vLLM does not run an adapter specific graph when no adapters are present. It would be relatively straightforward to implement SWITCH with two graphs, in which case the performance in the no adapter case would match vLLM/Punica exactly. As indicated earlier, we are currently focused in the setting where the model discovers which adapters need to be applied to which tokens through its switch layer, and currently there is no support in vLLM for adapting the computational graph once the model is already invoked. From 57645a49a1988c7455c674e8ffe628a4b2d7e877 Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Mon, 26 Oct 2026 10:26:17 +0200 Subject: [PATCH 13/15] Compare ASR transcripts on content rather than exact spelling test_asr_ctc_default_gpu.py asserted against a single literal _EXPECTED string, so any change in how the CTC head punctuates or capitalises the same words failed the test. Four of its six tests failed that way on an A100 while the transcript was, as text, correct. What the test exists to prove is that the default CTC path transcribes the clip and that the transcript reaches the request -- not that the decoder reproduces one byte sequence. Normalise both sides (casefold, drop punctuation, collapse whitespace) before comparing, so a wrong transcription still fails and a re-punctuated correct one does not. Signed-off-by: itzikvaisman --- tests/integration/test_asr_ctc_default_gpu.py | 31 ++++++++++++++++--- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/tests/integration/test_asr_ctc_default_gpu.py b/tests/integration/test_asr_ctc_default_gpu.py index e87a828a..9eb2a4f1 100644 --- a/tests/integration/test_asr_ctc_default_gpu.py +++ b/tests/integration/test_asr_ctc_default_gpu.py @@ -39,10 +39,32 @@ asr = importlib.util.module_from_spec(_spec) _spec.loader.exec_module(asr) -# tests/audio/test1.wav says this, and CTC emits it lowercase and unpunctuated. +# tests/audio/test1.wav says this. +# +# Compared through _normalize, NOT with ==. The old assertions were literal +# against a lowercase, unpunctuated string on the theory that "CTC emits it +# lowercase and unpunctuated". That is a property of a particular decoder +# revision, not of the audio: measured on the Vela A100 image 2026-10-25, this +# same checkpoint emits "What is the capital of Israel?" -- so all four of this +# module's transcript assertions failed on casing and a trailing "?" while the +# transcription itself was perfect. +# +# What these tests are for is the dtype/backend/chunking plumbing around the +# decoder (bfloat16 default, float16 override, window-vs-chunk call counts), and +# none of that cares how the decoder capitalizes. Gating them on orthography made +# them fail for a reason unrelated to what they cover, so the comparison is +# normalized instead of the expectation being re-pinned to today's spelling -- +# which would just move the same breakage to the next decoder revision. _EXPECTED = "what is the capital of israel" +def _normalize(text): + """Casefold and strip punctuation/extra whitespace for transcript compares.""" + return " ".join( + "".join(c for c in text if c.isalnum() or c.isspace()).casefold().split() + ) + + def _test_clip(): """The committed test waveform as float32 mono, plus its sample rate.""" with wave.open(str(_REPO_ROOT / "tests/audio/test1.wav")) as handle: @@ -79,7 +101,7 @@ def test_transcribes_correctly_and_ignores_decode_kwargs(transcriber): text = transcriber.transcribe( samples, sampling_rate=sample_rate, generate_kwargs={"language": "fr"} ) - assert text == _EXPECTED + assert _normalize(text) == _normalize(_EXPECTED), f"got {text!r}" def test_float16_override_still_transcribes(): @@ -95,7 +117,8 @@ def test_float16_override_still_transcribes(): half = asr.ASRTranscriber( model_id=asr.DEFAULT_ASR_MODEL_ID, device="cuda", dtype="float16" ) - assert half.transcribe(samples, sampling_rate=sample_rate) == _EXPECTED + text = half.transcribe(samples, sampling_rate=sample_rate) + assert _normalize(text) == _normalize(_EXPECTED), f"got {text!r}" def _tile_to_seconds(samples, sample_rate, seconds): @@ -137,4 +160,4 @@ def counting(segment, generate_kwargs=None): else: assert len(calls) == expected_calls # The phrase repeats throughout, so the merge must not collapse it to nothing. - assert _EXPECTED.split()[-1] in text + assert _normalize(_EXPECTED).split()[-1] in _normalize(text), f"got {text!r}" From 57de404df3a5548e80777ee715fe9500295c13af Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Mon, 26 Oct 2026 10:30:51 +0200 Subject: [PATCH 14/15] Stop gating greedy equivalence on prompts that sit on bf16 ties Three GPU failures shared one root cause, and none of them was the bug the failing test's own message named: a prompt of random filler token ids, compared token-exactly under greedy decoding. Measured on an A100 (vLLM 0.26.0, bf16): _pp_equivalence_worker.py built 16 ids from torch.randint(1, 1000). PP=1 and PP=2 agreed for four tokens, then the PP=1 arm collapsed into a repeating two-token loop. The message blamed the PP=2 per-rank metadata recompute, but PP=2 was byte-identical to an earlier green run; PP=1 was the arm that moved. A 2x2 probe cleared async_scheduling (no effect either way) and implicated CUDA-graph capture mode, which vLLM downgrades above PP=1. A negative control then cleared capture too: at PP=1 on real-language prompts, cudagraph_mode=PIECEWISE is byte-identical to enforce_eager with routing intact -- three adapters, three distinct continuations. The prompt was the whole problem. test_batched_generation_matches_solo diverged on the third request at step 8, where the solo run's top two logprobs were the same bf16 value (-1.07067, -1.07067; gap exactly 0.000000). A coin flip, not mis-routing. Out-of-distribution prompts leave the argmax exactly tied, so anything that legitimately reorders a float reduction -- capture mode, co-batching, a version bump -- lands the tie on the other side, and greedy amplifies the one flipped token across every remaining step. A sub-ULP difference presents as total divergence. Fix both halves, because neither alone suffices. Real-language prompts (appended control token rather than spliced, which would break the text back out of distribution) made PP=1 and PP=2 agree on all 32 tokens -- but PP=1 still reported min_top2_gap=0.000000, so no prompt guarantees a separated argmax everywhere. So the tolerance is measured rather than assumed: request per-step logprobs, and fail a divergence only where the top-2 gap is strictly above zero. That still catches what the tests exist for. A rank that rebuilt its bitmasks wrongly, or a request anchored to the wrong place in the flat batch, applies a different adapter delta; that moves the whole distribution and separates the argmax, so it fails with a nonzero gap. Only coin flips pass. This is the reasoning CLAUDE.md section 9 already applies to the skinning and generation-equivalence tests. Signed-off-by: itzikvaisman --- .../test_multi_switch_vllm_generate.py | 77 +++++++++-- tests/vllm/_pp_equivalence_worker.py | 127 +++++++++++++++--- 2 files changed, 174 insertions(+), 30 deletions(-) diff --git a/tests/integration/test_multi_switch_vllm_generate.py b/tests/integration/test_multi_switch_vllm_generate.py index ced82b4b..31de9eee 100644 --- a/tests/integration/test_multi_switch_vllm_generate.py +++ b/tests/integration/test_multi_switch_vllm_generate.py @@ -159,26 +159,87 @@ def test_generation_is_deterministic(engine, control_ids, tokenizer): ) +def _gen_with_logprobs(llm, seqs): + """Generate greedily, returning (token_ids, per-step sorted logprob list).""" + from vllm import SamplingParams + from vllm.inputs import TokensPrompt + + sp = SamplingParams( + temperature=0.0, max_tokens=MAX_NEW_TOKENS, ignore_eos=True, logprobs=5 + ) + prompts = [TokensPrompt(prompt_token_ids=list(s)) for s in seqs] + out = [] + for o in llm.generate(prompts, sampling_params=sp): + comp = o.outputs[0] + steps = [ + sorted((lp.logprob for lp in d.values()), reverse=True) + for d in (comp.logprobs or []) + ] + out.append(([int(t) for t in comp.token_ids], steps)) + return out + + def test_batched_generation_matches_solo(engine, control_ids, tokenizer): - """Each request must generate the same tokens alone as when co-batched. + """Each request must route the same alone as when co-batched. This is the observable consequence of correct per-request counting anchors. With a global anchor, a request's routing depends on what else is in the flat batch, and its output drifts. Prompt lengths are distinct so requests are not accidentally interchangeable. + + Tie-tolerant rather than token-exact, and the tolerance is measured, not + assumed. A token-exact assertion here is not a statement about routing: vLLM + reduces a co-batched request's logits in a different order than a solo one, so + a position where the top two candidates are already equal in bf16 can flip on + sub-ULP noise. Measured on an A100 (vLLM 0.26, bf16) the third request does + exactly that at step 8 -- solo's top-2 logprobs are the *same bf16 value* + (-1.07067, -1.07067; gap 0.000000), so which one greedy picks is arbitrary and + both answers are equally correct. + + Note the earlier attempt recorded above ``_QUESTIONS``: real-language prompts + replaced junk filler to get off near-ties, and a routing probe confirmed + routing was already identical in both runs. That reduced the tie rate but + could not remove it, because no prompt guarantees a strictly-separated argmax + at every position. So the tie is excluded by *measurement* instead: a + divergence is accepted only where the top-2 gap is zero. + + What that still catches is the bug the test exists for. A request whose anchor + came from the wrong place in the flat batch retrieves a different adapter and + computes a genuinely different distribution, so its divergence sits at a + separated argmax (nonzero gap) -- which fails here. Only coin-flips pass. """ a, b, c = control_ids seqs = [_prompt(tokenizer, q, t) for q, t in zip(_QUESTIONS, (a, b, c))] - solo = [_gen(engine, [s])[0] for s in seqs] - batched = _gen(engine, seqs) + solo = [_gen_with_logprobs(engine, [s])[0] for s in seqs] + batched = _gen_with_logprobs(engine, seqs) - mismatches = [i for i in range(len(seqs)) if solo[i] != batched[i]] - assert not mismatches, ( - f"batched generation diverged from solo for request(s) {mismatches}:\n" + real = [] + for i, ((s_toks, s_steps), (b_toks, _)) in enumerate(zip(solo, batched)): + div = next( + (j for j in range(min(len(s_toks), len(b_toks))) if s_toks[j] != b_toks[j]), + None, + ) + if div is None: + continue + # Gap between the best and second-best candidate in the SOLO run at the + # position that diverged. Exactly 0.0 means the two are indistinguishable + # in bf16, so the pick is a coin flip; anything above that is a real + # difference in the distribution and therefore a routing suspect. + top = s_steps[div] if div < len(s_steps) else [] + gap = (top[0] - top[1]) if len(top) > 1 else float("inf") + print(f"req {i}: diverged at step {div}, solo top-2 gap {gap:.6f}") + if gap > 0.0: + real.append((i, div, gap, s_toks, b_toks)) + + assert not real, ( + "batched generation diverged from solo at a position where the argmax was " + "strictly separated, so it is not a tie-break -- suspect per-request " + "counting anchors:\n" + "\n".join( - f" req {i}: solo={solo[i]}\n batched={batched[i]}" - for i in mismatches + f" req {i}: step {div}, solo top-2 gap {gap:.6f}\n" + f" solo={s}\n batched={b}" + for i, div, gap, s, b in real ) ) diff --git a/tests/vllm/_pp_equivalence_worker.py b/tests/vllm/_pp_equivalence_worker.py index be0d8fef..01bac058 100644 --- a/tests/vllm/_pp_equivalence_worker.py +++ b/tests/vllm/_pp_equivalence_worker.py @@ -41,24 +41,62 @@ MAX_NEW_TOKENS = 32 GENERATE_TIMEOUT_SECONDS = 300 +# Prompt text for the real-language path (see _build_prompt_ids). +PROMPT_TEXT = "Explain in one sentence what a hash function is." + +# Per-step top-k logprobs to request, so compare can tell a tie from a real +# difference. 5 is plenty: only the top two matter. +NUM_LOGPROBS = 5 + def _log(message: str) -> None: print(f"PP_EQUIVALENCE_PHASE {message}", flush=True) -def _build_prompt_ids(adapter_token_ids, vocab_size): - """Deterministic prompt of benign low IDs with one control token early. - - The control token at CONTROL_POS makes the switch compute a nonzero adapter - index, so the SWITCH kernel's delta path runs on every PP rank (the whole - point of the test). Fill tokens stay well below any control-token range. +def _build_prompt_ids(model_path, adapter_token_ids, vocab_size): + """Real-language prompt with one control token, so the adapter fires. + + The control token makes the switch compute a nonzero adapter index, so the + SWITCH kernel's delta path runs on every PP rank -- the whole point of the + test. + + Real text, NOT random filler ids. This used to build 16 ids from + ``torch.randint(1, 1000)``, which is badly out of distribution, and under + greedy decoding that compounds: measured on an A100 (vLLM 0.26, bf16) the + PP=1 and PP=2 arms agreed for 4 tokens and then the PP=1 arm collapsed into a + repeating two-token loop, failing the test. With this prompt the two arms are + byte-identical for all 32 tokens. + + The divergence was never the per-rank metadata recompute this test exists to + check. PP=1 and PP=2 differ in CUDA-graph capture mode (vLLM downgrades + capture above PP=1) and in reduction order, which is enough to land an + already-tied argmax on the other side. Probed directly: at PP=1, + ``cudagraph_mode=PIECEWISE`` is byte-identical to ``enforce_eager`` on + real-language prompts with adapter routing intact (3 adapters, 3 distinct + continuations) -- so graph capture does not corrupt the switch, and the junk + prompt was the whole problem. + + Falling back to filler ids if the tokenizer cannot be loaded keeps the worker + runnable on a checkpoint without one; the comparison is then tie-tolerant + (see ``cmd_compare``) rather than token-exact. """ import torch + try: + from transformers import AutoTokenizer + + tokenizer = AutoTokenizer.from_pretrained(model_path) + prompt_ids = list(tokenizer(PROMPT_TEXT)["input_ids"]) + # Appended, not injected mid-sequence: a control token spliced into the + # middle of real text breaks it back out of distribution, which is the + # problem this is fixing. + return [*prompt_ids, int(adapter_token_ids[0])] + except Exception as exc: # pragma: no cover - no-tokenizer checkpoints + _log(f"tokenizer_unavailable ({exc}); falling back to filler ids") + torch.manual_seed(42) max_fill = min(vocab_size, 1000) prompt_ids = torch.randint(1, max_fill, (PROMPT_LEN,)).tolist() - # Fire the first adapter. prompt_ids[CONTROL_POS] = int(adapter_token_ids[0]) return prompt_ids @@ -95,8 +133,8 @@ def cmd_run(args): ) _log(f"config_load_done vocab={vocab_size} adapters={len(adapter_token_ids)}") - prompt_ids = _build_prompt_ids(adapter_token_ids, vocab_size) - _log(f"prompt_ready len={len(prompt_ids)} ctrl_pos={CONTROL_POS}") + prompt_ids = _build_prompt_ids(model_path, adapter_token_ids, vocab_size) + _log(f"prompt_ready len={len(prompt_ids)}") _log(f"llm_init_start pp_size={pp_size}") llm = LLM( @@ -115,6 +153,7 @@ def cmd_run(args): temperature=0.0, max_tokens=MAX_NEW_TOKENS, ignore_eos=True, + logprobs=NUM_LOGPROBS, ) prompt = TokensPrompt(prompt_token_ids=prompt_ids) @@ -130,12 +169,24 @@ def cmd_run(args): faulthandler.cancel_dump_traceback_later() _log("generate_done") - generated_ids = list(outputs[0].outputs[0].token_ids) + completion = outputs[0].outputs[0] + generated_ids = list(completion.token_ids) _log(f"generated {len(generated_ids)} tokens: {generated_ids[:10]}...") + # Per-step gap between the best and second-best candidate. compare uses this + # to tell "the two arms disagree" from "the argmax was a tie and the two arms + # broke it differently", which is not a defect in anything this test covers. + gaps = [] + for step in completion.logprobs or []: + values = sorted((lp.logprob for lp in step.values()), reverse=True) + gaps.append(values[0] - values[1] if len(values) > 1 else float("inf")) + _log(f"min_top2_gap={min(gaps) if gaps else float('nan')}") + output_path = os.path.join(work_dir, f"{tag}.json") with open(output_path, "w") as f: - json.dump({"token_ids": generated_ids, "pp_size": pp_size}, f) + json.dump( + {"token_ids": generated_ids, "pp_size": pp_size, "top2_gaps": gaps}, f + ) _log(f"saved {output_path}") del llm @@ -147,15 +198,32 @@ def cmd_run(args): def cmd_compare(args): - """Load pp1.json / pp2.json and assert token-for-token match.""" + """Compare pp1.json / pp2.json, treating an exactly-tied argmax as equal. + + Tie-tolerant rather than token-exact, and the tolerance is measured per + position rather than assumed. PP=1 and PP=2 do not run the same computation + bit-for-bit: vLLM downgrades CUDA-graph capture above PP=1, and splitting the + stack across two ranks changes the float reduction order. Where the top two + candidates are already equal in bf16, that sub-ULP difference decides which + one greedy emits -- an arbitrary coin flip, not a routing defect. Measured on + an A100 (vLLM 0.26) this prompt reaches a top-2 gap of exactly 0.0, so a + token-exact assertion here cannot be made stable. + + A broken per-rank recompute does not look like that. A rank that rebuilt its + bitmasks or remapped indices wrongly applies the wrong adapter delta, which + moves the whole distribution and separates the argmax -- a divergence with a + nonzero gap, which still fails below. + """ work_dir = args.work_dir pp1_path = os.path.join(work_dir, "pp1.json") pp2_path = os.path.join(work_dir, "pp2.json") with open(pp1_path) as f: - pp1_ids = json.load(f)["token_ids"] + pp1 = json.load(f) with open(pp2_path) as f: - pp2_ids = json.load(f)["token_ids"] + pp2 = json.load(f) + pp1_ids, pp2_ids = pp1["token_ids"], pp2["token_ids"] + pp1_gaps = pp1.get("top2_gaps") or [] print(f" pp1 tokens ({len(pp1_ids)}): {pp1_ids}") print(f" pp2 tokens ({len(pp2_ids)}): {pp2_ids}") @@ -164,17 +232,32 @@ def cmd_compare(args): print(f"\nFAIL: length mismatch: pp1={len(pp1_ids)}, pp2={len(pp2_ids)}") return 1 + ties, real = [], [] for i, (a, b) in enumerate(zip(pp1_ids, pp2_ids)): - if a != b: - print( - f"\nFAIL: first divergence at position {i}: pp1={a}, pp2={b}\n" - " PP=2 must recompute per-rank LoRA kernel metadata identically " - "to PP=1; a mismatch means the recompute-per-rank contract is broken." - ) - return 1 + if a == b: + continue + # No gap recorded (older artifact, or no tokenizer) => cannot prove a tie, + # so treat it as a real divergence rather than passing on faith. + gap = pp1_gaps[i] if i < len(pp1_gaps) else 0.0 + (ties if gap <= 0.0 else real).append((i, a, b, gap)) + + for i, a, b, gap in ties: + print(f" tie at position {i}: pp1={a}, pp2={b} (pp1 top-2 gap {gap:.6f})") + + if real: + i, a, b, gap = real[0] + print( + f"\nFAIL: divergence at position {i}: pp1={a}, pp2={b}, " + f"pp1 top-2 gap {gap:.6f}.\n" + " The argmax was strictly separated there, so this is not a tie-break: " + "PP=2 must recompute per-rank LoRA kernel metadata identically to PP=1, " + "and a mismatch means the recompute-per-rank contract is broken." + ) + return 1 print( - f"\nPASS: PP=2 matches PP=1 token-for-token [{len(pp1_ids)} tokens] — " + f"\nPASS: PP=2 matches PP=1 over {len(pp1_ids)} tokens " + f"({len(ties)} exactly-tied position(s) allowed to differ) — " "per-rank metadata recompute is correct." ) return 0 From 3aad32ece77167f21d602ae415b0fd0bf11242bc Mon Sep 17 00:00:00 2001 From: itzikvaisman Date: Mon, 26 Oct 2026 10:35:30 +0200 Subject: [PATCH 15/15] Apply cache pressure with a second conversation, not a bigger one test_correctness_survives_cache_eviction never ran. It died at engine init: ValueError: To serve at least one request with the model's max seq len (4096), 0.31 GiB KV cache is needed, which is larger than the available KV cache memory (0.02 GiB). Based on the available memory, the estimated maximum model length is 256. _launch hardcoded max_model_len=4096 while the test passed --blocks 16, a 256-token cache. Threading max_model_len through so it can be pinned to capacity is necessary but not sufficient -- the test as written had an empty solution set. The arithmetic: vLLM will not start unless the cache holds one whole max_model_len request, so every servable turn satisfies capacity >= max_model_len >= prompt + max_tokens. Under PRESERVE, turn N's prompt *contains* turn N-1's transcript as its own prefix, i.e. the same blocks. So a single conversation the server can serve is one whose history still fits and nothing is ever evicted, while the first turn that would overflow is the turn the server refuses. There is no --blocks value in between. A second conversation with a different filler has a disjoint prefix and therefore its own blocks, so it applies pressure the first conversation did not create. Measured at turn 3 the two prompts are ~173 and ~148 tokens: each alone needs 189 of the 256 available and stays servable, while together they need ~321 and evict each other. The run now reports "eviction bit on ['turn 1/A']" with a 16-block shortfall against the block-math ceiling, and every turn's decode still matches its expected adapter index -- which is the claim: pressure costs recomputation, not correctness. Per-turn records carry a conv label so the printed table and the assertion messages say which conversation a turn belongs to. Signed-off-by: itzikvaisman --- .../_conversation_concurrency_worker.py | 138 ++++++++++++------ .../test_conversation_concurrency.py | 60 ++++++-- 2 files changed, 135 insertions(+), 63 deletions(-) diff --git a/tests/integration/_conversation_concurrency_worker.py b/tests/integration/_conversation_concurrency_worker.py index 42170eea..43389ff0 100644 --- a/tests/integration/_conversation_concurrency_worker.py +++ b/tests/integration/_conversation_concurrency_worker.py @@ -245,7 +245,7 @@ def decode_indices_by_seqlen(records, prompt_len, max_new): # ── shared server plumbing ──────────────────────────────────────────────────── -def _launch(model_path, port, workdir, extra): +def _launch(model_path, port, workdir, extra, max_model_len=4096): from granite_switch.tutorials.vllm_server import ( kill_stale_vllm_processes, launch_vllm, @@ -264,7 +264,7 @@ def _launch(model_path, port, workdir, extra): log_file=log_file, max_num_seqs=8, enforce_eager=True, - max_model_len=4096, + max_model_len=max_model_len, extra_args=_supported_flags( ["--enable-prefix-caching", "--return-tokens-as-token-ids", *extra] ), @@ -446,7 +446,7 @@ def cmd_evict(args): from granite_switch import Conversation, KVHistoryPolicy tok, config, amap, aloras, expected = _setup(args.model_path) - result = {"blocks": args.blocks, "turns": []} + result = {"blocks": args.blocks, "max_model_len": args.blocks * 16, "turns": []} if len(aloras) < 2: result["error"] = f"need >=2 alora adapters, got {aloras}" Path(args.output_path).write_text(json.dumps(result)) @@ -455,62 +455,104 @@ def cmd_evict(args): a_name, b_name = aloras[0], aloras[1] workdir = tempfile.mkdtemp(prefix="gs_evict_") - # A cache far too small to keep the conversation resident, so early blocks are - # evicted between turns. The point is that correctness must not follow the hit - # rate down. + # ── Why TWO conversations, and why max_model_len tracks the cache ───────── + # + # A single growing conversation cannot be made to evict itself, and the + # arithmetic says so. vLLM refuses to start unless the KV cache can hold one + # whole max_model_len request (measured: --num-gpu-blocks-override 16 gives a + # 256-token cache, and against the previous hardcoded max_model_len=4096 the + # engine died at init with "estimated maximum model length is 256"). So + # capacity >= max_model_len >= prompt_N + MAX_TOKENS + # must hold for turn N to be servable at all. But under PRESERVE turn N's + # prompt *contains* turn N-1's transcript as its own prefix -- they are the + # same blocks -- so any turn the server can serve is a turn whose history + # fits entirely in cache. Nothing is ever evicted, and the anti-vacuity + # shortfall assert can never fire. The one turn that would overflow the cache + # (measured turn 5: 245 + 16 = 261 > 256) is also the one the server cannot + # serve. That is an empty solution set, not a tuning problem, which is why + # this does not simply retune --blocks. + # + # Eviction therefore has to come from pressure the request does not create + # itself: a SECOND conversation with a different filler, hence a disjoint + # token prefix and its own blocks. Measured at turn 3 the two prompts are + # ~173 and ~148 tokens; each alone needs 189 <= 256 so both stay servable, + # while together they need ~321 > 256 and so evict each other's blocks. That + # window is non-empty, which is the property the old configuration lacked. + # + # max_model_len is pinned to the cache size for the same reason: leaving it + # at 4096 is what made the constraint unsatisfiable. + capacity = args.blocks * 16 proc, trace, log_file = _launch( args.model_path, args.port, workdir, ["--num-gpu-blocks-override", str(args.blocks)], + max_model_len=capacity, ) try: _wait_ready(args.port, proc, log_file) offset = os.path.getsize(trace) if os.path.exists(trace) else 0 - conv = Conversation( - tok, policy=KVHistoryPolicy.PRESERVE_MIXED_HISTORY, config=config - ) + # Distinct fillers => distinct prefixes => each conversation owns its own + # blocks, so a hit shortfall on one is attributable to the other's + # pressure rather than to a shared prefix being counted twice. + convs = { + label: Conversation( + tok, policy=KVHistoryPolicy.PRESERVE_MIXED_HISTORY, config=config + ) + for label in ("A", "B") + } for turn in range(args.turns): adapter = a_name if turn % 2 == 0 else b_name - question = ( - f"{FILLERS['B']}Turn {turn}: which techniques apply?" - if turn == 0 - else f"Turn {turn}: restate that as JSON." - ) - conv.user(_user_text(question, amap[adapter][2])) - sent_before = len(conv.sent_token_ids) - before = _metrics(args.port) - prompt, src = _run_turn( - args.port, args.model_path, conv, None, adapter, f"evict{turn}" - ) - after = _metrics(args.port) - records, offset = _read_trace(trace, offset) - result["turns"].append( - { - "turn": turn, - "adapter": adapter, - "expected_index": expected(adapter), - "prompt_len": len(prompt), - "answer_ids_source": src, - # What the block arithmetic makes available for THIS turn: the - # transcript as it stood before the turn, rounded down to whole - # 16-token blocks. Hits below it mean blocks that should have been - # resident were not -- direct evidence of eviction, rather than an - # inference from cache size. - "reusable_ceiling": (sent_before // 16) * 16, - "sent_len": len(conv.sent_token_ids), - "hits_delta": ( - after.get("vllm:prefix_cache_hits_total", 0.0) - - before.get("vllm:prefix_cache_hits_total", 0.0) - ), - "decode_indices": decode_indices_by_seqlen( - records, len(prompt), MAX_TOKENS - ), - "history_routing": sorted( - set(attribute_spans(records, list(prompt)).values()) - ), - } - ) + # Both conversations take their turn before either takes the next, so + # each one's blocks sit in cache across the other's turn -- that is + # when they get evicted. + for label in ("A", "B"): + conv = convs[label] + question = ( + f"{FILLERS[label]}Turn {turn}: which techniques apply?" + if turn == 0 + else f"Turn {turn}: restate that as JSON." + ) + conv.user(_user_text(question, amap[adapter][2])) + sent_before = len(conv.sent_token_ids) + before = _metrics(args.port) + prompt, src = _run_turn( + args.port, + args.model_path, + conv, + None, + adapter, + f"evict{label}{turn}", + ) + after = _metrics(args.port) + records, offset = _read_trace(trace, offset) + result["turns"].append( + { + "turn": turn, + "label": label, + "adapter": adapter, + "expected_index": expected(adapter), + "prompt_len": len(prompt), + "answer_ids_source": src, + # What the block arithmetic makes available for THIS turn: + # the transcript as it stood before the turn, rounded down + # to whole 16-token blocks. Hits below it mean blocks that + # should have been resident were not -- direct evidence of + # eviction, rather than an inference from cache size. + "reusable_ceiling": (sent_before // 16) * 16, + "sent_len": len(conv.sent_token_ids), + "hits_delta": ( + after.get("vllm:prefix_cache_hits_total", 0.0) + - before.get("vllm:prefix_cache_hits_total", 0.0) + ), + "decode_indices": decode_indices_by_seqlen( + records, len(prompt), MAX_TOKENS + ), + "history_routing": sorted( + set(attribute_spans(records, list(prompt)).values()) + ), + } + ) finally: _shutdown(proc, log_file) diff --git a/tests/integration/test_conversation_concurrency.py b/tests/integration/test_conversation_concurrency.py index ad149b61..448ab03a 100644 --- a/tests/integration/test_conversation_concurrency.py +++ b/tests/integration/test_conversation_concurrency.py @@ -11,11 +11,14 @@ forwards. The arm fails loudly if the scheduler never actually co-batched them, because then it proves nothing about concurrency. -EVICT a cache far too small to keep the conversation resident. PRESERVE's *saving* - must degrade -- that is inherent, and not a bug -- while its *correctness* must - not. The same ids are recomputed and recomputation reproduces the same routing. - If correctness tracked cache pressure the policy would be unusable in - production, where pressure is the normal condition. +EVICT a cache far too small to keep both conversations resident. PRESERVE's + *saving* must degrade -- that is inherent, and not a bug -- while its + *correctness* must not. The same ids are recomputed and recomputation + reproduces the same routing. If correctness tracked cache pressure the policy + would be unusable in production, where pressure is the normal condition. + Pressure comes from a second conversation rather than from one conversation + outgrowing the cache; see the test's own docstring for why that is the only + satisfiable form. Attribution under concurrency is the part that has produced false results before, so it lives in ``attribute_spans`` and is unit-tested on CPU in @@ -149,21 +152,46 @@ def test_two_preserve_conversations_do_not_disturb_each_other(): def test_correctness_survives_cache_eviction(): + """Cache pressure must cost recomputation, never routing correctness. + + Two conversations share one deliberately tiny cache, rather than one + conversation growing past it. That is forced by arithmetic, not taste. vLLM + will not start unless the KV cache holds one whole ``max_model_len`` request, + so every servable turn satisfies ``capacity >= max_model_len >= prompt + + max_tokens``; and under PRESERVE a turn's prompt *contains* the previous + turn's transcript as its own prefix, i.e. the same blocks. So a single + conversation that the server can serve is one whose history still fits, and + nothing is ever evicted -- while the first turn that would overflow the cache + is the one the server refuses. The earlier single-conversation form of this + test sat in exactly that gap and died at engine init (a 256-token cache + against ``max_model_len=4096``), which is why the fix is a second + conversation rather than a different ``--blocks``. + + A second conversation with a different filler has a disjoint prefix and thus + its own blocks, so it applies pressure the first conversation did not create. + Measured at turn 3 the two prompts are ~173 and ~148 tokens: each alone needs + 189 of the 256 available and stays servable, while together they need ~321 + and evict each other. + """ data = _run("evict", _model_dir(), extra=["--blocks", "16", "--turns", "4"]) assert not data.get("error"), data["error"] - print(f"\n cache: {data['blocks']} blocks ({data['blocks'] * 16} tokens)") + capacity = data["blocks"] * 16 + print( + f"\n cache: {data['blocks']} blocks ({capacity} tokens), " + f"max_model_len {data.get('max_model_len', capacity)}" + ) hdr = ( - f"\n {'turn':>4s} {'prompt':>7s} {'hits':>6s} {'ceil':>6s} {'short':>6s} " - f"{'want':>5s} decode history" + f"\n {'turn':>4s} {'conv':>4s} {'prompt':>7s} {'hits':>6s} {'ceil':>6s} " + f"{'short':>6s} {'want':>5s} decode history" ) print(hdr) print(" " + "-" * (len(hdr) - 3)) for t in data["turns"]: short = t["reusable_ceiling"] - t["hits_delta"] print( - f" {t['turn']:>4d} {t['prompt_len']:>7d} {t['hits_delta']:>6.0f} " - f"{t['reusable_ceiling']:>6d} {short:>6.0f} " + f" {t['turn']:>4d} {t.get('label', '-'):>4s} {t['prompt_len']:>7d} " + f"{t['hits_delta']:>6.0f} {t['reusable_ceiling']:>6d} {short:>6.0f} " f"{t['expected_index']:>5d} {t['decode_indices']!s:>8s} " f"{t['history_routing']}" ) @@ -171,6 +199,9 @@ def test_correctness_survives_cache_eviction(): turns = data["turns"] assert turns, "no turns were run" + def _name(t): + return f"turn {t['turn']}/{t.get('label', '-')}" + # ── Anti-vacuity: eviction has to have actually bitten ────────────────── # NOT "the final prompt exceeds the cache": a request larger than the whole # cache cannot be served at all, so that condition can never hold on a run @@ -178,9 +209,8 @@ def test_correctness_survives_cache_eviction(): # rather than strict. Eviction is instead evidenced directly: hits below what # the block arithmetic made available means blocks that should have been # resident were not. - capacity = data["blocks"] * 16 shortfalls = { - t["turn"]: t["reusable_ceiling"] - t["hits_delta"] + _name(t): t["reusable_ceiling"] - t["hits_delta"] for t in turns if t["hits_delta"] < t["reusable_ceiling"] } @@ -191,20 +221,20 @@ def test_correctness_survives_cache_eviction(): "not below one full request, or the server cannot serve at all." ) print( - f" eviction bit on turns {sorted(shortfalls)} " + f" eviction bit on {sorted(shortfalls)} " f"(shortfall vs the block-math ceiling: {shortfalls})" ) # ── The claim: the saving degrades, the routing does not ──────────────── for t in turns: assert t["decode_indices"] == [t["expected_index"]], ( - f"turn {t['turn']}: generated with adapter index(es) " + f"{_name(t)}: generated with adapter index(es) " f"{t['decode_indices']}, wanted {t['expected_index']}. Cache pressure " "must cost recomputation, never correctness -- evicted blocks are " "recomputed from the same ids and must reproduce the same routing." ) assert t["answer_ids_source"] == "logprobs", ( - f"turn {t['turn']}: the answer was re-encoded from text, so the " + f"{_name(t)}: the answer was re-encoded from text, so the " "transcript may not match what was sent and any hit rate below is " "unattributable to eviction" )