Skip to content

[OMNIML-5899] Add IQ1_S and IQ2_XS quantization and unified checkpoint export - #2381

Closed
ChenhanYu wants to merge 8 commits into
mainfrom
iq2xs-unified-export
Closed

ChenhanYu wants to merge 8 commits into
mainfrom
iq2xs-unified-export

Conversation

@ChenhanYu

@ChenhanYu ChenhanYu commented Sep 11, 2026

Copy link
Copy Markdown
Collaborator

What does this PR do?

Type of change: new feature

Adds IQ1_S and IQ2_XS weight quantization and unified Hugging Face and Megatron Core checkpoint export (TP=1). The implementation provides deterministic CPU reference encoders/decoders, optional CUDA encoders, the canonical format codebooks, and GGML-compatible 50-byte IQ1_S or 74-byte IQ2_XS blocks for every 256 logical weights. Unified safetensors export writes each supported weight as packed_weights plus weight_shape, with format metadata and reusable PTQ recipes.

Usage

python examples/hf_ptq/hf_ptq.py \
  --pyt_ckpt_path <huggingface_model_card_or_path> \
  --recipe general/ptq/iq1_s \
  --export_path <quantized_checkpoint_path>

Use general/ptq/iq2_xs for IQ2_XS.

Testing

  • uv run --frozen pytest -q tests/unit/torch/quantization/test_iq1_s.py tests/unit/torch/quantization/test_iq2_xs.py tests/unit/torch/export/test_export_weight.py tests/unit/torch/export/test_get_quantization.py tests/unit/recipe/test_recipe_docs.py — 35 passed.
  • uv run --frozen pytest -q tests/unit/recipe/test_recipe_docs.py tests/unit/recipe/test_loader.py — 337 passed.
  • uvx --from pre-commit==4.6.0 pre-commit run --from-ref origin/main --to-ref HEAD — all applicable hooks passed.
  • Added CUDA determinism/layout tests for both formats; they were not run locally because this environment has no GPU or nvcc.

Before your PR is "Ready for review"

Make sure you read and follow Contributor guidelines and your commits are signed (git commit -s -S).

Make sure you read and follow the Security Best Practices (e.g. avoiding hardcoded trust_remote_code=True, torch.load(..., weights_only=False), pickle, etc.).

  • Is this change backward compatible?: ✅
  • If you copied code from any other sources or added a new PIP dependency, did you follow guidance in CONTRIBUTING.md: ❌ — llama.cpp MIT attribution and third-party notices are included; the internal OSRB/NVBug record remains to be updated before ready-for-review.
  • Did you write any new necessary tests?: ✅
  • Did you update Changelog?: ✅
  • Did you get Claude approval on this PR?: ❌ — pending /claude review.

Additional Information

The canonical IQ1_S and IQ2_XS codebooks are adapted from llama.cpp at revision 9b05354e under its MIT license. The search behavior is aligned with the PSX-LUTS search_impl="auto" path used for QAT. Runtime integration in vLLM is intentionally out of scope for this PR.

@copy-pr-bot

copy-pr-bot Bot commented Sep 11, 2026

Copy link
Copy Markdown

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@coderabbitai

coderabbitai Bot commented Sep 11, 2026

Copy link
Copy Markdown
Contributor

Important

Draft PR not reviewed

Draft PRs are not automatically reviewed by default.

  • Trigger a manual review

To automatically review draft PRs, update your CodeRabbit configuration:

reviews:
  auto_review:
    drafts: true

Comment @coderabbitai help to get the list of available commands.

@codecov

codecov Bot commented Sep 11, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 78.94737% with 72 lines in your changes missing coverage. Please review.
✅ Project coverage is 69.80%. Comparing base (d69e93a) to head (73baffe).
⚠️ Report is 18 commits behind head on main.

Files with missing lines Patch % Lines
modelopt/torch/export/unified_export_megatron.py 6.52% 43 Missing ⚠️
modelopt/torch/quantization/extensions.py 37.50% 10 Missing ⚠️
modelopt/torch/quantization/ggml/iq1_s.py 94.44% 6 Missing ⚠️
modelopt/torch/quantization/ggml/iq2_xs.py 94.69% 6 Missing ⚠️
modelopt/torch/quantization/ggml/backend.py 54.54% 5 Missing ⚠️
modelopt/torch/export/quant_utils.py 83.33% 1 Missing ⚠️
modelopt/torch/quantization/ggml/common.py 93.75% 1 Missing ⚠️
Additional details and impacted files
@@            Coverage Diff             @@
##             main    #2381      +/-   ##
==========================================
- Coverage   70.96%   69.80%   -1.17%     
==========================================
  Files         542      595      +53     
  Lines       63784    68763    +4979     
==========================================
+ Hits        45266    47998    +2732     
- Misses      18518    20765    +2247     
Flag Coverage Δ
unit 57.96% <78.94%> (+0.79%) ⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@ChenhanYu ChenhanYu changed the title [OMNIML-5899] Add IQ2_XS quantization and unified checkpoint export to ModelOpt [OMNIML-5899] Add IQ1_S and IQ2_XS quantization and unified checkpoint export Sep 11, 2026
@ChenhanYu

Copy link
Copy Markdown
Collaborator Author

/claude review

@copy-pr-bot

copy-pr-bot Bot commented Sep 11, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

# Save the merged weights
if merged_weight_scale is None:
if qformat in (QUANTIZATION_IQ1_S, QUANTIZATION_IQ2_XS):
self._state_dict.update(self._get_iq_weight_state(prefix, merged_weight, qformat))

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[IMPORTANT Export] _get_iq_weight_state appends a . to prefix, but at this call site (and at the _pack_name_remapping_gpt_oss twin around line 2025) prefix is a complete tensor key, not a module prefix.

The non-IQ branches right below treat it that way:

elif merged_weight_scale is None:
    self._state_dict[prefix] = merged_weight
else:
    self._state_dict[prefix] = to_quantized_weight(...)
    self._state_dict[prefix + "_weight_scale"] = merged_weight_scale

and the rule book confirms it — mcore_llama.py:72 maps local_experts.linear_fc1 to "language_model.model.layers.{}.feed_forward.experts.gate_up_proj" and mcore_gptoss.py:41 to "model.layers.{}.mlp.experts.gate_up_proj", neither with a trailing dot. Since _get_iq_weight_state does prefix = prefix.rstrip(".") + ".", packed MoE experts emit

  • ...experts.gate_up_proj.packed_weights
  • ...experts.gate_up_proj.weight_shape

next to siblings named ...experts.gate_up_proj_weight_scale / ..._input_scale. Every other IQ call site (1232, 1280/1283, 1486, 1653, 1775) passes a real module prefix ending in ., so those come out right — only the two pack_name_remapping paths (Llama4 / GPT-OSS MoE) produce this inconsistent key shape. A consumer that resolves gate_up_proj as a parameter name plus suffixes will not find these tensors.

Suggested fix: don't normalize inside the helper; take the two conventions explicitly at the call site, e.g.

# in _get_iq_weight_state, replace the rstrip/append with:
#   (no prefix rewriting; caller supplies the exact base)
return {
    f"{prefix}_packed_weights": packed_weights.detach().cpu(),
    f"{prefix}_weight_shape": weight_shape.detach().cpu(),
}

and have the module-prefix call sites pass prefix + "weight" so they still yield ...weight_packed_weights-style keys — or, if the dotted form is the intended wire format, pass prefix + "gate_up_proj." style bases here instead of the bare parameter key. Either way the two packed-expert paths should match the naming the rest of the exporter uses for that tensor.


@torch.no_grad()
def quantize_iq1_s(
weight: torch.Tensor, *, block_chunk_size: int = 4

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[IMPORTANT Performance] block_chunk_size=4 makes the CPU/eager encoder launch-bound by ~16×. _encode_blocks runs a fixed Python nest of 16 grid tiles × 2 shifts × 8 local scales = 256 iterations plus the matmuls, per chunk, independent of chunk size. With 4 blocks per chunk, a 4096×4096 weight (65 536 blocks) needs 16 384 chunks ⇒ ~4.2 M tiny tensor ops; at 64 blocks it's 1 024 chunks for the same arithmetic.

Memory isn't the constraint: the largest temporary is error/shifted_dot at chunk × 32 × 16 floats, ~1 MB at chunk 64 — and quantize_iq2_xs in this same PR already defaults to 64 with a larger grid (512×8 with per-entry sign search). The asymmetry looks unintentional.

Suggested change
weight: torch.Tensor, *, block_chunk_size: int = 4
weight: torch.Tensor, *, block_chunk_size: int = 64

Comment on lines +187 to +191
error = (
xnorm.unsqueeze(-1) - 2 * scale * shifted_dot + scale.square() * shifted_norm
)
tile_error, tile_index = error.min(dim=-1)
replace = tile_error < best_error[:, :, choice]

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[IMPORTANT Compatibility] The IQ1_S eager error is left unclamped, while iq1_s.cu clamps it in both search stages (fmaxf(error, 0.0f)) and the sibling iq2_xs.py:186 applies .clamp_min_(0). So the two IQ1_S implementations minimize different objectives.

The clamp is not cosmetic for argmin: xnorm - 2·scale·dot + scale²·norm is an exact-arithmetic-nonnegative expression evaluated in floating point, so near-perfect fits land at small values of either sign. When one candidate rounds to -1e-7 and another to +1e-7, unclamped CPU prefers the negative one and clamped CUDA sees a tie broken by lowest entry index — different entry, different payload byte. The same applies at the group_error.argmin stage, where 4 clamped-vs-unclamped terms are summed.

Fix: mirror the kernel (and IQ2_XS) here —

                error = (
                    xnorm.unsqueeze(-1)
                    - 2 * scale * shifted_dot
                    + scale.square() * shifted_norm
                ).clamp_min_(0)

Two related gaps worth closing in the same pass (details in a top-level comment): the CPU↔CUDA payloads are never asserted equal by any test, and get_cuda_ext_iq1_s/get_cuda_ext_iq2_xs build with --use_fast_math, which is a second independent source of divergence.

Comment on lines +25 to +34
packed, shape = quantize_iq1_s(weight)
packed_again, _ = quantize_iq1_s(weight)
reconstructed = dequantize_iq1_s(packed, shape)

assert packed.shape == (8, 1, 50)
assert torch.equal(packed, packed_again)
normalized_mse = (
reconstructed.float() - weight.float()
).square().mean() / weight.float().square().mean()
assert normalized_mse < 0.25

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[IMPORTANT Compatibility] These GPU tests only check that the CUDA encoder agrees with itself (packed == packed_again) and lands under an MSE bound. Nothing in the suite asserts that the CUDA payload equals the eager payload for the same input — yet the two are presented as interchangeable implementations behind search_impl: auto, so which one produced a checkpoint depends on where the export ran.

Three concrete divergence sources, all live today:

  1. The eager error at iq1_s.py:187 is unclamped; iq1_s.cu applies fmaxf(..., 0.0f) in both stages (flagged inline).
  2. extensions.py builds both IQ kernels with --use_fast_math, i.e. imprecise div/sqrt and denormal flush-to-zero. find_scale computes d and rounds it to fp16 — a single ULP of drift in amax / _IQ1_S_NATIVE_MAX * 0.61 before the __float2half can tip d to the neighbouring fp16 value, after which the entire 50-byte block differs, not one field.
  3. Reduction order differs by construction: the kernel accumulates dot with sequential fmaf across a warp shuffle tree, the eager path uses matmul/sum(dim=-1).

Suggested test to add here (the useful assertion is bitwise, since the payload is what ships):

def test_iq1_s_cuda_pack_matches_eager_pack():
    generator = torch.Generator().manual_seed(1234)
    weight = torch.randn((8, 512), generator=generator, dtype=torch.bfloat16)

    eager, eager_shape = quantize_iq1_s(weight, search_impl="eager")
    cuda, cuda_shape = quantize_iq1_s(weight.cuda())

    assert torch.equal(cuda.cpu(), eager)
    assert torch.equal(cuda_shape.cpu(), eager_shape)

(adjust to whatever knob forces the eager path — backend_extra_args={"search_impl": ...} today only reaches it through the quantizer). If exact equality is not a goal, that's a legitimate design choice, but then it should be stated in the module docstring, because it means the exported GGML bytes are host-dependent; a tolerance-based check on the dequantized values plus a documented note would be the alternative.

Comment on lines +49 to +69
def test_iq1_s_dequantizes_ggml_metadata_bit_fields():
packed = torch.zeros((1, 1, 50), dtype=torch.uint8)
d = torch.tensor([2.0], dtype=torch.float16).view(torch.uint8)
packed[0, 0, :2] = d
entries = torch.tensor([0, 256, 511, 2047], dtype=torch.int64)
packed[0, 0, 2:6] = (entries & 0xFF).to(torch.uint8)
qh = (
((entries[0] >> 8) & 7)
| (((entries[1] >> 8) & 7) << 3)
| (((entries[2] >> 8) & 7) << 6)
| (((entries[3] >> 8) & 7) << 9)
| (3 << 12)
| (1 << 15)
)
packed[0, 0, 34] = (qh & 0xFF).to(torch.uint8)
packed[0, 0, 35] = (qh >> 8).to(torch.uint8)

decoded = dequantize_iq1_s(packed, torch.tensor([1, 256]), dtype=torch.float32)
expected = (iq1_s_grid()[entries] - 0.125) * 14.0

assert torch.equal(decoded[0, :32].reshape(4, 8), expected)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[SUGGESTION] This is the closest thing to a format-compatibility test, and it's self-referential: expected is computed with iq1_s_grid() and the same (g - 0.125) * d formula the decoder uses, so an error in the grid table itself or in the scale ladder passes. The whole point of IQ1_S/IQ2_XS is that llama.cpp can read the bytes.

Worth adding a handful of golden vectors — a fixed input array with the 50-byte (resp. 74-byte) payload produced by llama-quantize/ggml_quantize_chunk, checked in as a literal, plus the expected dequantize_row_iq1_s output. That pins the grid contents, the 2*ls+1 ladder, the ±0.125 delta, the qh field packing, and (for IQ2_XS) the ksigns_iq2xs even-parity convention against the reference in one assertion. Byte-exact encoder equality is a strong ask given the search heuristics differ; asserting the decoder reproduces reference GGML blocks is cheap and hermetic (tests/unit stays offline — just constants).

)
if not weight.is_floating_point():
raise TypeError(f"{format_name} requires a floating-point weight, got {weight.dtype}")
if not torch.isfinite(weight).all():

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[SUGGESTION] torch.isfinite(weight).all() in a Python if forces a device→host sync on every quantize_iq* call, and it reads the full tensor. CONTRIBUTING's coding standards call out avoiding CPU-GPU syncs specifically. It's once per weight rather than per block, so the cost is bounded — but the check also doesn't buy much: amax of a tensor containing inf/NaN already produces a non-finite d that clamp(max=65504) turns into either 65504 or NaN, and a NaN weight is a bug upstream. Either drop it or gate it behind a debug flag, so the hot export path doesn't serialize on it.

Comment on lines +120 to +132
elif quant_algo in ("IQ1_S", "IQ2_XS"):
effective_bits, payload_bytes = (1.5625, 50) if quant_algo == "IQ1_S" else (2.3125, 74)
return {
"weights": {
"dynamic": False,
"num_bits": 1 if quant_algo == "IQ1_S" else 2,
"effective_bits": effective_bits,
"type": "int",
"group_size": 256,
"packing": "ggml",
"block_payload_bytes": payload_bytes,
}
}

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[SUGGESTION] This branch hardcodes "group_size": 256 and silently ignores the group_size argument (gs) that every other branch honours. That's correct for GGML — the 256-value block is baked into the 50/74-byte layout — but a caller who passes group_size=128 gets a config claiming 256 with no warning, and the mismatch surfaces only when a consumer tries to unpack. Worth asserting the contract instead of dropping the argument:

    elif quant_algo in ("IQ1_S", "IQ2_XS"):
        if gs not in (None, 256):
            warnings.warn(f"{quant_algo} uses a fixed GGML block of 256; ignoring group_size={gs}")
        effective_bits, payload_bytes = (1.5625, 50) if quant_algo == "IQ1_S" else (2.3125, 74)

Also "type": "int" with "num_bits": 1 is a lossy description of the format for anything reading this generically — the values are ternary {-1,0,1} (IQ1_S) / magnitude-grid codes (IQ2_XS) with a shared fp16 block scale, not 1-bit integers. effective_bits and packing: "ggml" carry the real information; consider a comment noting num_bits is nominal here so a future reader doesn't try to derive the size from it.

@claude

claude Bot commented Sep 11, 2026

Copy link
Copy Markdown

Claude review — full review of [OMNIML-5899] Add IQ1_S and IQ2_XS quantization and unified checkpoint export

Counts: 0 CRITICAL · 4 IMPORTANT · 3 SUGGESTION

What I verified as correct

I checked the bit-level format work against llama.cpp semantics rather than taking it on faith, and it holds up:

  • IQ1_S — 50-byte block, 2048×8 ternary grid, ±0.125 delta, 2*ls+1 scale ladder, qh field packing (4×3-bit high entry bits + 3-bit local scale + 1-bit shift). Matches dequantize_row_iq1_s.
  • IQ2_XS — 74-byte block, 512×8 magnitude grid {8,25,43}, db = 0.25*(0.5+ls), scale-nibble ordering, and the 7-bit sign field with the 8th sign implied by even parity (ksigns_iq2xs convention). The decoder recomputing bit 7 as the parity of bits 0..6 is consistent with the encoder's odd-parity dot correction and weakest-element sign flip.
  • State-dict plumbingpacked_weights / _iq_weight_shape traced end-to-end. The _iq_ prefix trick genuinely works: _BASE_SKIP_KEYS matches weight_shape by substring and would drop the tensor, but _KV_CACHE_REPLACEMENTS["_iq_weight_shape"] = "weight_shape" renames it after the skip test, in both _postprocess_single_tensor and postprocess_state_dict. The Megatron exporter doesn't call postprocess_state_dict, so its direct weight_shape keys survive too.
  • CUDA kernels — shared-memory budget for IQ2_XS (512×8 grid + 512 norms ≈ 18.7 KB, fits 48 KB), __syncthreads() ordering around the payload[66+tid] nibble packing, uniform early returns on the d == 0 path, and the packed (error_bits<<32)|index argmin key giving deterministic lowest-index tie-breaking. All sound.
  • Backend dispatch_fake_quantize routes to psx_luts before any amax/block_sizes handling, so the IQ path correctly needs no calibration. pass_through_bwd: true in the numerics YAMLs is therefore redundant (harmless).

Most impactful findings

1. _get_iq_weight_state prefix convention mismatch (IMPORTANT Export) — inline

The helper does prefix = prefix.rstrip(".") + ".", which is right for the six module-prefix call sites but wrong for the two pack_name_remapping sites (:1912, :2025), where prefix is a complete tensor key — confirmed against mcore_llama.py:72 and mcore_gptoss.py:41, neither of which has a trailing dot. Llama4 / GPT-OSS packed MoE experts therefore emit ...experts.gate_up_proj.packed_weights beside ...experts.gate_up_proj_weight_scale. Structural fix, so it goes here rather than inline:

    @staticmethod
    def _get_iq_weight_state(
        prefix: str, weight: torch.Tensor, qformat: str
    ) -> dict[str, torch.Tensor]:
        quantize_iq = quantize_iq1_s if qformat == QUANTIZATION_IQ1_S else quantize_iq2_xs
        packed_weights, weight_shape = quantize_iq(weight)
        # `prefix` is the exact base of the tensor key; callers append the separator
        # they need ("." for a module prefix, "_" for a packed-parameter key) so this
        # matches the `prefix + "_weight_scale"` convention used by the non-IQ paths.
        return {
            prefix + "packed_weights": packed_weights.detach().cpu(),
            prefix + "weight_shape": weight_shape.detach().cpu(),
        }

with the two packed-expert call sites becoming self._get_iq_weight_state(prefix + "_", merged_weight, qformat) and the module-prefix sites keeping their existing trailing-dot prefixes. (Whichever direction you pick, the two conventions should be explicit at the call site instead of normalized away inside the helper.)

2. CPU and CUDA encoders are not asserted to agree, and one known asymmetry exists (IMPORTANT Compatibility) — inline: clamp · inline: test gap

iq1_s.py:187 leaves the search error unclamped; iq1_s.cu applies fmaxf(error, 0.0f) in both stages, and iq2_xs.py:186 uses .clamp_min_(0). That is IQ1_S eager minimizing a different objective from IQ1_S CUDA, and it matters for argmin: the expression is nonnegative in exact arithmetic, so good fits land at small values of either sign and the sign decides the winner. On top of that, both kernels build with --use_fast_math (imprecise div/sqrt, FTZ) and use a different reduction order, and d is rounded to fp16 — one ULP of drift there changes the whole 50-byte block, not one field. The GPU test only checks packed == packed_again, so none of this is currently detectable. If bitwise CPU↔CUDA equality is not a goal, that is a defensible call, but it should be documented, because it makes the exported GGML bytes host-dependent.

3. quantize_iq1_s(block_chunk_size=4) (IMPORTANT Performance) — inline

_encode_blocks runs a fixed 16-tile × 2-shift × 8-scale Python nest per chunk regardless of chunk size, so chunk 4 means ~16× the launches for the same arithmetic (~4.2 M tiny ops for a 4096×4096 weight vs ~260 K at chunk 64). Not memory-motivated — the largest temporary is ~1 MB at chunk 64 — and quantize_iq2_xs already defaults to 64 with a larger grid.

Suggestions (non-blocking)

  • No llama.cpp golden byte vectors; test_iq1_s_dequantizes_ggml_metadata_bit_fields builds expected from the same grid and formula the decoder uses, so the tests self-verify the encoder/decoder pair without pinning GGML compatibility.
  • validate_weight's torch.isfinite(weight).all() forces a CPU-GPU sync per call and adds little.
  • The IQ branch of _quant_algo_to_group_config silently ignores its group_size argument, and "type": "int", "num_bits": 1 is a lossy generic description of a ternary-grid format.

Overall risk

Moderate, and concentrated in export rather than in the quantizers. The math and the byte layouts are the parts I would have expected to be risky, and they are the parts that check out — the encoders, decoders, and both kernels are internally consistent and consistent with GGML. What is unresolved is (a) two MoE export paths emitting keys in a shape nothing else in the exporter uses, which will bite Llama4/GPT-OSS users specifically and silently, and (b) two implementations of the same format that are asserted to be interchangeable but never compared, with a concrete asymmetry already visible in the source. Both are fixable without touching the format work. Nothing here risks producing wrong values — a mismatched key or a differently-searched-but-valid block still dequantizes correctly — which is why none of it is CRITICAL.

Coverage

Reviewed modelopt/ in full (all ggml modules, both .cu kernels, extensions.py, quant_format.py, quant_utils.py, convert_hf_config.py, unified_export_hf.py, unified_export_megatron.py, tensor_quantizer.py, config.py), the iq1_s recipe trio, and the IQ1_S unit + GPU tests. Not opened: .pre-commit-config.yaml, LICENSE, CHANGELOG.rst, modelopt_recipes/ptq.md, the iq2_xs.yaml recipe trio, the two .cpp binding files, and the test_iq2_xs* / test_unified_export_megatron.py / test_export_weight.py / test_get_quantization.py test files. Findings 2 and 3 likely have IQ2_XS analogues in the files I did not open — worth a look when addressing them.

I also could not obtain the true merge base (shallow checkout; origin/main has advanced 3 commits past base d69e93a), so I worked from the authoritative file list in gh pr view --json files and read the working tree directly. Reverse hunks from main's divergence (e.g. kv_cache_auto_quant.py, _has_large_fp8_scale) were treated as not part of this PR.

@claude claude Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Claude review complete — 0 CRITICAL, 4 IMPORTANT, 3 SUGGESTION. Not approving; the IMPORTANT findings should be resolved first. Full write-up in the summary comment above.

IMPORTANT:

  1. Export_get_iq_weight_state normalizes prefix to end in ., but the two pack_name_remapping call sites (unified_export_megatron.py:1912, :2025) pass a complete tensor key, not a module prefix. Llama4 / GPT-OSS packed MoE experts end up with ...experts.gate_up_proj.packed_weights next to ...experts.gate_up_proj_weight_scale. The other six IQ call sites are fine.
  2. Compatibility — the IQ1_S eager search error is unclamped (ggml/iq1_s.py:187) while iq1_s.cu applies fmaxf(error, 0.0f) in both stages and iq2_xs.py:186 uses .clamp_min_(0). The two IQ1_S implementations minimize different objectives, which changes argmin outcomes near zero error.
  3. Compatibility — nothing asserts CPU/CUDA payload equality; test_iq1_s_cuda.py only checks self-determinism. Combined with --use_fast_math on both kernels and the fp16 rounding of d, exported GGML bytes can depend on which host produced them.
  4. Performancequantize_iq1_s defaults to block_chunk_size=4, ~16× more launches than needed for identical arithmetic; quantize_iq2_xs already uses 64 with a larger grid.

SUGGESTION (non-blocking): no llama.cpp golden byte vectors pinning GGML compatibility; validate_weight's torch.isfinite(...).all() forces a CPU-GPU sync; the IQ branch of _quant_algo_to_group_config silently ignores its group_size argument.

The format work itself — both bit layouts, scale ladders, the IQ2_XS parity convention, the state-dict rename-after-skip mechanism, and both kernels' shared-memory and synchronization structure — checked out against llama.cpp and against itself. The open items are in export plumbing and test coverage, not in the quantizers.

@github-actions

github-actions Bot commented Sep 11, 2026

Copy link
Copy Markdown
Contributor
PR Preview Action v1.8.1
Preview removed because the pull request was closed.
2026-09-16 00:32 UTC

Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Add TP=1 Megatron unified checkpoint export for IQ1_S and IQ2_XS, including final-layout packing and coverage for the scale-free payload schema.

Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Store IQ payloads under canonical weight keys and remove redundant logical-shape tensors.\n\nSigned-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>

Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
Register packed quantized expert weights as buffers so unified export retains them. Add regression coverage for buffer-backed fused-expert projections.

Signed-off-by: Chenhan D. Yu <5185878+ChenhanYu@users.noreply.github.com>
@hychiang-git

Copy link
Copy Markdown
Contributor

This large PR is superseded by three focused draft PRs:

The export and CUDA PRs will be retargeted to main after #2443 merges.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants