CompressedGemma commited on
Commit
2e643f2
·
verified ·
1 Parent(s): 82b0808

Exponential speed-up

Browse files
Files changed (1) hide show
  1. hexstate_quantize.c +137 -80
hexstate_quantize.c CHANGED
@@ -5466,7 +5466,7 @@ static double hex_quad(const float *e, const float *H, int64_t n)
5466
 
5467
  /* Optimise one output row of length n_in (a whole number of Q2_K blocks). */
5468
  static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5469
- BlockQ2K *out, int n_sweeps,
5470
  float *E, float *Gd, uint8_t *q,
5471
  uint8_t *Ls, uint8_t *Lm, float *dv, float *mv)
5472
  {
@@ -5519,11 +5519,15 @@ static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5519
  float ds = dv[b] * (float)Ls[j], ms = mv[b] * (float)Lm[j];
5520
  E[i] = w[i] - (ds * (float)q[i] - ms);
5521
  }
5522
- for (int64_t i = 0; i < n_in; i++) {
5523
- const float *Hi = H + i * n_in;
5524
- double s = 0.0;
5525
- for (int64_t k = 0; k < n_in; k++) s += (double)Hi[k] * (double)E[k];
5526
- Gd[i] = (float)s;
 
 
 
 
5527
  }
5528
 
5529
  for (int sweep = 0; sweep < n_sweeps; sweep++) {
@@ -5550,71 +5554,83 @@ static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5550
  q[j] = (uint8_t)bestq;
5551
  E[j] -= bestd;
5552
  changed++;
5553
- if (HEX_HESS_FLUSH <= 1) {
 
 
 
5554
  const float *Hj = H + j * n_in;
5555
  for (int64_t k = 0; k < n_in; k++) Gd[k] -= bestd * Hj[k];
5556
- } else {
5557
- pend++;
5558
- }
5559
- }
5560
- if (HEX_HESS_FLUSH > 1 && pend && ((j + 1) % HEX_HESS_FLUSH == 0)) {
5561
- for (int64_t i = 0; i < n_in; i++) {
5562
- const float *Hi = H + i * n_in;
5563
- double s = 0.0;
5564
- for (int64_t k = 0; k < n_in; k++) s += (double)Hi[k] * (double)E[k];
5565
- Gd[i] = (float)s;
5566
  }
5567
- pend = 0;
5568
- }
5569
- }
5570
- if (HEX_HESS_FLUSH > 1) {
5571
- for (int64_t i = 0; i < n_in; i++) {
5572
- const float *Hi = H + i * n_in;
5573
- double s = 0.0;
5574
- for (int64_t k = 0; k < n_in; k++) s += (double)Hi[k] * (double)E[k];
5575
- Gd[i] = (float)s;
5576
  }
5577
  }
 
5578
  }
5579
 
5580
  /* ── move B: joint (ls, lm) per sub-block, full 16x16 ── */
5581
  for (int64_t sj = 0; sj < n_in / 16; sj++) {
5582
  int64_t b = (sj * 16) / QK_K, c0 = sj * 16;
5583
  float dd = dv[b], mm = mv[b];
5584
- float old[16], del[16], bestdel[16];
5585
- for (int k = 0; k < 16; k++)
5586
- old[k] = dd * (float)Ls[sj] * (float)q[c0 + k] - mm * (float)Lm[sj];
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
5587
  float bestg = 0.0f; int bl = Ls[sj], bm = Lm[sj]; int found = 0;
 
5588
  for (int tl = 0; tl <= 15; tl++) {
 
5589
  for (int tm = 0; tm <= 15; tm++) {
5590
- float lin = 0.0f, quad = 0.0f;
5591
- for (int k = 0; k < 16; k++) {
5592
- float nw = dd * (float)tl * (float)q[c0 + k] - mm * (float)tm;
5593
- del[k] = nw - old[k];
5594
- lin += del[k] * Gd[c0 + k];
5595
- }
5596
- for (int a = 0; a < 16; a++) {
5597
- if (del[a] == 0.0f) continue;
5598
- const float *Ha = H + (c0 + a) * n_in + c0;
5599
- float s = 0.0f;
5600
- for (int c = 0; c < 16; c++) s += Ha[c] * del[c];
5601
- quad += del[a] * s;
5602
- }
5603
- float gain = 2.0f * lin - quad;
5604
- if (gain > bestg + 1e-20f) {
5605
- bestg = gain; bl = tl; bm = tm; found = 1;
5606
- memcpy(bestdel, del, sizeof(del));
5607
  }
5608
  }
5609
  }
 
 
 
 
5610
  if (found) {
5611
  Ls[sj] = (uint8_t)bl; Lm[sj] = (uint8_t)bm;
5612
  for (int k = 0; k < 16; k++) E[c0 + k] -= bestdel[k];
5613
- for (int64_t i = 0; i < n_in; i++) {
5614
- const float *Hi = H + i * n_in + c0;
5615
- float s = 0.0f;
5616
- for (int k = 0; k < 16; k++) s += Hi[k] * bestdel[k];
5617
- Gd[i] -= s;
 
 
 
 
 
5618
  }
5619
  changed++;
5620
  }
@@ -5634,16 +5650,25 @@ static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5634
  bvv[i] = (float)Lm[b * NSl + j];
5635
  tv[i] = w[c0 + i];
5636
  }
5637
- for (int i = 0; i < QK_K; i++) {
5638
- const float *Hi = H + (c0 + i) * n_in + c0;
5639
- double sa=0, sb=0, st=0;
5640
- for (int k = 0; k < QK_K; k++) {
5641
- sa += (double)Hi[k] * av[k];
5642
- sb += (double)Hi[k] * bvv[k];
5643
- st += (double)Hi[k] * tv[k];
 
 
 
 
 
 
 
 
 
 
 
5644
  }
5645
- Saa += av[i]*sa; Sab += av[i]*sb; Sbb += bvv[i]*sb;
5646
- Sat += av[i]*st; Sbt += bvv[i]*st;
5647
  }
5648
  double det = Saa * Sbb - Sab * Sab;
5649
  if (fabs(det) < 1e-30) continue;
@@ -5653,31 +5678,49 @@ static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5653
  if (mn < 0.0) mn = 0.0;
5654
  float dtry = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)dn));
5655
  float mtry = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)mn));
5656
- /* exact accept test on the block's contribution */
5657
- float eo[QK_K], en[QK_K];
 
 
 
 
 
 
5658
  for (int i = 0; i < QK_K; i++) {
5659
  eo[i] = E[c0 + i];
5660
  en[i] = w[c0 + i] - (dtry * av[i] - mtry * bvv[i]);
 
5661
  }
5662
- double Lo=0, Ln=0;
5663
- for (int i = 0; i < QK_K; i++) {
5664
- const float *Hi = H + (c0 + i) * n_in + c0;
5665
- double so=0, sn=0;
5666
- for (int k = 0; k < QK_K; k++) { so += Hi[k]*eo[k]; sn += Hi[k]*en[k]; }
5667
- Lo += eo[i]*so; Ln += en[i]*sn;
5668
- }
5669
- if (Ln < Lo) {
5670
- dv[b] = dtry; mv[b] = mtry;
5671
  for (int i = 0; i < QK_K; i++) {
5672
- float d0 = en[i] - E[c0 + i];
5673
- E[c0 + i] = en[i];
5674
- (void)d0;
 
 
5675
  }
5676
- for (int64_t i = 0; i < n_in; i++) {
5677
- const float *Hi = H + i * n_in;
5678
- double s = 0.0;
5679
- for (int64_t k = 0; k < n_in; k++) s += (double)Hi[k] * (double)E[k];
5680
- Gd[i] = (float)s;
 
 
 
 
 
 
 
 
 
 
 
 
5681
  }
5682
  changed++;
5683
  }
@@ -5713,6 +5756,20 @@ static void quantize_tensor_q2k_analog(const float *weights, int64_t n_elements,
5713
  int64_t nb_row = n_in / QK_K;
5714
  double total_err = 0.0;
5715
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
5716
  #pragma omp parallel for schedule(dynamic, 1) reduction(+:total_err)
5717
  for (int64_t r = 0; r < n_rows; r++) {
5718
  float *E = (float *)malloc((size_t)n_in * sizeof(float));
@@ -5725,7 +5782,7 @@ static void quantize_tensor_q2k_analog(const float *weights, int64_t n_elements,
5725
 
5726
  const float *w = weights + r * n_in;
5727
  BlockQ2K *o = output + r * nb_row;
5728
- hex_analog_row(w, n_in, H, o, n_sweeps, E, Gd, q, Ls, Lm, dv, mv);
5729
 
5730
  for (int64_t b = 0; b < nb_row; b++)
5731
  total_err += (double)gguf_q2_k_block_error(w + b * QK_K, &o[b]);
 
5466
 
5467
  /* Optimise one output row of length n_in (a whole number of Q2_K blocks). */
5468
  static void hex_analog_row(const float *w, int64_t n_in, const float *H,
5469
+ int h_diag, BlockQ2K *out, int n_sweeps,
5470
  float *E, float *Gd, uint8_t *q,
5471
  uint8_t *Ls, uint8_t *Lm, float *dv, float *mv)
5472
  {
 
5519
  float ds = dv[b] * (float)Ls[j], ms = mv[b] * (float)Lm[j];
5520
  E[i] = w[i] - (ds * (float)q[i] - ms);
5521
  }
5522
+ if (h_diag) {
5523
+ for (int64_t i = 0; i < n_in; i++) Gd[i] = H[i * n_in + i] * E[i];
5524
+ } else {
5525
+ for (int64_t i = 0; i < n_in; i++) {
5526
+ const float *Hi = H + i * n_in;
5527
+ double s = 0.0;
5528
+ for (int64_t k = 0; k < n_in; k++) s += (double)Hi[k] * (double)E[k];
5529
+ Gd[i] = (float)s;
5530
+ }
5531
  }
5532
 
5533
  for (int sweep = 0; sweep < n_sweeps; sweep++) {
 
5554
  q[j] = (uint8_t)bestq;
5555
  E[j] -= bestd;
5556
  changed++;
5557
+ if (h_diag) {
5558
+ /* diagonal H: only Gd[j] depends on E[j]. O(1). */
5559
+ Gd[j] = Hjj * E[j];
5560
+ } else {
5561
  const float *Hj = H + j * n_in;
5562
  for (int64_t k = 0; k < n_in; k++) Gd[k] -= bestd * Hj[k];
 
 
 
 
 
 
 
 
 
 
5563
  }
 
 
 
 
 
 
 
 
 
5564
  }
5565
  }
5566
+ (void)pend;
5567
  }
5568
 
5569
  /* ── move B: joint (ls, lm) per sub-block, full 16x16 ── */
5570
  for (int64_t sj = 0; sj < n_in / 16; sj++) {
5571
  int64_t b = (sj * 16) / QK_K, c0 = sj * 16;
5572
  float dd = dv[b], mm = mv[b];
5573
+ /* del(tl,tm) is AFFINE in the candidate:
5574
+ * del = a*u - b*e, a = tl-Ls, b = tm-Lm
5575
+ * u[k] = d*q[k], e[k] = dmin (constant vector)
5576
+ * so the gain expands with five precomputed scalars
5577
+ * gain = 2(a*Lu - b*Le) - (a^2*Quu - 2ab*Que + b^2*Qee)
5578
+ * and each of the 256 candidates costs O(1) instead of the
5579
+ * O(256) 16x16 quadratic form. This inner loop dominated the
5580
+ * whole routine: ~65k ops per sub-block became ~3k. */
5581
+ float u[16], del[16], bestdel[16];
5582
+ for (int k = 0; k < 16; k++) u[k] = dd * (float)q[c0 + k];
5583
+ const float ev = mm;
5584
+
5585
+ double Quu = 0, Que = 0, Qee = 0, Lu = 0, Le = 0;
5586
+ if (h_diag) {
5587
+ for (int k = 0; k < 16; k++) {
5588
+ double h = H[(c0 + k) * n_in + c0 + k];
5589
+ Quu += h * u[k] * u[k];
5590
+ Que += h * u[k] * ev;
5591
+ Qee += h * ev * ev;
5592
+ }
5593
+ } else {
5594
+ for (int a = 0; a < 16; a++) {
5595
+ const float *Ha = H + (c0 + a) * n_in + c0;
5596
+ double su = 0, se = 0;
5597
+ for (int c = 0; c < 16; c++) { su += Ha[c] * u[c]; se += Ha[c] * ev; }
5598
+ Quu += u[a] * su; Que += u[a] * se; Qee += ev * se;
5599
+ }
5600
+ }
5601
+ for (int k = 0; k < 16; k++) { Lu += u[k] * Gd[c0 + k]; Le += ev * Gd[c0 + k]; }
5602
+
5603
  float bestg = 0.0f; int bl = Ls[sj], bm = Lm[sj]; int found = 0;
5604
+ double ba = 0.0, bb2 = 0.0;
5605
  for (int tl = 0; tl <= 15; tl++) {
5606
+ double a = (double)tl - (double)Ls[sj];
5607
  for (int tm = 0; tm <= 15; tm++) {
5608
+ double b = (double)tm - (double)Lm[sj];
5609
+ double gain = 2.0 * (a * Lu - b * Le)
5610
+ - (a*a*Quu - 2.0*a*b*Que + b*b*Qee);
5611
+ if (gain > bestg + 1e-20) {
5612
+ bestg = (float)gain; bl = tl; bm = tm; found = 1;
5613
+ ba = a; bb2 = b;
 
 
 
 
 
 
 
 
 
 
 
5614
  }
5615
  }
5616
  }
5617
+ if (found)
5618
+ for (int k = 0; k < 16; k++)
5619
+ bestdel[k] = (float)(ba * u[k] - bb2 * ev);
5620
+ (void)del;
5621
  if (found) {
5622
  Ls[sj] = (uint8_t)bl; Lm[sj] = (uint8_t)bm;
5623
  for (int k = 0; k < 16; k++) E[c0 + k] -= bestdel[k];
5624
+ if (h_diag) {
5625
+ for (int k = 0; k < 16; k++)
5626
+ Gd[c0 + k] = H[(c0 + k) * n_in + c0 + k] * E[c0 + k];
5627
+ } else {
5628
+ for (int64_t i = 0; i < n_in; i++) {
5629
+ const float *Hi = H + i * n_in + c0;
5630
+ float s = 0.0f;
5631
+ for (int k = 0; k < 16; k++) s += Hi[k] * bestdel[k];
5632
+ Gd[i] -= s;
5633
+ }
5634
  }
5635
  changed++;
5636
  }
 
5650
  bvv[i] = (float)Lm[b * NSl + j];
5651
  tv[i] = w[c0 + i];
5652
  }
5653
+ if (h_diag) {
5654
+ for (int i = 0; i < QK_K; i++) {
5655
+ double h = H[(c0 + i) * n_in + c0 + i];
5656
+ Saa += h*av[i]*av[i]; Sab += h*av[i]*bvv[i];
5657
+ Sbb += h*bvv[i]*bvv[i];
5658
+ Sat += h*av[i]*tv[i]; Sbt += h*bvv[i]*tv[i];
5659
+ }
5660
+ } else {
5661
+ for (int i = 0; i < QK_K; i++) {
5662
+ const float *Hi = H + (c0 + i) * n_in + c0;
5663
+ double sa=0, sb=0, st=0;
5664
+ for (int k = 0; k < QK_K; k++) {
5665
+ sa += (double)Hi[k] * av[k];
5666
+ sb += (double)Hi[k] * bvv[k];
5667
+ st += (double)Hi[k] * tv[k];
5668
+ }
5669
+ Saa += av[i]*sa; Sab += av[i]*sb; Sbb += bvv[i]*sb;
5670
+ Sat += av[i]*st; Sbt += bvv[i]*st;
5671
  }
 
 
5672
  }
5673
  double det = Saa * Sbb - Sab * Sab;
5674
  if (fabs(det) < 1e-30) continue;
 
5678
  if (mn < 0.0) mn = 0.0;
5679
  float dtry = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)dn));
5680
  float mtry = gguf_fp16_to_fp32(gguf_fp32_to_fp16((float)mn));
5681
+ /* Accept test. E changes only inside this block, by dl = en - eo,
5682
+ * so the exact objective change is
5683
+ * dL = 2*dl^T Gd + dl^T H_blk dl
5684
+ * -- the SAME gain form used by moves A and B. The previous
5685
+ * version compared eo^T H_blk eo against en^T H_blk en, which
5686
+ * silently dropped the cross-block coupling carried in Gd and so
5687
+ * could accept a move that raised the true objective. */
5688
+ float eo[QK_K], en[QK_K], dl[QK_K];
5689
  for (int i = 0; i < QK_K; i++) {
5690
  eo[i] = E[c0 + i];
5691
  en[i] = w[c0 + i] - (dtry * av[i] - mtry * bvv[i]);
5692
+ dl[i] = en[i] - eo[i];
5693
  }
5694
+ double lin = 0.0, quad = 0.0;
5695
+ for (int i = 0; i < QK_K; i++) lin += (double)dl[i] * Gd[c0 + i];
5696
+ if (h_diag) {
5697
+ for (int i = 0; i < QK_K; i++)
5698
+ quad += (double)dl[i]*dl[i]*H[(c0 + i) * n_in + c0 + i];
5699
+ } else {
 
 
 
5700
  for (int i = 0; i < QK_K; i++) {
5701
+ if (dl[i] == 0.0f) continue;
5702
+ const float *Hi = H + (c0 + i) * n_in + c0;
5703
+ double sn = 0.0;
5704
+ for (int k = 0; k < QK_K; k++) sn += (double)Hi[k] * dl[k];
5705
+ quad += (double)dl[i] * sn;
5706
  }
5707
+ }
5708
+ if (2.0 * lin + quad < 0.0) { /* strict decrease */
5709
+ dv[b] = dtry; mv[b] = mtry;
5710
+ for (int i = 0; i < QK_K; i++) E[c0 + i] = en[i];
5711
+ /* Banded gradient update: Gd += H[:, blk] @ dl.
5712
+ * Was a full O(n_in^2) recompute -- 10x waste at n_in=2560,
5713
+ * 24x at 6144, and it dominated the whole runtime. */
5714
+ if (h_diag) {
5715
+ for (int i = 0; i < QK_K; i++)
5716
+ Gd[c0 + i] = H[(c0 + i) * n_in + c0 + i] * E[c0 + i];
5717
+ } else {
5718
+ for (int64_t i = 0; i < n_in; i++) {
5719
+ const float *Hi = H + i * n_in + c0;
5720
+ double sm2 = 0.0;
5721
+ for (int k = 0; k < QK_K; k++) sm2 += (double)Hi[k] * dl[k];
5722
+ Gd[i] += (float)sm2;
5723
+ }
5724
  }
5725
  changed++;
5726
  }
 
5756
  int64_t nb_row = n_in / QK_K;
5757
  double total_err = 0.0;
5758
 
5759
+ /* Detect a diagonal Hessian once. An imatrix-derived H is diagonal, and
5760
+ * then Gd = diag(H) .* E has NO coupling: every gradient update collapses
5761
+ * from O(n_in) or O(n_in^2) to O(1). Same output, orders of magnitude
5762
+ * less work -- worth the one-time O(n_in^2) scan. */
5763
+ int h_diag = 1;
5764
+ for (int64_t i = 0; i < n_in && h_diag; i++) {
5765
+ const float *Hi = H + i * n_in;
5766
+ for (int64_t k = 0; k < n_in; k++)
5767
+ if (k != i && Hi[k] != 0.0f) { h_diag = 0; break; }
5768
+ }
5769
+ if (verbose)
5770
+ printf(" [analog] Hessian is %s -> %s path\n",
5771
+ h_diag ? "DIAGONAL" : "dense", h_diag ? "fast" : "full");
5772
+
5773
  #pragma omp parallel for schedule(dynamic, 1) reduction(+:total_err)
5774
  for (int64_t r = 0; r < n_rows; r++) {
5775
  float *E = (float *)malloc((size_t)n_in * sizeof(float));
 
5782
 
5783
  const float *w = weights + r * n_in;
5784
  BlockQ2K *o = output + r * nb_row;
5785
+ hex_analog_row(w, n_in, H, h_diag, o, n_sweeps, E, Gd, q, Ls, Lm, dv, mv);
5786
 
5787
  for (int64_t b = 0; b < nb_row; b++)
5788
  total_err += (double)gguf_q2_k_block_error(w + b * QK_K, &o[b]);