IQ2_XS

hexstate_quantize.c

@@ -1793,11 +1793,13 @@ void hexstate_set_spectral_params(float dc_lambda, float vw_lambda, float dc_dec
         g_hex_dc_decay = dc_decay;
 }
 
-/* Spectral penalty: DC² + Σ vesica². Always ≥ 0 (no wave credit). */
-static inline float hex_spectral_penalty(const float *e, int n)
+/* Spectral penalty: (Σe + dc_carry)² + Σ vesica². dc_carry is the
+ * rolling residual we want this block to cancel (Σe ≈ −dc_carry).
+ * Reconstruction target is always the true weight; we never quantize x−bias. */
+static inline float hex_spectral_penalty_ex(const float *e, int n, float dc_carry)
 {
     if (HEX_DC_LAMBDA == 0.0f && HEX_VW_LAMBDA == 0.0f) return 0.0f;
-    float dc = 0.0f, ves = 0.0f;
+    float dc = dc_carry, ves = 0.0f;
     int half = n / 2;
     for (int i = 0; i < half; i++) {
         float v = e[i] + e[i + half];
@@ -1808,6 +1810,190 @@ static inline float hex_spectral_penalty(const float *e, int n)
          + (HEX_VW_LAMBDA / (float)n) * ves;
 }
 
+static inline float hex_spectral_penalty(const float *e, int n)
+{
+    return hex_spectral_penalty_ex(e, n, 0.0f);
+}
+
+/* Relative SSE we may spend to hit Σe ≈ −carry. 5e-4 ⇒ ~0.025% RMSE.
+ * Runtime-tunable (hexstate_set_sse_budget) — codebook formats need more. */
+static float g_hex_sse_budget = 5.0e-4f;
+#define HEX_DC_SSE_BUDGET (g_hex_sse_budget)
+void hexstate_set_sse_budget(float rel) { if (rel >= 0.0f) g_hex_sse_budget = rel; }
+
+static inline void hex_q2k_unpack_L(const BlockQ2K *b, uint8_t L[QK_K])
+{
+    for (int j = 0; j < QK_K; j += 128) {
+        for (int l = 0; l < 32; l++) {
+            uint8_t p = b->qs[j / 4 + l];
+            L[j + l]      = (uint8_t)( p        & 3);
+            L[j + l + 32] = (uint8_t)((p >> 2)  & 3);
+            L[j + l + 64] = (uint8_t)((p >> 4)  & 3);
+            L[j + l + 96] = (uint8_t)((p >> 6)  & 3);
+        }
+    }
+}
+
+static inline void hex_q2k_pack_L(BlockQ2K *b, const uint8_t L[QK_K])
+{
+    for (int j = 0; j < QK_K; j += 128) {
+        for (int l = 0; l < 32; l++) {
+            b->qs[j / 4 + l] = (uint8_t)(L[j + l]
+                | (L[j + l + 32] << 2)
+                | (L[j + l + 64] << 4)
+                | (L[j + l + 96] << 6));
+        }
+    }
+}
+
+static inline float hex_q2k_el_w(const float *imat, int64_t blk, int i)
+{
+    return imat ? imat[blk * QK_K + i] : 1.0f;
+}
+
+/* Frozen codes: put (d, dmin) on Σ(x−deq) = −carry if SSE stays in budget. */
+static int hex_q2k_hit_dc_carry(const float *x, const uint8_t *L,
+                                const uint8_t *scales, float *dm, float *mm,
+                                float dc_carry, const float *w256)
+{
+    float d0 = *dm, m0 = *mm;
+    double A = 0.0, B = 0.0, Sx = 0.0, sse0 = 0.0, dc0 = 0.0;
+    for (int i = 0; i < QK_K; i++) {
+        int j = i >> 4;
+        float a = (float)(scales[j] & 0xF) * (float)L[i];
+        float b = (float)(scales[j] >> 4);
+        float wi = w256 ? w256[i] : 1.0f;
+        float e = x[i] - (d0 * a - m0 * b);
+        sse0 += (double)wi * e * e;
+        dc0  += e;
+        A += a; B += b; Sx += x[i];
+    }
+    double T = Sx + (double)dc_carry;
+    float d_try = d0, m_try = m0;
+
+    if (fabs(B) > 1e-12) {
+        double k = A / B, tB = T / B;
+        double Szz = 0.0, Syz = 0.0;
+        for (int i = 0; i < QK_K; i++) {
+            int j = i >> 4;
+            float a = (float)(scales[j] & 0xF) * (float)L[i];
+            float b = (float)(scales[j] >> 4);
+            float wi = w256 ? w256[i] : 1.0f;
+            double z = (double)a - k * (double)b;
+            double y = (double)x[i] - tB * (double)b;
+            Szz += (double)wi * z * z;
+            Syz += (double)wi * y * z;
+        }
+        if (Szz < 1e-30) return 0;
+        double d_ref = Syz / Szz;
+        double m_ref = (A * d_ref - T) / B;
+        if (d_ref <= 0.0 || m_ref < 0.0) return 0;
+        d_try = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)d_ref));
+        m_try = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)m_ref));
+    } else if (fabs(A) > 1e-12) {
+        double d_ref = T / A;
+        if (d_ref <= 0.0) return 0;
+        double num = 0.0, den = 0.0;
+        for (int i = 0; i < QK_K; i++) {
+            int j = i >> 4;
+            float a = (float)(scales[j] & 0xF) * (float)L[i];
+            float b = (float)(scales[j] >> 4);
+            float wi = w256 ? w256[i] : 1.0f;
+            num += (double)wi * ((double)x[i] - d_ref * (double)a) * (double)b;
+            den += (double)wi * (double)b * (double)b;
+        }
+        if (den < 1e-30) return 0;
+        double m_ref = -num / den;
+        if (m_ref < 0.0) return 0;
+        d_try = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)d_ref));
+        m_try = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)m_ref));
+    } else {
+        return 0;
+    }
+    if (d_try <= 0.0f || m_try < 0.0f) return 0;
+
+    double sse1 = 0.0, dc1 = 0.0;
+    for (int i = 0; i < QK_K; i++) {
+        int j = i >> 4;
+        float a = (float)(scales[j] & 0xF) * (float)L[i];
+        float b = (float)(scales[j] >> 4);
+        float wi = w256 ? w256[i] : 1.0f;
+        float e = x[i] - (d_try * a - m_try * b);
+        sse1 += (double)wi * e * e;
+        dc1  += e;
+    }
+    double off0 = dc0 + (double)dc_carry;
+    double off1 = dc1 + (double)dc_carry;
+    if (fabs(off1) >= fabs(off0) - 1e-12) return 0;
+    if (sse1 > sse0 * (1.0 + (double)HEX_DC_SSE_BUDGET)) return 0;
+    *dm = d_try;
+    *mm = m_try;
+    return 1;
+}
+
+/* q±1 toward Σe ≈ −carry, spending at most HEX_DC_SSE_BUDGET extra SSE. */
+static void hex_q2k_dc_nudge_codes(const float *x, uint8_t *L,
+                                   const uint8_t *scales, float dm, float mm,
+                                   float dc_carry, const float *w256)
+{
+    float e[QK_K];
+    float sse = 0.0f, dc = 0.0f;
+    for (int i = 0; i < QK_K; i++) {
+        int j = i >> 4;
+        float d_s = dm * (float)(scales[j] & 0xF);
+        float m_s = mm * (float)(scales[j] >> 4);
+        e[i] = x[i] - (d_s * (float)L[i] - m_s);
+        float wi = w256 ? w256[i] : 1.0f;
+        sse += e[i] * e[i] * wi;
+        dc  += e[i];
+    }
+    float cap = sse * (1.0f + HEX_DC_SSE_BUDGET);
+    float median_step = dm * 4.0f;
+    if (median_step < 1e-15f) median_step = 1e-15f;
+
+    for (int pass = 0; pass < 64; pass++) {
+        float off = dc + dc_carry;
+        if (fabsf(off) <= median_step) break;
+        int best_i = -1, best_q = 0;
+        float best_ratio = 0.0f;
+        for (int i = 0; i < QK_K; i++) {
+            int j = i >> 4;
+            float d_s = dm * (float)(scales[j] & 0xF);
+            float m_s = mm * (float)(scales[j] >> 4);
+            if (d_s < 1e-15f) continue;
+            int q_cur = (int)L[i];
+            int q_try = (off > 0.0f) ? q_cur + 1 : q_cur - 1;
+            if (q_try < 0 || q_try > 3) continue;
+            float e_new = x[i] - (d_s * (float)q_try - m_s);
+            float dc_red = fabsf(off) - fabsf(off + (e_new - e[i]));
+            if (dc_red <= 0.0f) continue;
+            float wi = w256 ? w256[i] : 1.0f;
+            float sse_new = sse + wi * (e_new * e_new - e[i] * e[i]);
+            if (sse_new > cap) continue;
+            float sse_cost = sse_new - sse;
+            if (sse_cost < 0.0f) sse_cost = 0.0f;
+            float ratio = dc_red / (sse_cost + 1e-20f);
+            if (ratio > best_ratio) {
+                best_ratio = ratio;
+                best_i = i;
+                best_q = q_try;
+            }
+        }
+        if (best_i < 0) break;
+        {
+            int j = best_i >> 4;
+            float d_s = dm * (float)(scales[j] & 0xF);
+            float m_s = mm * (float)(scales[j] >> 4);
+            float e_new = x[best_i] - (d_s * (float)best_q - m_s);
+            float wi = w256 ? w256[best_i] : 1.0f;
+            sse += wi * (e_new * e_new - e[best_i] * e[best_i]);
+            dc  += (e_new - e[best_i]);
+            e[best_i] = e_new;
+            L[best_i] = (uint8_t)best_q;
+        }
+    }
+}
+
 /* Robust temperature estimator for the HExState measurement model.
  *
  * The old path estimated T from the mean of each block's MAXIMUM candidate
@@ -3607,49 +3793,29 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
     }
 
     /* ══════════════════════════════════════════════════════════════════
-     * PHASE 3.9 — ROLLING DC BOUNDARY CONDITION PRE-PASS
-     *
-     * Transforms the tensor from a collection of isolated 256-element
-     * Q2_K superblocks into a single, continuous error-cancelling waveform.
-     *
-     * After Phase 3 has selected the optimal (d, dmin) candidate for every
-     * block, this sequential pass computes the net DC residual left by each
-     * block using a cheap round-nearest forward quantization, then feeds the
-     * negated, exponentially-decayed residual as a correction bias into the
-     * WLS solver of the immediately following block.
-     *
-     * Mathematically, for block N with final DC residual R_N = Σ εᵢ:
-     *
-     *   dc_bias[N+1] = −DC_DECAY × R_N / QK_K      (per-element offset)
+     * PHASE 3.9 — ROLLING DC *RESIDUAL CARRY* (does NOT shift weights)
      *
-     * Block N+1's WLS targets become x′ᵢ = xᵢ − dc_bias[N+1], steering the
-     * quantizer toward codes whose reconstruction deq ≈ x′, so that
+     * Bug that killed PPL: we used to quantize x′ = x − bias so llama
+     * stored a DC-shifted matrix. Cancellation belongs on the residual
+     * e = x − deq of an unshifted reconstruction:
      *
-     *   Σ (xᵢ − deqᵢ) ≈ dc_bias[N+1] × QK_K = −DC_DECAY × R_N
+     *   carry[N] = DC_DECAY · Σ e_{N−1}
+     *   prefer Σ e_N ≈ −carry[N]
      *
-     * The accumulated cross-block DC collapses geometrically:
-     *
-     *   R₀, DC_DECAY·R₀, DC_DECAY²·R₀, …  → 0
-     *
-     * The result is written into block_dc_bias[n_blocks].  Phase 4 reads
-     * this array (safe: written sequentially before the parallel loop).
+     * via the spectral term (��e + carry)², while every code/scale is
+     * still chosen against the true x.
      * ══════════════════════════════════════════════════════════════════ */
 
     #define DC_DECAY (g_hex_dc_decay)
 
-    float *block_dc_bias = (float *)calloc(n_blocks, sizeof(float));
+    float *block_dc_carry = (float *)calloc(n_blocks, sizeof(float));
 
-    if (block_dc_bias) {
+    if (block_dc_carry) {
         float rolling_dc = 0.0f;
-        /* Row-boundary awareness: blocks_per_row = row_width / QK_K.
-         * Rows are independent dot products — DC residual from the end of
-         * one row must NOT leak into the start of the next.  When row_width
-         * is 0 (unknown), fall back to the old flat-stream behaviour. */
         int64_t blocks_per_row = (row_width > 0 && row_width % QK_K == 0)
                                  ? row_width / QK_K : 0;
 
         for (int64_t blk = 0; blk < n_blocks; blk++) {
-            /* Reset DC at every row boundary */
             if (blocks_per_row > 0 && (blk % blocks_per_row) == 0)
                 rolling_dc = 0.0f;
 
@@ -3664,29 +3830,25 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
             hex_derive_subscales(seeds[blk].scales, seeds[blk].mins,
                                  dm0, mm0, dc_Ls, dc_Lm);
 
-            /* Bias applied to THIS block's WLS targets */
-            float dc_bias       = (DC_DECAY * rolling_dc) / (float)QK_K;
-            block_dc_bias[blk]  = dc_bias;
+            /* Residual-space carry: next block should cancel this, not
+             * reconstruct a shifted x. */
+            block_dc_carry[blk] = DC_DECAY * rolling_dc;
 
-            /* Quick round-nearest quant to estimate DC residual for NEXT block.
-             * We quantize the adjusted target x′ = x − dc_bias, then measure
-             * the residual of the ORIGINAL weight against the chosen code. */
             float dc_res = 0.0f;
             int   j, k;
             for (j = 0; j < N_SUB; j++) {
                 float d_sub = dm0 * (float)dc_Ls[j];
                 float m_sub = mm0 * (float)dc_Lm[j];
                 for (k = 0; k < 16; k++) {
-                    float x_adj = bx[16*j + k] - dc_bias;
+                    float x = bx[16*j + k];
                     int q = 0;
                     if (d_sub >= 1e-15f) {
-                        q = gguf_nearest_int((x_adj + m_sub) / d_sub);
+                        q = gguf_nearest_int((x + m_sub) / d_sub);
                         if (q < 0) q = 0;
                         if (q > 3) q = 3;
                     }
                     float deq = d_sub * (float)q - m_sub;
-                    /* Residual against ORIGINAL weight (not adjusted) */
-                    dc_res += bx[16*j + k] - deq;
+                    dc_res += x - deq;
                 }
             }
             rolling_dc = dc_res;
@@ -3710,21 +3872,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
         const float *block_x = weights + blk * QK_K;
         int cidx = best_candidate[blk];
         uint8_t Ls_blk[16], Lm_blk[16];
-
-        /* ── Rolling DC boundary condition ──────────────────────────────
-         * dc_adj shifts every WLS target in this block so that the net
-         * quantisation error steers toward cancelling the previous block's
-         * DC residual (written by the sequential Phase 3.9 pre-pass). */
-        float dc_adj = (block_dc_bias) ? block_dc_bias[blk] : 0.0f;
-
-        /* Adjusted weight view — WLS and sieve work on this array;
-         * the final error is always reported against the original block_x. */
-        float adj_block_x[QK_K];
-        {
-            int _i;
-            for (_i = 0; _i < QK_K; _i++)
-                adj_block_x[_i] = block_x[_i] - dc_adj;
-        }
+        const float dc_carry = (block_dc_carry) ? block_dc_carry[blk] : 0.0f;
 
         uint16_t base_c_d16, base_c_m16;
         hex_candidate_pair(seeds[blk].base_dm, seeds[blk].base_mm, cidx, &base_c_d16, &base_c_m16);
@@ -3742,7 +3890,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
             float state_err[N_SUB][6];
 
             for (int j = 0; j < N_SUB; j++) {
-                const float *sx = adj_block_x + 16 * j;
+                const float *sx = block_x + 16 * j;
                 for (int v = 0; v < 6; v++) state_err[j][v] = 1e30f;
 
                 for (int try_ls = 0; try_ls <= 15; try_ls++) {
@@ -3850,7 +3998,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     continue;
                 }
                 for (int k = 0; k < 16; k++) {
-                    int q = gguf_nearest_int((adj_block_x[16*j+k] + m_sub) / d_sub);
+                    int q = gguf_nearest_int((block_x[16*j+k] + m_sub) / d_sub);
                     if (q < 0) q = 0; if (q > 3) q = 3;
                     L[16*j+k] = (uint8_t)q;
                 }
@@ -3861,7 +4009,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                 float ls_f = (float)Ls_blk[j];
                 float lm_f = (float)Lm_blk[j];
                 for (int k = 0; k < 16; k++) {
-                    float x = adj_block_x[16*j+k];
+                    float x = block_x[16*j+k];
                     float w = (imat_importance) ?
                               imat_importance[blk * QK_K + 16*j+k] : 1.0f;
                     float a = ls_f * (float)L[16*j+k];
@@ -3917,7 +4065,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                         float d_sub = trial_dm * (float)Ls_blk[j];
                         float m_sub = trial_mm * (float)Lm_blk[j];
                         for (int k = 0; k < 16; k++) {
-                            float x = adj_block_x[16*j+k];
+                            float x = block_x[16*j+k];
                             float w = (imat_importance) ?
                                       imat_importance[blk * QK_K + 16*j+k] : 1.0f;
                             int q;
@@ -3943,7 +4091,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
         }
 
         for (int j = 0; j < N_SUB; j++) {
-            const float *sx = adj_block_x + 16 * j;
+            const float *sx = block_x + 16 * j;
             float best_sub_err = 1e30f;
             uint8_t best_ls = Ls_blk[j], best_lm = Lm_blk[j];
             for (int try_ls = 0; try_ls <= 15; try_ls++) {
@@ -4010,8 +4158,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     q_cont_all[i] = 0.0f;
                     q_base_all[i] = 0;
                 } else {
-                    /* Quantize the DC-adjusted target */
-                    float qc = (adj_block_x[i] + m_s) / d_s;
+                    float qc = (block_x[i] + m_s) / d_s;
                     q_cont_all[i] = qc;
                     int qr = gguf_nearest_int(qc);
                     if (qr < 0) qr = 0; if (qr > 3) qr = 3;
@@ -4021,17 +4168,19 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
             memcpy(q_shaped_all, q_base_all, QK_K * sizeof(int));
 
             float e_live[QK_K];
-            float dc_cur = 0.0f;
+            float dc_cur = dc_carry;
+            float sse_live = 0.0f;
             for (int i = 0; i < QK_K; i++) {
                 int   jj  = i >> 4;
                 float d_s = dm * (float)(output[blk].scales[jj] & 0xF);
                 float m_s = mm * (float)(output[blk].scales[jj] >> 4);
                 float deq = d_s * (float)q_shaped_all[i] - m_s;
-                /* Vesica on ORIGINAL residuals. DC on the rolling-adjusted
-                 * residual so this block still absorbs Phase 3.9's target. */
                 e_live[i] = block_x[i] - deq;
-                dc_cur   += adj_block_x[i] - deq;
+                dc_cur   += e_live[i];
+                float w = (imat_importance) ? imat_importance[blk * QK_K + i] : 1.0f;
+                sse_live += e_live[i] * e_live[i] * w;
             }
+            const float sse_cap = sse_live * (1.0f + HEX_DC_SSE_BUDGET);
 
             float v_live[QK_K / 2];
             float vesica_cur = 0.0f;
@@ -4064,6 +4213,9 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                         if (q_try == q_cur) continue;
 
                         float e_new = block_x[k] - (d_s * (float)q_try - m_s);
+                        float w = (imat_importance) ? imat_importance[blk * QK_K + k] : 1.0f;
+                        float sse_alt = sse_live + w * (e_new * e_new - e_live[k] * e_live[k]);
+                        if (sse_alt > sse_cap) continue;
                         float de    = e_new - e_live[k];
 
                         float v_new = v_live[pi] + de;
@@ -4096,29 +4248,35 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     metric_cur    = 4.0f * vesica_cur + dc_cur * dc_cur;
                     v_live[pi_c]  = v_new_c;
                     e_live[best_k]= e_new_c;
+                    {
+                        float w_c = (imat_importance) ?
+                                    imat_importance[blk * QK_K + best_k] : 1.0f;
+                        sse_live += w_c * (e_new_c * e_new_c
+                            - (e_new_c - de_c) * (e_new_c - de_c));
+                    }
                 }
             }
 
-            /* Choose base vs shaped on the EXTENDED objective vs originals */
-            float err_base = 0.0f, err_shaped = 0.0f;
+            float sse_base = 0.0f, sse_shaped = 0.0f;
             float e_qb[QK_K], e_qs[QK_K];
             for (int i = 0; i < QK_K; i++) {
                 int   jj  = i >> 4;
                 float d_s = dm * (float)(output[blk].scales[jj] & 0xF);
                 float m_s = mm * (float)(output[blk].scales[jj] >> 4);
                 float w   = (imat_importance) ? imat_importance[blk * QK_K + i] : 1.0f;
-                float deq_b = d_s * (float)q_base_all[i]   - m_s;  /* −m_s when d_s==0 */
+                float deq_b = d_s * (float)q_base_all[i]   - m_s;
                 float deq_s = d_s * (float)q_shaped_all[i] - m_s;
-                float xv    = block_x[i];   /* original weight for error report */
+                float xv    = block_x[i];
                 e_qb[i] = xv - deq_b;
                 e_qs[i] = xv - deq_s;
-                err_base   += e_qb[i] * e_qb[i] * w;
-                err_shaped += e_qs[i] * e_qs[i] * w;
+                sse_base   += e_qb[i] * e_qb[i] * w;
+                sse_shaped += e_qs[i] * e_qs[i] * w;
             }
-            err_base   += hex_spectral_penalty(e_qb, QK_K);
-            err_shaped += hex_spectral_penalty(e_qs, QK_K);
+            float err_base   = sse_base   + hex_spectral_penalty_ex(e_qb, QK_K, dc_carry);
+            float err_shaped = sse_shaped + hex_spectral_penalty_ex(e_qs, QK_K, dc_carry);
             {
-                int use_shaped = (err_shaped <= err_base);
+                int use_shaped = (sse_shaped <= sse_base * (1.0f + HEX_DC_SSE_BUDGET)
+                                  && err_shaped <= err_base);
                 for (int i = 0; i < QK_K; i++)
                     L[i] = (uint8_t)(use_shaped ? q_shaped_all[i] : q_base_all[i]);
             }
@@ -4157,7 +4315,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                 for (fs_k = 0; fs_k < 16; fs_k++) {
                     int   idx    = base + fs_k;
                     float x_orig = block_x[idx];
-                    float x_adj  = adj_block_x[idx] + carry;  /* adjusted + diffused */
+                    float x_adj  = block_x[idx] + carry;  /* adjusted + diffused */
 
                     /* Propose new code from diffused target */
                     int q_fs = gguf_nearest_int((x_adj + m_s) / d_s);
@@ -4178,7 +4336,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     /* Propagate 7/16 of the residual (adj target vs committed code) */
                     {
                         float deq_final = d_s * (float)L[idx] - m_s;
-                        float residual  = (adj_block_x[idx] - deq_final);
+                        float residual  = (block_x[idx] - deq_final);
                         carry = (fs_k < 15) ? residual * (7.0f / 16.0f) : 0.0f;
                     }
                 }
@@ -4199,7 +4357,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
         {
             /* Compute current whole-block DC and per-weight residuals */
             float wb_e[QK_K];
-            float wb_dc = 0.0f;
+            float wb_dc = 0.0f, wb_sse = 0.0f;
             for (int i = 0; i < QK_K; i++) {
                 int   jj  = i >> 4;
                 float d_s = dm * (float)(output[blk].scales[jj] & 0xF);
@@ -4207,21 +4365,21 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                 float deq = d_s * (float)L[i] - m_s;
                 wb_e[i] = block_x[i] - deq;
                 wb_dc += wb_e[i];
+                float w = (imat_importance) ? imat_importance[blk * QK_K + i] : 1.0f;
+                wb_sse += wb_e[i] * wb_e[i] * w;
             }
 
-            /* Target: reduce |DC| to ≤ 1 quantization step of the median
-             * sub-block scale, which is the noise floor for DC. */
-            float median_step = dm * 4.0f;  /* rough: d * median(Ls) ≈ d*4 */
+            float wb_off = wb_dc + dc_carry;
+            float median_step = dm * 4.0f;
             if (median_step < 1e-15f) median_step = 1e-15f;
+            float wb_cap = wb_sse * (1.0f + HEX_DC_SSE_BUDGET);
 
             for (int dc_pass = 0; dc_pass < 32; dc_pass++) {
-                if (fabsf(wb_dc) <= median_step) break;
+                if (fabsf(wb_off) <= median_step) break;
 
-                /* Find the weight whose nudge toward reducing DC costs the
-                 * least in weighted SSE. */
                 int   best_i  = -1;
                 int   best_q  = 0;
-                float best_ratio = 0.0f;  /* |dc_reduction| / sse_cost */
+                float best_ratio = 0.0f;
 
                 for (int i = 0; i < QK_K; i++) {
                     int   jj  = i >> 4;
@@ -4229,28 +4387,22 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     float m_s = mm * (float)(output[blk].scales[jj] >> 4);
                     if (d_s < 1e-15f) continue;
 
-                    /* Nudge direction: if DC > 0, we need negative residual
-                     * change (increase q to reduce x−deq), and vice versa. */
                     int q_cur = (int)L[i];
-                    int q_try = (wb_dc > 0.0f) ? q_cur + 1 : q_cur - 1;
+                    int q_try = (wb_off > 0.0f) ? q_cur + 1 : q_cur - 1;
                     if (q_try < 0 || q_try > 3) continue;
 
                     float deq_new = d_s * (float)q_try - m_s;
                     float e_new   = block_x[i] - deq_new;
-                    float dc_reduction = fabsf(wb_dc) - fabsf(wb_dc + (e_new - wb_e[i]));
+                    float dc_reduction = fabsf(wb_off) - fabsf(wb_off + (e_new - wb_e[i]));
                     if (dc_reduction <= 0.0f) continue;
 
                     float w = (imat_importance) ?
                               imat_importance[blk * QK_K + i] : 1.0f;
-                    float sse_old = wb_e[i] * wb_e[i] * w;
-                    float sse_new = e_new * e_new * w;
-                    float sse_cost = sse_new - sse_old;
-                    if (sse_cost < 0.0f) sse_cost = 0.0f;  /* free improvement */
-
-                    /* Ratio: how much DC reduction per unit SSE cost.
-                     * Free moves (sse_cost ≤ 0) get infinite ratio. */
-                    float ratio = (sse_cost < 1e-30f) ? 1e30f
-                                  : dc_reduction / sse_cost;
+                    float sse_new = wb_sse + w * (e_new * e_new - wb_e[i] * wb_e[i]);
+                    if (sse_new > wb_cap) continue;
+                    float sse_cost = sse_new - wb_sse;
+                    if (sse_cost < 0.0f) sse_cost = 0.0f;
+                    float ratio = dc_reduction / (sse_cost + 1e-20f);
                     if (ratio > best_ratio) {
                         best_ratio = ratio;
                         best_i     = i;
@@ -4258,17 +4410,19 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     }
                 }
 
-                if (best_i < 0) break;  /* no improving nudge found */
+                if (best_i < 0) break;
 
-                /* Accept: the minimum-cost nudge toward DC reduction.
-                 * Only accept if SSE cost is modest relative to DC gain. */
                 {
                     int   jj  = best_i >> 4;
                     float d_s = dm * (float)(output[blk].scales[jj] & 0xF);
                     float m_s = mm * (float)(output[blk].scales[jj] >> 4);
                     float deq_new = d_s * (float)best_q - m_s;
                     float e_new   = block_x[best_i] - deq_new;
-                    wb_dc += (e_new - wb_e[best_i]);
+                    float w = (imat_importance) ?
+                              imat_importance[blk * QK_K + best_i] : 1.0f;
+                    wb_sse += w * (e_new * e_new - wb_e[best_i] * wb_e[best_i]);
+                    wb_dc  += (e_new - wb_e[best_i]);
+                    wb_off  = wb_dc + dc_carry;
                     wb_e[best_i] = e_new;
                     L[best_i] = (uint8_t)best_q;
                 }
@@ -4311,9 +4465,10 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
              * λ_dc/n; vesica/wave handled by the extended-E acceptance. */
             {
                 double rw = (double)HEX_DC_LAMBDA / (double)QK_K;
+                double rSt = rS + (double)dc_carry;
                 rSaa += rw * rA * rA;  rSab += rw * rA * rB;
-                rSbb += rw * rB * rB;  rSxa += rw * rS * rA;
-                rSxb += rw * rS * rB;
+                rSbb += rw * rB * rB;  rSxa += rw * rSt * rA;
+                rSxb += rw * rSt * rB;
             }
             double rdet = rSaa * rSbb - rSab * rSab;
             if (fabs(rdet) > 1e-30) {
@@ -4343,9 +4498,12 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                             err_try += e_rt[idx] * e_rt[idx] * w;
                         }
                     }
-                    err_cur += hex_spectral_penalty(e_rc, QK_K);
-                    err_try += hex_spectral_penalty(e_rt, QK_K);
-                    if (err_try < err_cur) { dm = dm_try; mm = mm_try; }
+                    float sse_cur = err_cur, sse_try = err_try;
+                    err_cur += hex_spectral_penalty_ex(e_rc, QK_K, dc_carry);
+                    err_try += hex_spectral_penalty_ex(e_rt, QK_K, dc_carry);
+                    if (sse_try <= sse_cur && err_try < err_cur) {
+                        dm = dm_try; mm = mm_try;
+                    }
                 }
             }
             output[blk].d    = gguf_fp32_to_fp16(dm);
@@ -4357,7 +4515,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
          * Objective-function mismatch fix: the final passes that commit the
          * 2-bit codes — the 16×16 (ls, lm) sub-block search, the ±8 ULP
          * (d, dmin) neighborhood search, and the greedy-descent error shaping
-         * — all minimise error against the DC-ADJUSTED target adj_block_x.
+         * — all minimise error against the DC-ADJUSTED target block_x.
          * The reported RMSE, however, is measured against the ORIGINAL
          * weights. The codes are therefore stranded at the optimum of a
          * SHIFTED objective, while only the scalar (d, dmin) refit above
@@ -4443,7 +4601,8 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
 
                     /* Extended score of the CURRENT committed state */
                     float best_sub = sub_sse[j]
-                        + (HEX_DC_LAMBDA / (float)QK_K) * dc_tot * dc_tot
+                        + (HEX_DC_LAMBDA / (float)QK_K)
+                          * (dc_tot + dc_carry) * (dc_tot + dc_carry)
                         + (HEX_VW_LAMBDA / (float)QK_K) * ves_tot;
                     int     best_ls = -1, best_lm = 0;
                     uint8_t best_q[16];
@@ -4480,9 +4639,11 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                                 if (sub_err >= best_sub) { aborted = 1; break; }
                             }
                             if (aborted) continue;
+                            if (sub_err > sub_sse[j]) continue;
+                            float dcc_tot = dc_rest + dcc + dc_carry;
                             float score = sub_err
                                 + (HEX_DC_LAMBDA / (float)QK_K)
-                                  * (dc_rest + dcc) * (dc_rest + dcc)
+                                  * dcc_tot * dcc_tot
                                 + (HEX_VW_LAMBDA / (float)QK_K)
                                   * (ves_rest + vesc);
                             if (score < best_sub) {
@@ -4536,9 +4697,10 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     }
                     {
                         double pw = (double)HEX_DC_LAMBDA / (double)QK_K;
+                        double pSt = pS + (double)dc_carry;
                         pSaa += pw * pA * pA;  pSab += pw * pA * pB;
-                        pSbb += pw * pB * pB;  pSxa += pw * pS * pA;
-                        pSxb += pw * pS * pB;
+                        pSbb += pw * pB * pB;  pSxa += pw * pSt * pA;
+                        pSxb += pw * pSt * pB;
                     }
                     double pdet = pSaa * pSbb - pSab * pSab;
                     if (fabs(pdet) > 1e-30) {
@@ -4570,9 +4732,10 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                                     err_try += e_pt[idx] * e_pt[idx] * w;
                                 }
                             }
-                            err_cur += hex_spectral_penalty(e_pc, QK_K);
-                            err_try += hex_spectral_penalty(e_pt, QK_K);
-                            if (err_try < err_cur) {
+                            float sse_cur = err_cur, sse_try = err_try;
+                            err_cur += hex_spectral_penalty_ex(e_pc, QK_K, dc_carry);
+                            err_try += hex_spectral_penalty_ex(e_pt, QK_K, dc_carry);
+                            if (sse_try <= sse_cur && err_try < err_cur) {
                                 dm = dm_try;
                                 mm = mm_try;
                                 pol_improved = 1;
@@ -4609,7 +4772,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                             cur_err += e_u[idx] * e_u[idx] * w;
                         }
                     }
-                    cur_err += hex_spectral_penalty(e_u, QK_K);
+                    cur_err += hex_spectral_penalty_ex(e_u, QK_K, dc_carry);
 
                     float    best_err = cur_err;
                     uint16_t best_d16 = base_d16, best_m16 = base_m16;
@@ -4639,7 +4802,7 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                                 if (err >= best_err) { pruned = 1; break; }
                             }
                             if (pruned) continue;
-                            err += hex_spectral_penalty(e_u, QK_K);
+                            err += hex_spectral_penalty_ex(e_u, QK_K, dc_carry);
                             if (err < best_err) {
                                 best_err = err;
                                 best_d16 = (uint16_t)cd16;
@@ -4690,8 +4853,8 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
                     fin_err += e_f[idx] * e_f[idx] * w;
                 }
             }
-            fin_err += hex_spectral_penalty(e_f, QK_K);
-
+            /* Floor is reconstruction SSE only. Spectral must not replace a
+             * better W with a worse candidate just to zero DC. */
             float g_best = candidate_errors[blk][0];
             int   g_cand = 0;
             for (int c = 1; c < TOTAL_SCALE_CANDIDATES; c++) {
@@ -4748,11 +4911,53 @@ static void quantize_tensor_q2k_hpc(const float *weights, int64_t n_elements,
         total_err += berr;
     }
 
+    /* ── PHASE 4.8: sequential TRUE residual carry ──────────────────────
+     * Phase 3.9 carry is a nearest-round estimate on the seed (d,dmin);
+     * Phase 4 then rewrites codes in parallel, so that carry is stale.
+     * Walk each row in order, measure the encoded Σe, and spend a tiny
+     * SSE budget on (d,dmin) + q±1 to hit Σe ≈ −decay·R_prev. */
+    if (HEX_DC_LAMBDA > 0.0f || DC_DECAY > 0.0f) {
+        int64_t bpr = (row_width > 0 && row_width % QK_K == 0)
+                      ? row_width / QK_K : 0;
+        float rolling_dc = 0.0f;
+        uint8_t L8[QK_K];
+        float w256[QK_K];
+        for (int64_t blk = 0; blk < n_blocks; blk++) {
+            if (bpr > 0 && (blk % bpr) == 0)
+                rolling_dc = 0.0f;
+            const float *bx = weights + blk * QK_K;
+            float carry = DC_DECAY * rolling_dc;
+            hex_q2k_unpack_L(&output[blk], L8);
+            float dm8 = gguf_fp16_to_fp32(output[blk].d);
+            float mm8 = gguf_fp16_to_fp32(output[blk].dmin);
+            for (int i = 0; i < QK_K; i++)
+                w256[i] = hex_q2k_el_w(imat_importance, blk, i);
+            if (hex_q2k_hit_dc_carry(bx, L8, output[blk].scales, &dm8, &mm8,
+                                     carry, w256)) {
+                output[blk].d    = gguf_fp32_to_fp16(dm8);
+                output[blk].dmin = gguf_fp32_to_fp16(mm8);
+            }
+            hex_q2k_dc_nudge_codes(bx, L8, output[blk].scales, dm8, mm8,
+                                   carry, w256);
+            hex_q2k_pack_L(&output[blk], L8);
+
+            float deq[QK_K];
+            gguf_dequantize_q2_k_block(&output[blk], deq);
+            float dc_res = 0.0f;
+            for (int i = 0; i < QK_K; i++)
+                dc_res += bx[i] - deq[i];
+            rolling_dc = dc_res;
+        }
+        total_err = 0.0f;
+        for (int64_t blk = 0; blk < n_blocks; blk++)
+            total_err += gguf_q2_k_block_error(weights + blk * QK_K, &output[blk]);
+    }
+
     for (int _ti = 0; _ti < _n_omp_threads; _ti++)
         hpc_destroy(_tl_graphs[_ti]);
     free(_tl_graphs);
 
-    free(block_dc_bias);
+    free(block_dc_carry);
     free(seeds);
     free(candidate_errors);
     free(best_candidate);
@@ -5336,6 +5541,418 @@ write_fail:
     return -1;
 }
 
+/* ═══════════════════════════════════════════════════════════════════════════
+ * IQ2_XS — E8-CODEBOOK QUANTIZER WITH FOLD/DC SHAPING (2.3125 bpw)
+ *
+ * Each 8-weight group is one of 512 codewords (magnitudes {8,25,43}) times
+ * a sign pattern with EVEN parity (7 bits stored, 8th = parity). Sub-block
+ * (16) scale db = d·(ls+0.5)/4, ls ∈ 0..15, d fp16 per 256-block.
+ *
+ * Why the codebook is where fold finally pays: for a given group there are
+ * several codewords within a hair of the nearest one (E8 shells are dense),
+ * so residual shaping — Σe ≈ −carry across blocks, e_i + e_{i+128} small —
+ * can pick among them at ~zero SSE cost. Q2_K only had ±1 scalar steps.
+ *
+ * Over ggml's encoder: exact grid magnitudes in the objective, sign-parity
+ * flip chosen jointly with the codeword, codewords re-picked after the
+ * 4-bit scale quantisation, d candidate search, ls ±1 descent, and the
+ * rolling residual carry. Reconstruction target is always the true x.
+ * ═══════════════════════════════════════════════════════════════════════════ */
+
+#include "iq2xs_grid.h"
+
+#define IQ2XS_NGRID   512
+#define IQ2XS_TOPK    8
+#define IQ2XS_NGROUP  (QK_K / 8)       /* 32 */
+#define IQ2XS_NSUB    (QK_K / 16)      /* 16 */
+
+static float g_iq2xs_gridf[IQ2XS_NGRID][8];
+static int   g_iq2xs_grid_ready = 0;
+
+static void iq2xs_prepare_grid(void)
+{
+    if (g_iq2xs_grid_ready) return;
+    for (int k = 0; k < IQ2XS_NGRID; k++)
+        for (int j = 0; j < 8; j++)
+            g_iq2xs_gridf[k][j] = (float)((iq2xs_grid[k] >> (8 * j)) & 0xFF);
+    g_iq2xs_grid_ready = 1;
+}
+
+typedef struct {
+    uint16_t code;      /* grid | signs7 << 9 */
+    float    sse;       /* Σ w (x − deq)² */
+    float    deq[8];
+} IQ2Cand;
+
+/* Top-K codewords for 8 weights at magnitude scale db (db > 0).
+ * Sign parity is enforced by the cheapest flip *per codeword*. */
+static int iq2xs_group_candidates(const float *x, const float *w, float db,
+                                  int K, IQ2Cand *out)
+{
+    float   ax[8], wa[8];
+    uint8_t s = 0; int par = 0;
+    for (int i = 0; i < 8; i++) {
+        ax[i] = fabsf(x[i]);
+        wa[i] = w[i];
+        if (x[i] < 0.0f) { s |= (uint8_t)(1u << i); par ^= 1; }
+    }
+    int n = 0;
+    for (int k = 0; k < IQ2XS_NGRID; k++) {
+        const float *g = g_iq2xs_gridf[k];
+        float e0 = 0.0f;
+        for (int i = 0; i < 8; i++) {
+            float d = ax[i] - db * g[i];
+            e0 += wa[i] * d * d;
+        }
+        /* Parity: odd sign count needs one flip. Offer the two cheapest
+         * flips as separate candidates — flipping a different weight is
+         * the cheapest DC lever this offset-free format has. */
+        int   flips[2] = { -1, -1 };
+        float costs[2] = { 0.0f, 0.0f };
+        int   nf = 1;
+        if (par) {
+            float c1 = 1e30f, c2 = 1e30f; int f1 = -1, f2 = -1;
+            for (int i = 0; i < 8; i++) {
+                float c = 4.0f * wa[i] * ax[i] * db * g[i];
+                if (c < c1)      { c2 = c1; f2 = f1; c1 = c; f1 = i; }
+                else if (c < c2) { c2 = c;  f2 = i; }
+            }
+            flips[0] = f1; costs[0] = c1;
+            flips[1] = f2; costs[1] = c2;
+            nf = (K > 1 && f2 >= 0) ? 2 : 1;
+        }
+        for (int f = 0; f < nf; f++) {
+            float e = e0 + costs[f];
+            if (n == K && e >= out[K - 1].sse) continue;
+            int pos = n < K ? n : K - 1;
+            while (pos > 0 && out[pos - 1].sse > e) { out[pos] = out[pos - 1]; pos--; }
+            uint8_t sf = s;
+            if (flips[f] >= 0) sf ^= (uint8_t)(1u << flips[f]);
+            out[pos].code = (uint16_t)(k | ((sf & 127) << 9));
+            out[pos].sse  = e;
+            for (int i = 0; i < 8; i++)
+                out[pos].deq[i] = db * g[i] * ((sf >> i) & 1 ? -1.0f : 1.0f);
+            if (n < K) n++;
+        }
+    }
+    return n;
+}
+
+/* Decode one group from its 16-bit code at scale db. */
+static inline void iq2xs_decode_group(uint16_t code, float db, float *deq)
+{
+    const float *g = g_iq2xs_gridf[code & 511];
+    uint8_t signs = ksigns_iq2xs[code >> 9];
+    for (int j = 0; j < 8; j++)
+        deq[j] = db * g[j] * ((signs >> j) & 1 ? -1.0f : 1.0f);
+}
+
+/* Best codes for a 16-weight sub-block at fixed db; returns weighted SSE. */
+static float iq2xs_sub_pick(const float *x, const float *w, float db,
+                            uint16_t code[2], float deq[16])
+{
+    if (db <= 0.0f) {
+        code[0] = code[1] = 0;
+        float e = 0.0f;
+        for (int i = 0; i < 16; i++) { deq[i] = 0.0f; e += w[i] * x[i] * x[i]; }
+        return e;
+    }
+    IQ2Cand c;
+    float e = 0.0f;
+    for (int k = 0; k < 2; k++) {
+        iq2xs_group_candidates(x + 8 * k, w + 8 * k, db, 1, &c);
+        code[k] = c.code;
+        memcpy(deq + 8 * k, c.deq, sizeof(c.deq));
+        e += c.sse;
+    }
+    return e;
+}
+
+/* Float scale search for one sub-block: candidate db grid + LS refit. */
+static float iq2xs_sub_fit(const float *x, const float *w, float *db_out)
+{
+    float amax = 0.0f;
+    for (int i = 0; i < 16; i++) amax = fmaxf(amax, fabsf(x[i]));
+    if (amax < 1e-12f) { *db_out = 0.0f; return 0.0f; }
+
+    float best_e = 1e30f, best_db = amax / 43.0f;
+    uint16_t code[2]; float deq[16];
+    for (int is = -10; is <= 10; is++) {
+        /* amax lands on grid value 43·(1+0.035·is): includes clipped maxima */
+        float db = amax / (43.0f * (1.0f + 0.035f * (float)is));
+        float e  = iq2xs_sub_pick(x, w, db, code, deq);
+        /* LS refit of db with codes fixed: deq = db·ĝ */
+        double num = 0.0, den = 0.0;
+        for (int i = 0; i < 16; i++) {
+            double gh = deq[i] / db;
+            num += (double)w[i] * x[i] * gh;
+            den += (double)w[i] * gh * gh;
+        }
+        if (den > 0.0 && num > 0.0) {
+            float db2 = (float)(num / den);
+            float e2  = iq2xs_sub_pick(x, w, db2, code, deq);
+            if (e2 < e) { e = e2; db = db2; }
+        }
+        if (e < best_e) { best_e = e; best_db = db; }
+    }
+    *db_out = best_db;
+    return best_e;
+}
+
+/* Whole-block encode at a given d: ls from float sub-scales, re-pick codes.
+ * Returns weighted SSE. */
+static float iq2xs_block_at_d(const float *x, const float *w, float d,
+                              const float *db_f, uint8_t ls[IQ2XS_NSUB],
+                              uint16_t code[IQ2XS_NGROUP], float deq[QK_K])
+{
+    float e = 0.0f;
+    for (int ib = 0; ib < IQ2XS_NSUB; ib++) {
+        int l = (d > 0.0f) ? gguf_nearest_int(db_f[ib] * 4.0f / d - 0.5f) : 0;
+        if (l < 0) l = 0; if (l > 15) l = 15;
+        ls[ib] = (uint8_t)l;
+        float db = d * ((float)l + 0.5f) * 0.25f;
+        e += iq2xs_sub_pick(x + 16 * ib, w + 16 * ib, db, code + 2 * ib, deq + 16 * ib);
+    }
+    return e;
+}
+
+static inline float iq2xs_sub_db(float d, uint8_t ls)
+{
+    return d * ((float)ls + 0.5f) * 0.25f;
+}
+
+/* Greedy re-selection among top-K codewords per group on
+ *   SSE + (λ_dc/n)(Σe + carry)² + (λ_vw/n) Σ_p (e_p + e_{p+128})²
+ * with a block SSE cap. This is the fold-through-codebook step. */
+static void iq2xs_shape_block(const float *x, const float *w, float d,
+                              const uint8_t ls[IQ2XS_NSUB],
+                              uint16_t code[IQ2XS_NGROUP], float dc_carry)
+{
+    if (HEX_DC_LAMBDA == 0.0f && HEX_VW_LAMBDA == 0.0f) return;
+
+    IQ2Cand cands[IQ2XS_NGROUP][IQ2XS_TOPK];
+    int   ncand[IQ2XS_NGROUP], cur[IQ2XS_NGROUP];
+    float e[QK_K];
+    float sse = 0.0f, dc = dc_carry;
+
+    for (int g = 0; g < IQ2XS_NGROUP; g++) {
+        float db = iq2xs_sub_db(d, ls[g >> 1]);
+        if (db <= 0.0f) { ncand[g] = 0; cur[g] = -1;
+            for (int j = 0; j < 8; j++) { e[8*g+j] = x[8*g+j]; sse += w[8*g+j]*x[8*g+j]*x[8*g+j]; dc += e[8*g+j]; }
+            continue; }
+        ncand[g] = iq2xs_group_candidates(x + 8*g, w + 8*g, db, IQ2XS_TOPK, cands[g]);
+        cur[g] = 0;
+        for (int c = 0; c < ncand[g]; c++)
+            if (cands[g][c].code == code[g]) { cur[g] = c; break; }
+        const IQ2Cand *cc = &cands[g][cur[g]];
+        for (int j = 0; j < 8; j++) { e[8*g+j] = x[8*g+j] - cc->deq[j]; dc += e[8*g+j]; }
+        sse += cc->sse;
+    }
+    const float cap = sse * (1.0f + HEX_DC_SSE_BUDGET);
+    float v[QK_K / 2], ves = 0.0f;
+    for (int p = 0; p < QK_K / 2; p++) { v[p] = e[p] + e[p + QK_K/2]; ves += v[p] * v[p]; }
+    const float ldc = HEX_DC_LAMBDA / (float)QK_K, lvw = HEX_VW_LAMBDA / (float)QK_K;
+    float metric = sse + ldc * dc * dc + lvw * ves;
+
+    for (int pass = 0; pass < 96; pass++) {
+        int best_g = -1, best_c = 0; float best_m = metric;
+        for (int g = 0; g < IQ2XS_NGROUP; g++) {
+            if (cur[g] < 0) continue;
+            const IQ2Cand *co = &cands[g][cur[g]];
+            for (int c = 0; c < ncand[g]; c++) {
+                if (c == cur[g]) continue;
+                const IQ2Cand *cn = &cands[g][c];
+                float sse2 = sse - co->sse + cn->sse;
+                if (sse2 > cap) continue;
+                float dc2 = dc, ves2 = ves;
+                for (int j = 0; j < 8; j++) {
+                    int i  = 8*g + j;
+                    float de = co->deq[j] - cn->deq[j];     /* e_new − e_old */
+                    dc2 += de;
+                    int p = (i < QK_K/2) ? i : i - QK_K/2;
+                    float vn = v[p] + de;
+                    ves2 += vn * vn - v[p] * v[p];
+                }
+                float m2 = sse2 + ldc * dc2 * dc2 + lvw * ves2;
+                if (m2 < best_m) { best_m = m2; best_g = g; best_c = c; }
+            }
+        }
+        if (best_g < 0) break;
+        const IQ2Cand *co = &cands[best_g][cur[best_g]];
+        const IQ2Cand *cn = &cands[best_g][best_c];
+        sse += cn->sse - co->sse;
+        for (int j = 0; j < 8; j++) {
+            int i = 8*best_g + j;
+            float de = co->deq[j] - cn->deq[j];
+            e[i] += de; dc += de;
+            int p = (i < QK_K/2) ? i : i - QK_K/2;
+            ves += (v[p] + de) * (v[p] + de) - v[p] * v[p];
+            v[p] += de;
+        }
+        cur[best_g] = best_c;
+        code[best_g] = cn->code;
+        metric = best_m;
+    }
+}
+
+static void iq2xs_pack(BlockIQ2XS *b, float d, const uint8_t ls[IQ2XS_NSUB],
+                       const uint16_t code[IQ2XS_NGROUP])
+{
+    b->d = gguf_fp32_to_fp16(d);
+    for (int ib = 0; ib < IQ2XS_NSUB; ib += 2)
+        b->scales[ib / 2] = (uint8_t)(ls[ib] | (ls[ib + 1] << 4));
+    memcpy(b->qs, code, sizeof(uint16_t) * IQ2XS_NGROUP);
+}
+
+static void iq2xs_dequant_block(const BlockIQ2XS *b, float *out)
+{
+    iq2xs_prepare_grid();
+    float d = gguf_fp16_to_fp32(b->d);
+    for (int g = 0; g < IQ2XS_NGROUP; g++) {
+        uint8_t ls = (g & 2) ? (b->scales[g >> 2] >> 4) : (b->scales[g >> 2] & 0xF);
+        iq2xs_decode_group(b->qs[g], iq2xs_sub_db(d, ls), out + 8 * g);
+    }
+}
+
+static void quantize_tensor_iq2_xs_hpc(const float *weights, int64_t n_elements,
+                                       BlockIQ2XS *output, float *out_total_error,
+                                       const float *imat_importance, int verbose,
+                                       int64_t row_width)
+{
+    if (!weights || !output || n_elements <= 0 || n_elements % QK_K != 0) {
+        if (out_total_error) *out_total_error = -1.0f;
+        return;
+    }
+    iq2xs_prepare_grid();
+    const int64_t n_blocks = n_elements / QK_K;
+    static const float d_mult[] = { 1.0f, 0.97f, 1.03f, 0.94f, 1.06f, 0.90f, 1.10f, 0.85f, 1.15f };
+    const int n_dm = (int)(sizeof(d_mult) / sizeof(d_mult[0]));
+
+    /* Experimental knobs, both OFF by default — measured on a controlled
+     * splice A/B (SmolLM2 ffn_down, 64×512-token PPL, imatrix = E[a²]):
+     *  HEX_IQ2_WMODE=1   ggml's w = imat·sqrt(σ²_blk + x²)  → PPL +6%  (worse)
+     *  HEX_IQ2_INFLATE=α scale d by (1+α) after the fit     → PPL +1..4% (worse)
+     * Plain imatrix-weighted SSE with the exact grid is the best objective. */
+    const char *wm = getenv("HEX_IQ2_WMODE");
+    const int wmode = wm ? atoi(wm) : 0;
+    const char *inf = getenv("HEX_IQ2_INFLATE");
+    const float inflate = inf ? (float)atof(inf) : 0.0f;
+
+    #pragma omp parallel for schedule(dynamic, 16)
+    for (int64_t blk = 0; blk < n_blocks; blk++) {
+        const float *x = weights + blk * QK_K;
+        float w[QK_K];
+        float sigma2 = 0.0f;
+        for (int i = 0; i < QK_K; i++) sigma2 += x[i] * x[i];
+        sigma2 /= (float)QK_K;
+        for (int i = 0; i < QK_K; i++) {
+            float base = imat_importance ? imat_importance[blk * QK_K + i] : 1.0f;
+            w[i] = (wmode == 1) ? base * sqrtf(sigma2 + x[i] * x[i]) : base;
+        }
+
+        /* 1. float sub-block scales */
+        float db_f[IQ2XS_NSUB], db_max = 0.0f;
+        for (int ib = 0; ib < IQ2XS_NSUB; ib++) {
+            iq2xs_sub_fit(x + 16 * ib, w + 16 * ib, &db_f[ib]);
+            db_max = fmaxf(db_max, db_f[ib]);
+        }
+        if (db_max <= 0.0f) { memset(&output[blk], 0, sizeof(BlockIQ2XS)); continue; }
+
+        /* 2. d candidate search (fp16-exact), codes re-picked at quantised db */
+        float    d0 = db_max * 4.0f / 15.5f;
+        float    best_e = 1e30f, best_d = d0;
+        uint8_t  ls[IQ2XS_NSUB], ls_t[IQ2XS_NSUB];
+        uint16_t code[IQ2XS_NGROUP], code_t[IQ2XS_NGROUP];
+        float    deq[QK_K];
+        for (int c = 0; c < n_dm; c++) {
+            float d = gguf_fp16_to_fp32(gguf_fp32_to_fp16(d0 * d_mult[c]));
+            if (d <= 0.0f) continue;
+            float e = iq2xs_block_at_d(x, w, d, db_f, ls_t, code_t, deq);
+            if (e < best_e) { best_e = e; best_d = d;
+                memcpy(ls, ls_t, sizeof(ls)); memcpy(code, code_t, sizeof(code)); }
+        }
+
+        /* 3. per-sub-block ls ±1 coordinate descent at fixed d */
+        float sub_e[IQ2XS_NSUB];
+        for (int ib = 0; ib < IQ2XS_NSUB; ib++)
+            sub_e[ib] = iq2xs_sub_pick(x + 16*ib, w + 16*ib, iq2xs_sub_db(best_d, ls[ib]),
+                                       code + 2*ib, deq + 16*ib);
+        for (int it = 0; it < 3; it++) {
+            int moved = 0;
+            for (int ib = 0; ib < IQ2XS_NSUB; ib++) {
+                for (int dl = -1; dl <= 1; dl += 2) {
+                    int l = (int)ls[ib] + dl;
+                    if (l < 0 || l > 15) continue;
+                    uint16_t ct[2]; float dq[16];
+                    float e = iq2xs_sub_pick(x + 16*ib, w + 16*ib,
+                                             iq2xs_sub_db(best_d, (uint8_t)l), ct, dq);
+                    if (e < sub_e[ib]) {
+                        sub_e[ib] = e; ls[ib] = (uint8_t)l;
+                        code[2*ib] = ct[0]; code[2*ib+1] = ct[1];
+                        memcpy(deq + 16*ib, dq, sizeof(dq));
+                        moved = 1;
+                    }
+                }
+            }
+            if (!moved) break;
+        }
+
+        /* 4. fold/DC shaping among near-equivalent codewords (carry = 0 here;
+         *    the sequential pass below applies the true residual carry). */
+        iq2xs_shape_block(x, w, best_d, ls, code, 0.0f);
+
+        float d_out = best_d;
+        if (inflate != 0.0f)
+            d_out = gguf_fp16_to_fp32(gguf_fp32_to_fp16(best_d * (1.0f + inflate)));
+        iq2xs_pack(&output[blk], d_out, ls, code);
+    }
+
+    /* 5. sequential TRUE residual carry along each row */
+    if (HEX_DC_LAMBDA > 0.0f && g_hex_dc_decay > 0.0f) {
+        int64_t bpr = (row_width > 0 && row_width % QK_K == 0) ? row_width / QK_K : 0;
+        float rolling = 0.0f, w[QK_K], deq[QK_K];
+        uint8_t ls[IQ2XS_NSUB]; uint16_t code[IQ2XS_NGROUP];
+        for (int64_t blk = 0; blk < n_blocks; blk++) {
+            if (bpr > 0 && (blk % bpr) == 0) rolling = 0.0f;
+            const float *x = weights + blk * QK_K;
+            BlockIQ2XS *b = &output[blk];
+            float d = gguf_fp16_to_fp32(b->d);
+            if (d > 0.0f) {
+                float sigma2 = 0.0f;
+                for (int i = 0; i < QK_K; i++) sigma2 += x[i] * x[i];
+                sigma2 /= (float)QK_K;
+                for (int i = 0; i < QK_K; i++) {
+                    float base = imat_importance ? imat_importance[blk * QK_K + i] : 1.0f;
+                    w[i] = (wmode == 1) ? base * sqrtf(sigma2 + x[i] * x[i]) : base;
+                }
+                for (int ib = 0; ib < IQ2XS_NSUB; ib++)
+                    ls[ib] = (ib & 1) ? (b->scales[ib >> 1] >> 4) : (b->scales[ib >> 1] & 0xF);
+                memcpy(code, b->qs, sizeof(code));
+                iq2xs_shape_block(x, w, d, ls, code, g_hex_dc_decay * rolling);
+                memcpy(b->qs, code, sizeof(code));
+            }
+            iq2xs_dequant_block(b, deq);
+            float r = 0.0f;
+            for (int i = 0; i < QK_K; i++) r += x[i] - deq[i];
+            rolling = r;
+        }
+    }
+
+    /* 6. exact reconstruction SSE */
+    double tot = 0.0;
+    #pragma omp parallel for reduction(+:tot)
+    for (int64_t blk = 0; blk < n_blocks; blk++) {
+        float deq[QK_K];
+        iq2xs_dequant_block(&output[blk], deq);
+        const float *x = weights + blk * QK_K;
+        for (int i = 0; i < QK_K; i++) { double e = x[i] - deq[i]; tot += e * e; }
+    }
+    if (out_total_error) *out_total_error = (float)tot;
+    if (verbose)
+        printf("  [IQ2_XS·Sieve] blocks=%lld rmse=%.4e\n", (long long)n_blocks,
+               sqrt(tot / (double)n_elements));
+}
+
 /* ═══════════════════════════════════════════════════════════════════════════
  * LIBRARY API — Exported functions for Python ctypes integration
  *
@@ -5442,6 +6059,27 @@ void hexstate_quantize_tensor_q8_0_hpc(const float *weights, int64_t n_elements,
                              imat_importance, verbose);
 }
 
+/* IQ2_XS (74 bytes / 256 weights). Row-aware residual carry like Q2_K. */
+int hexstate_iq2xs_block_bytes(void)    { return (int)sizeof(BlockIQ2XS); }
+int hexstate_iq2xs_block_elements(void) { return QK_K; }
+
+void hexstate_quantize_tensor_iq2_xs_hpc(const float *weights, int64_t n_elements,
+                                         void *output, float *out_error,
+                                         const float *imat_importance, int verbose,
+                                         int64_t row_width)
+{
+    hexstate_init();
+    quantize_tensor_iq2_xs_hpc(weights, n_elements, (BlockIQ2XS *)output,
+                               out_error, imat_importance, verbose, row_width);
+}
+
+void hexstate_dequant_iq2_xs(const void *blocks, int64_t n_blocks, float *out)
+{
+    const BlockIQ2XS *b = (const BlockIQ2XS *)blocks;
+    for (int64_t i = 0; i < n_blocks; i++)
+        iq2xs_dequant_block(&b[i], out + i * QK_K);
+}
+
 #ifndef HEXSTATE_LIBRARY
 /* ═══════════════════════════════════════════════════════════════════════════
  * MAIN

hexstate_requantize.py

@@ -114,8 +114,44 @@ def _load_hexstate_lib():
             lib.hexstate_set_spectral_params.restype = None
             lib.hexstate_set_spectral_params.argtypes = [
                 ctypes.c_float, ctypes.c_float, ctypes.c_float]
+        if hasattr(lib, 'hexstate_set_sse_budget'):
+            lib.hexstate_set_sse_budget.restype = None
+            lib.hexstate_set_sse_budget.argtypes = [ctypes.c_float]
+
+        # IQ2_XS (E8 codebook) quantizer + dequant
+        if hasattr(lib, 'hexstate_quantize_tensor_iq2_xs_hpc'):
+            lib.hexstate_quantize_tensor_iq2_xs_hpc.restype = None
+            lib.hexstate_quantize_tensor_iq2_xs_hpc.argtypes = [
+                ctypes.POINTER(ctypes.c_float),  # weights
+                ctypes.c_int64,                   # n_elements
+                ctypes.c_void_p,                  # output
+                ctypes.POINTER(ctypes.c_float),   # out_error
+                ctypes.POINTER(ctypes.c_float),   # imat_importance (can be NULL)
+                ctypes.c_int,                     # verbose
+                ctypes.c_int64,                   # row_width
+            ]
+            lib.hexstate_dequant_iq2_xs.restype = None
+            lib.hexstate_dequant_iq2_xs.argtypes = [
+                ctypes.c_void_p, ctypes.c_int64, ctypes.POINTER(ctypes.c_float)]
 
         lib.hexstate_init()
+        dc_l = os.environ.get('HEX_DC_LAMBDA')
+        vw_l = os.environ.get('HEX_VW_LAMBDA')
+        dc_d = os.environ.get('HEX_DC_DECAY')
+        if hasattr(lib, 'hexstate_set_spectral_params') and (dc_l or vw_l or dc_d):
+            lib.hexstate_set_spectral_params(
+                ctypes.c_float(float(dc_l) if dc_l else 1.0),
+                ctypes.c_float(float(vw_l) if vw_l else 1.0),
+                ctypes.c_float(float(dc_d) if dc_d else 0.85),
+            )
+        # Relative SSE the DC/vesica shaper may spend. Q2_K default 5e-4;
+        # IQ2_XS has no per-sub-block offset, so cancelling DC means swapping
+        # codewords — it needs ~2e-2 (≈ +0.2% RMSE) to be effective.
+        budget = os.environ.get('HEX_SSE_BUDGET')
+        if budget is None and LOWBIT_FORMAT == 'iq2xs':
+            budget = '0.02'
+        if budget is not None and hasattr(lib, 'hexstate_set_sse_budget'):
+            lib.hexstate_set_sse_budget(ctypes.c_float(float(budget)))
         _HEXSTATE_LIB = lib
         return lib
     except Exception as e:
@@ -198,12 +234,20 @@ def read_imatrix(path):
                     base = name[:-len('.counts')]
                     counts_data[base] = data
 
-            # Compute normalized importance: sqrt(in_sum2 / counts) / mean
+            # Importance = in_sum2 / counts = E[a²] per input column, i.e. the
+            # weight in Σ E[a_i²]·e_i² (output-error variance). This is exactly
+            # ggml's quant_weights and what the legacy .dat branch below returns.
+            # The earlier sqrt() here under-weighted important columns: in a
+            # controlled splice A/B (SmolLM2, ffn_down only) it cost +4.7% PPL
+            # for Q2_K and +7.8% for IQ2_XS. HEX_IMAT_SQRT=1 restores it.
+            use_sqrt = os.environ.get('HEX_IMAT_SQRT', '0') == '1'
             for base_name in sum2_data:
                 in_sum2 = sum2_data[base_name]
                 count = counts_data.get(base_name, np.array([1.0]))[0]
                 if count > 0:
-                    importance = np.sqrt(in_sum2 / count)
+                    importance = in_sum2 / count
+                    if use_sqrt:
+                        importance = np.sqrt(importance)
                 else:
                     importance = np.ones_like(in_sum2)
                 mean = importance.mean()
@@ -304,24 +348,31 @@ GGML_TYPE_F16   = 1
 GGML_TYPE_Q4_0  = 2
 GGML_TYPE_Q8_0  = 8
 GGML_TYPE_Q2_K  = 10
+GGML_TYPE_IQ2_XS = 17
 GGML_TYPE_BF16  = 30
 
+IQ2_XS_BLOCK_BYTES = 74   # d(fp16) + 32×u16 codes + 8 scale bytes
+
+# Low-bit target for the "Q2_K plan" tensors: 'q2k' (default) or 'iq2xs'
+# (--iq2xs: E8 codebook, 2.3125 bpw, native llama.cpp decode).
+LOWBIT_FORMAT = 'q2k'
+
 TYPE_NAME = {
     0: "F32", 1: "F16", 2: "Q4_0", 3: "Q4_1", 6: "Q5_0", 7: "Q5_1",
     8: "Q8_0", 9: "Q8_1", 10: "Q2_K", 11: "Q3_K", 12: "Q4_K",
-    13: "Q5_K", 14: "Q6_K", 15: "Q8_K", 30: "BF16",
+    13: "Q5_K", 14: "Q6_K", 15: "Q8_K", 17: "IQ2_XS", 30: "BF16",
 }
 
 # Block sizes and byte sizes for each type
 TYPE_BLOCK_SIZE = {
     0: 1, 1: 1, 2: 32, 3: 32, 6: 32, 7: 32,
     8: 32, 9: 32, 10: 256, 11: 256, 12: 256,
-    13: 256, 14: 256, 15: 256, 30: 1,
+    13: 256, 14: 256, 15: 256, 17: 256, 30: 1,
 }
 TYPE_BLOCK_BYTES = {
     0: 4, 1: 2, 2: 18, 3: 20, 6: 20, 7: 22,
     8: 34, 9: 36, 10: 84, 11: 110, 12: 144,
-    13: 176, 14: 210, 15: 292, 30: 2,
+    13: 176, 14: 210, 15: 292, 17: IQ2_XS_BLOCK_BYTES, 30: 2,
 }
 
 
@@ -784,8 +835,42 @@ def _copy_bytes(fin, fout, abs_offset, n_bytes):
     return written
 
 
+def quantize_tensor_iq2xs_hpc(f32_data, importance=None, row_width=0):
+    """IQ2_XS (E8 codebook, 74 B / 256 w) via the HPC C library.
+    Returns (bytes, n_blocks, dequant_f32)."""
+    lib = _load_hexstate_lib()
+    if lib is None or not hasattr(lib, 'hexstate_quantize_tensor_iq2_xs_hpc'):
+        raise RuntimeError('libhexstate_q2k.so lacks IQ2_XS support — rebuild')
+    f32 = np.ascontiguousarray(f32_data, dtype=np.float32).reshape(-1)
+    n = int(f32.size)
+    if n % QK_K != 0:
+        raise ValueError(f'IQ2_XS needs a multiple of {QK_K} elements, got {n}')
+    n_blocks = n // QK_K
+    out = np.zeros(n_blocks * IQ2_XS_BLOCK_BYTES, dtype=np.uint8)
+    err = ctypes.c_float(0.0)
+    imat_ptr = None
+    if importance is not None:
+        imat_c = np.ascontiguousarray(importance, dtype=np.float32).reshape(-1)
+        if imat_c.size == n:
+            imat_ptr = imat_c.ctypes.data_as(ctypes.POINTER(ctypes.c_float))
+    lib.hexstate_quantize_tensor_iq2_xs_hpc(
+        f32.ctypes.data_as(ctypes.POINTER(ctypes.c_float)),
+        ctypes.c_int64(n),
+        out.ctypes.data_as(ctypes.c_void_p),
+        ctypes.byref(err),
+        imat_ptr,
+        ctypes.c_int(0),
+        ctypes.c_int64(int(row_width)),
+    )
+    deq = np.zeros(n, dtype=np.float32)
+    lib.hexstate_dequant_iq2_xs(
+        out.ctypes.data_as(ctypes.c_void_p), ctypes.c_int64(n_blocks),
+        deq.ctypes.data_as(ctypes.POINTER(ctypes.c_float)))
+    return out.tobytes(), n_blocks, deq
+
+
 def _stream_quantize(fin, fout, ti, abs_offset, kind, imatrix_data, use_hpc):
-    """Quantize one tensor in row chunks. kind: 'q2k' | 'q4' | 'q8'.
+    """Quantize one tensor in row chunks. kind: 'q2k' | 'iq2xs' | 'q4' | 'q8'.
     Returns (n_out_bytes, rmse_or_None, sigma_or_None).
     """
     ttype = ti['type']
@@ -798,7 +883,7 @@ def _stream_quantize(fin, fout, ti, abs_offset, kind, imatrix_data, use_hpc):
 
     d0 = int(ti['dims'][0])
     n_rows = int(ti['n_elements']) // d0
-    align = QK_K if kind == 'q2k' else 32
+    align = QK_K if kind in ('q2k', 'iq2xs') else 32
     if d0 % align != 0:
         raise ValueError(f'{ti["name"]} dim0={d0} not aligned to {align}')
 
@@ -816,7 +901,15 @@ def _stream_quantize(fin, fout, ti, abs_offset, kind, imatrix_data, use_hpc):
         total_ss += float(np.vdot(f32, f32))
         total_n += n_valid
 
-        if kind == 'q2k':
+        if kind == 'iq2xs':
+            qbytes, n_blocks, deq = quantize_tensor_iq2xs_hpc(
+                f32, importance=imp, row_width=d0)
+            fout.write(qbytes)
+            written += len(qbytes)
+            diff = f32.reshape(-1)[:n_valid] - deq[:n_valid]
+            total_se += float(np.sum(diff ** 2))
+            del qbytes, deq
+        elif kind == 'q2k':
             if use_hpc:
                 qbytes, n_blocks = quantize_tensor_q2k_hpc(
                     f32, opt_mode=2, importance=imp, row_width=d0)
@@ -981,16 +1074,29 @@ def should_quantize(name, n_dims, dims, tied_embeddings=False):
 def main():
     if len(sys.argv) < 3:
         print("Usage: python3 hexstate_requantize.py <input.gguf> <output.gguf>"
-              " [--keep-metadata] [--imatrix FILE] [--keep-embd] [--q2all]")
+              " [--keep-metadata] [--imatrix FILE] [--keep-embd] [--q2all] [--iq2xs]")
+        print("  --iq2xs          low-bit tensors → IQ2_XS (E8 codebook, 2.3125 bpw) instead of Q2_K")
         print("  HEX_CHUNK_ELEMS  max f32 elements per tensor chunk (default 2000000)")
+        print("  HEX_DC_LAMBDA    DC residual weight (default 1)")
+        print("  HEX_VW_LAMBDA    vesica weight (default 1)")
+        print("  HEX_DC_DECAY     rolling residual carry 0..1 (default 0.85)")
+        print("  HEX_SSE_BUDGET   relative SSE the shaper may spend (Q2_K 5e-4, IQ2_XS 2e-2)")
         sys.exit(1)
 
+    global LOWBIT_FORMAT
     input_path = sys.argv[1]
     output_path = sys.argv[2]
     keep_metadata = '--keep-metadata' in sys.argv
     quantize_none = '--quantize-none' in sys.argv
     q2all = '--q2all' in sys.argv
     keep_embd = '--keep-embd' in sys.argv   # keep tied embedding at source precision instead of Q8_0
+    if '--iq2xs' in sys.argv:
+        LOWBIT_FORMAT = 'iq2xs'
+    lowbit_kind = LOWBIT_FORMAT                       # 'q2k' | 'iq2xs'
+    lowbit_name = 'IQ2_XS' if lowbit_kind == 'iq2xs' else 'Q2_K'
+    lowbit_type = GGML_TYPE_IQ2_XS if lowbit_kind == 'iq2xs' else GGML_TYPE_Q2_K
+    lowbit_bytes = IQ2_XS_BLOCK_BYTES if lowbit_kind == 'iq2xs' else 84
+    lowbit_file_type = 20 if lowbit_kind == 'iq2xs' else 10   # LLAMA_FTYPE_MOSTLY_*
 
     # Check for imatrix
     imatrix_data = None
@@ -1006,11 +1112,19 @@ def main():
 
     # Check for HPC C library
     use_hpc = _load_hexstate_lib() is not None
+    if lowbit_kind == 'iq2xs':
+        lib = _load_hexstate_lib()
+        if lib is None or not hasattr(lib, 'hexstate_quantize_tensor_iq2_xs_hpc'):
+            print("  ERROR: --iq2xs needs libhexstate_q2k.so with IQ2_XS support "
+                  "(make -f makefile.quantize.c)")
+            sys.exit(1)
 
     print()
     print("  ╔════════════════════════════════════════════════════════════════��")
     print("  ║  HExState GGUF Re-Quantizer                                  ║")
-    print("  ║  GGUF → Q2_K GGUF with metadata passthrough                  ║")
+    print(f"  ║  GGUF → {lowbit_name:6s} GGUF with metadata passthrough                ║")
+    if lowbit_kind == 'iq2xs':
+        print("  ║  Low-bit: IQ2_XS  E8 codebook · fold/DC shaping · 2.3125 bpw ║")
     if q2all:
         print("  ║  Mode: --q2all  ALL eligible tensors → Q2_K (test mode)      ║")
     if use_hpc and imatrix_data:
@@ -1178,9 +1292,9 @@ def main():
                     out_size = n_blocks * 18
                     print(f"  Q4_0: {ti['name']} (dims[0]={dim0})")
                 elif quant_plan[i] is True and q2k_row_compatible(ti['n_dims'], ti['dims']):
-                    out_type = GGML_TYPE_Q2_K
+                    out_type = lowbit_type
                     n_blocks = ti['n_elements'] // QK_K
-                    out_size = n_blocks * 84
+                    out_size = n_blocks * lowbit_bytes
                 else:
                     out_type = ti['type']
                     out_size = ti['data_size']
@@ -1210,8 +1324,8 @@ def main():
         else:
             for key, vtype, raw_value in kv_pairs:
                 if key == 'general.file_type' and vtype == 4:  # UINT32
-                    # file_type=10 means Q2_K in llama.cpp
-                    updated_kv.append((key, vtype, struct.pack('<I', 10)))
+                    # LLAMA_FTYPE_MOSTLY_Q2_K = 10, MOSTLY_IQ2_XS = 20
+                    updated_kv.append((key, vtype, struct.pack('<I', lowbit_file_type)))
                 elif key == 'general.quantization_version' and vtype == 4:
                     updated_kv.append((key, vtype, struct.pack('<I', 2)))
                 elif key == 'tokenizer.ggml.token_type' and vtype == 9:
@@ -1369,14 +1483,14 @@ def main():
 
                 elif plan:
                     nbytes, rmse, sigma = _stream_quantize(
-                        fin, fout, ti, abs_offset, 'q2k', imatrix_data, use_hpc)
+                        fin, fout, ti, abs_offset, lowbit_kind, imatrix_data, use_hpc)
                     if rmse is not None:
                         q2k_rmse_sum += rmse
                         q2k_tensor_count += 1
-                        print(f"\n  [Q2_K] {ti['name'][:50]}  RMSE={rmse:.6e}"
+                        print(f"\n  [{lowbit_name}] {ti['name'][:50]}  RMSE={rmse:.6e}"
                               f"  σ={sigma:.4f}  rel={rmse / max(sigma, 1e-30):.4f}")
                     else:
-                        print(f"\n  [Q2_K] {ti['name'][:55]}  RMSE=n/a")
+                        print(f"\n  [{lowbit_name}] {ti['name'][:55]}  RMSE=n/a")
                     quant_count += 1
                     total_quant_bytes += nbytes
 
@@ -1402,16 +1516,16 @@ def main():
     print("  ╔════════════════════════════════════════════════════════════════╗")
     print("  ║  RE-QUANTIZATION SUMMARY                                     ║")
     print("  ╠════════════════════════════════════════════════════════════════╣")
-    print(f"  ║  Tensors quantized (Q2_K): {quant_count:<33d} ║")
+    print(f"  ║  Tensors quantized ({lowbit_name:6s}): {quant_count:<31d} ║")
     print(f"  ║  Tensors kept as-is:       {total_keep:<33d} ║")
-    print(f"  ║  Q2_K data:         {total_quant_bytes:>12,} bytes ({total_quant_bytes/1024**2:>7.1f} MB) ║")
+    print(f"  ║  {lowbit_name:6s} data:       {total_quant_bytes:>12,} bytes ({total_quant_bytes/1024**2:>7.1f} MB) ║")
     print(f"  ║  Kept data:         {total_keep_bytes:>12,} bytes ({total_keep_bytes/1024**2:>7.1f} MB) ║")
     print(f"  ║  Original size:     {file_size:>12,} bytes ({file_size/1024**3:>7.2f} GB) ║")
     print(f"  ║  Output size:       {final_size:>12,} bytes ({final_size/1024**3:>7.2f} GB) ║")
     print(f"  ║  Compression:       {compression:>42.1f}x ║")
     if q2k_tensor_count > 0:
         mean_rmse = q2k_rmse_sum / q2k_tensor_count
-        print(f"  ║  Mean Q2_K RMSE:                            {mean_rmse:>12.6e} ║")
+        print(f"  ║  Mean {lowbit_name:6s} RMSE:                          {mean_rmse:>12.6e} ║")
     print(f"  ║  Total time:        {elapsed:>39.1f} sec ║")
     print("  ╚════════════════════════════════════════════════════════════════╝")
     print()