CompressedGemma commited on
Commit
698bf6e
·
verified ·
1 Parent(s): cf98511

Upload 6 files

Browse files
Files changed (6) hide show
  1. __main__.py +9 -0
  2. config.py +3 -0
  3. gguf.py +3 -2
  4. gui.py +66 -15
  5. patch_llamacpp_hpcq_seg.py +21 -8
  6. segmentation.py +36 -24
__main__.py CHANGED
@@ -59,6 +59,8 @@ def _config_from_args(args) -> HPCConfig:
59
  cfg.segment_k = int(args.segment_k)
60
  if getattr(args, "segment_iters", None) is not None:
61
  cfg.segment_iters = int(args.segment_iters)
 
 
62
  if args.no_q4_fallback:
63
  cfg.enable_q4_fallback = False
64
  if args.keep_embd:
@@ -329,6 +331,13 @@ def main() -> None:
329
  metavar="N",
330
  help="proxy fixed-point rounds for superblock segmentation (default 4)",
331
  )
 
 
 
 
 
 
 
332
  parser.add_argument(
333
  "--format",
334
  choices=("mixed", "mixed_q1", "custom", "q1_0", "q3_k", "q4_0", "q2_k"),
 
59
  cfg.segment_k = int(args.segment_k)
60
  if getattr(args, "segment_iters", None) is not None:
61
  cfg.segment_iters = int(args.segment_iters)
62
+ if getattr(args, "segment_table", None) is not None:
63
+ cfg.segment_table_size = int(args.segment_table)
64
  if args.no_q4_fallback:
65
  cfg.enable_q4_fallback = False
66
  if args.keep_embd:
 
331
  metavar="N",
332
  help="proxy fixed-point rounds for superblock segmentation (default 4)",
333
  )
334
+ parser.add_argument(
335
+ "--segment-table",
336
+ type=int,
337
+ default=None,
338
+ metavar="M",
339
+ help="generalized permutation lookup table size per tensor (default 256 = +0.031 bpw; 0 = full per-block)",
340
+ )
341
  parser.add_argument(
342
  "--format",
343
  choices=("mixed", "mixed_q1", "custom", "q1_0", "q3_k", "q4_0", "q2_k"),
config.py CHANGED
@@ -113,6 +113,7 @@ class HPCConfig:
113
  segment: bool = False
114
  segment_k: int = 16
115
  segment_iters: int = 4
 
116
 
117
  # ---- seed clamp -----------------------------------------------------
118
  clamp_kurtosis_off: float = 6.0
@@ -176,6 +177,8 @@ class HPCConfig:
176
  raise ValueError("segment_k must lie in [2, 256]")
177
  if self.segment_iters < 1:
178
  raise ValueError("segment_iters must be >= 1")
 
 
179
  if self.device not in ("auto", "cuda", "cpu") and not self.device.startswith("cuda:"):
180
  raise ValueError("device must be 'auto', 'cuda', 'cuda:N', or 'cpu'")
181
  if self.quant_format not in ("mixed", "mixed_q1", "custom", "q1_0", "q3_k", "q4_0", "q2_k"):
 
113
  segment: bool = False
114
  segment_k: int = 16
115
  segment_iters: int = 4
116
+ segment_table_size: int = 256
117
 
118
  # ---- seed clamp -----------------------------------------------------
119
  clamp_kurtosis_off: float = 6.0
 
177
  raise ValueError("segment_k must lie in [2, 256]")
178
  if self.segment_iters < 1:
179
  raise ValueError("segment_iters must be >= 1")
180
+ if self.segment_table_size < 0 or (0 < self.segment_table_size < 4) or self.segment_table_size > 65535:
181
+ raise ValueError("segment_table_size must be 0 (full per-block) or lie in [4, 65535]")
182
  if self.device not in ("auto", "cuda", "cpu") and not self.device.startswith("cuda:"):
183
  raise ValueError("device must be 'auto', 'cuda', 'cuda:N', or 'cpu'")
184
  if self.quant_format not in ("mixed", "mixed_q1", "custom", "q1_0", "q3_k", "q4_0", "q2_k"):
gguf.py CHANGED
@@ -798,6 +798,7 @@ def write_quantized_gguf(
798
  custom_map = getattr(cfg, "custom_quant_map", None) or None
799
  seg_active = bool(getattr(cfg, "segment", False))
800
  seg_k = max(2, int(getattr(cfg, "segment_k", 16)))
 
801
 
802
  plan: list[str | None] = []
803
  quantized_count = 0
@@ -907,7 +908,7 @@ def write_quantized_gguf(
907
  for ti, fmt in zip(reader.tensors, plan):
908
  if fmt in ("Q2_K", "Q3_K"):
909
  key = SEG_KV_PREFIX + ti.name
910
- content_len = label_channel_bytes(ti.n_elements, seg_k)
911
  seg_reserve[key] = content_len
912
  # GGUF uint8 array: u32 element type + u64 count + zeros
913
  kv_pairs.append(
@@ -936,7 +937,7 @@ def write_quantized_gguf(
936
  fout.write(raw)
937
 
938
  def _patch_seg_kv(fout, key: str, labels: np.ndarray) -> None:
939
- packed = pack_label_channel(labels, seg_k)
940
  if len(packed) != seg_reserve.get(key, -1):
941
  raise RuntimeError(f"{key}: label channel size {len(packed)} != reserved")
942
  pos = fout.tell()
 
798
  custom_map = getattr(cfg, "custom_quant_map", None) or None
799
  seg_active = bool(getattr(cfg, "segment", False))
800
  seg_k = max(2, int(getattr(cfg, "segment_k", 16)))
801
+ seg_table_size = int(getattr(cfg, "segment_table_size", 256))
802
 
803
  plan: list[str | None] = []
804
  quantized_count = 0
 
908
  for ti, fmt in zip(reader.tensors, plan):
909
  if fmt in ("Q2_K", "Q3_K"):
910
  key = SEG_KV_PREFIX + ti.name
911
+ content_len = label_channel_bytes(ti.n_elements, seg_k, seg_table_size)
912
  seg_reserve[key] = content_len
913
  # GGUF uint8 array: u32 element type + u64 count + zeros
914
  kv_pairs.append(
 
937
  fout.write(raw)
938
 
939
  def _patch_seg_kv(fout, key: str, labels: np.ndarray) -> None:
940
+ packed = pack_label_channel(labels, seg_k, seg_table_size)
941
  if len(packed) != seg_reserve.get(key, -1):
942
  raise RuntimeError(f"{key}: label channel size {len(packed)} != reserved")
943
  pos = fout.tell()
gui.py CHANGED
@@ -70,10 +70,12 @@ from hpc_quantize_v2 import (
70
  read_imatrix,
71
  )
72
  from hpc_quantize_v2.gguf import (
 
73
  is_q1_0_eligible,
74
  is_q2k_eligible,
75
  is_q3_k_eligible,
76
  is_q4_0_eligible,
 
77
  target_quant_type,
78
  tensor_type_key,
79
  write_quantized_gguf,
@@ -187,6 +189,7 @@ class HPCQuantizeGUI:
187
  self.segment_var = tk.BooleanVar(value=False)
188
  self.segment_k_var = tk.IntVar(value=16)
189
  self.segment_iters_var = tk.IntVar(value=4)
 
190
 
191
  self.cli_preview_var = tk.StringVar()
192
  self.status_var = tk.StringVar(value="Ready.")
@@ -225,8 +228,11 @@ class HPCQuantizeGUI:
225
  self.segment_var,
226
  self.segment_k_var,
227
  self.segment_iters_var,
 
228
  ):
229
  v.trace_add("write", lambda *_: self._update_cli_preview())
 
 
230
 
231
  def _apply_theme(self) -> None:
232
  style = ttk.Style(self.root)
@@ -504,10 +510,17 @@ class HPCQuantizeGUI:
504
  ttk.Spinbox(
505
  spec_f, from_=1, to=12, increment=1, textvariable=self.segment_iters_var, width=8
506
  ).grid(row=4, column=3, sticky=tk.EW, padx=4, pady=2)
 
 
 
 
 
 
 
507
  ttk.Label(
508
  spec_f, text="(λ multiplier ≥1 during thermal anneal; ~1e2–1e3 ≈ zero-coherent-residual '0-error' anneal)",
509
  foreground="#93c5fd",
510
- ).grid(row=5, column=0, columnspan=4, sticky=tk.W, padx=4, pady=(0, 2))
511
 
512
  # Right: GGUF Tensor Browser
513
  browser_frame = ttk.LabelFrame(right_pane, text="3. GGUF Tensor Inspector & Live Status", padding=8)
@@ -840,6 +853,7 @@ class HPCQuantizeGUI:
840
  segment=bool(self.segment_var.get()),
841
  segment_k=max(2, int(self.segment_k_var.get())),
842
  segment_iters=max(1, int(self.segment_iters_var.get())),
 
843
  quant_format=self.quant_format_var.get(),
844
  custom_quant_map=self._custom_map_for_plan() or {},
845
  enable_q4_fallback=bool(self.q4_fallback_var.get()) and self.quant_format_var.get() not in ("q2_k", "q3_k"),
@@ -899,6 +913,8 @@ class HPCQuantizeGUI:
899
  cmd += ["--segment-k", str(int(self.segment_k_var.get()))]
900
  if int(self.segment_iters_var.get()) != 4:
901
  cmd += ["--segment-iters", str(int(self.segment_iters_var.get()))]
 
 
902
  cmd.append("--verbose")
903
  return " ".join(shlex.quote(part) for part in cmd)
904
 
@@ -962,6 +978,43 @@ class HPCQuantizeGUI:
962
  self._refresh_plan_view()
963
  self._update_cli_preview()
964
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
965
  def _refresh_plan_view(self) -> None:
966
  if self._reader is None:
967
  return
@@ -969,7 +1022,8 @@ class HPCQuantizeGUI:
969
  scope_selected = self.mode_var.get() == "selected"
970
  n_q2, n_q3, n_q4, n_copy = 0, 0, 0, 0
971
  n_q1 = 0
972
- planned_bytes = 0
 
973
  for ti in self._reader.tensors:
974
  tgt = self._plan_target(ti, chosen, scope_selected)
975
  if tgt == "Q2_K":
@@ -993,6 +1047,9 @@ class HPCQuantizeGUI:
993
  planned_bytes += (ti.n_elements // QK1_0) * Q1_0_BLOCK_BYTES
994
  else:
995
  planned_bytes += ti.n_bytes
 
 
 
996
 
997
  if not self.tree.exists(ti.name):
998
  continue
@@ -1010,12 +1067,7 @@ class HPCQuantizeGUI:
1010
  vals[6] = state_str
1011
  self.tree.item(ti.name, tags=(tgt_label,), values=vals)
1012
 
1013
- if planned_bytes >= 1024 ** 3:
1014
- planned_text = f"{planned_bytes / (1024 ** 3):.2f} GiB"
1015
- elif planned_bytes >= 1024 ** 2:
1016
- planned_text = f"{planned_bytes / (1024 ** 2):.1f} MiB"
1017
- else:
1018
- planned_text = f"{planned_bytes / 1024:.1f} KiB"
1019
  scope_text = "selected" if scope_selected else "all"
1020
  self.status_var.set(
1021
  f"Plan ({self.quant_format_var.get().upper()}, {scope_text}): {n_q1} Q1_0 | {n_q2} Q2_K | {n_q3} Q3_K | {n_q4} Q4_0 | {n_copy} Copy "
@@ -1032,7 +1084,8 @@ class HPCQuantizeGUI:
1032
  plan_filt = self.plan_filter_var.get()
1033
  n_q2, n_q3, n_q4, n_copy = 0, 0, 0, 0
1034
  n_q1 = 0
1035
- planned_bytes = 0
 
1036
  for ti in self._reader.tensors:
1037
  tgt = self._plan_target(ti, chosen, scope_selected)
1038
  if tgt == "Q2_K":
@@ -1056,6 +1109,9 @@ class HPCQuantizeGUI:
1056
  planned_bytes += (ti.n_elements // QK1_0) * Q1_0_BLOCK_BYTES
1057
  else:
1058
  planned_bytes += ti.n_bytes
 
 
 
1059
 
1060
  tgt_label = tgt if tgt else "Copy"
1061
  if plan_filt != "ALL" and tgt_label.upper() != plan_filt:
@@ -1097,12 +1153,7 @@ class HPCQuantizeGUI:
1097
  keep_selected = [n for n in chosen if self.tree.exists(n)]
1098
  if keep_selected:
1099
  self.tree.selection_set(keep_selected)
1100
- if planned_bytes >= 1024 ** 3:
1101
- planned_text = f"{planned_bytes / (1024 ** 3):.2f} GiB"
1102
- elif planned_bytes >= 1024 ** 2:
1103
- planned_text = f"{planned_bytes / (1024 ** 2):.1f} MiB"
1104
- else:
1105
- planned_text = f"{planned_bytes / 1024:.1f} KiB"
1106
  scope_text = "selected" if scope_selected else "all"
1107
  self.status_var.set(
1108
  f"Plan ({self.quant_format_var.get().upper()}, {scope_text}): {n_q1} Q1_0 | {n_q2} Q2_K | {n_q3} Q3_K | {n_q4} Q4_0 | {n_copy} Copy "
 
70
  read_imatrix,
71
  )
72
  from hpc_quantize_v2.gguf import (
73
+ SEG_KV_PREFIX,
74
  is_q1_0_eligible,
75
  is_q2k_eligible,
76
  is_q3_k_eligible,
77
  is_q4_0_eligible,
78
+ label_channel_bytes,
79
  target_quant_type,
80
  tensor_type_key,
81
  write_quantized_gguf,
 
189
  self.segment_var = tk.BooleanVar(value=False)
190
  self.segment_k_var = tk.IntVar(value=16)
191
  self.segment_iters_var = tk.IntVar(value=4)
192
+ self.segment_table_var = tk.IntVar(value=256)
193
 
194
  self.cli_preview_var = tk.StringVar()
195
  self.status_var = tk.StringVar(value="Ready.")
 
228
  self.segment_var,
229
  self.segment_k_var,
230
  self.segment_iters_var,
231
+ self.segment_table_var,
232
  ):
233
  v.trace_add("write", lambda *_: self._update_cli_preview())
234
+ for v in (self.segment_var, self.segment_k_var, self.segment_table_var):
235
+ v.trace_add("write", lambda *_: self._refresh_plan_view())
236
 
237
  def _apply_theme(self) -> None:
238
  style = ttk.Style(self.root)
 
510
  ttk.Spinbox(
511
  spec_f, from_=1, to=12, increment=1, textvariable=self.segment_iters_var, width=8
512
  ).grid(row=4, column=3, sticky=tk.EW, padx=4, pady=2)
513
+ ttk.Label(spec_f, text="Seg Lookup Table (0=full):").grid(row=5, column=0, columnspan=2, sticky=tk.W, padx=4, pady=2)
514
+ ttk.Combobox(
515
+ spec_f,
516
+ textvariable=self.segment_table_var,
517
+ values=(0, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096),
518
+ width=8,
519
+ ).grid(row=5, column=2, columnspan=2, sticky=tk.EW, padx=4, pady=2)
520
  ttk.Label(
521
  spec_f, text="(λ multiplier ≥1 during thermal anneal; ~1e2–1e3 ≈ zero-coherent-residual '0-error' anneal)",
522
  foreground="#93c5fd",
523
+ ).grid(row=6, column=0, columnspan=4, sticky=tk.W, padx=4, pady=(0, 2))
524
 
525
  # Right: GGUF Tensor Browser
526
  browser_frame = ttk.LabelFrame(right_pane, text="3. GGUF Tensor Inspector & Live Status", padding=8)
 
853
  segment=bool(self.segment_var.get()),
854
  segment_k=max(2, int(self.segment_k_var.get())),
855
  segment_iters=max(1, int(self.segment_iters_var.get())),
856
+ segment_table_size=max(0, int(self.segment_table_var.get())),
857
  quant_format=self.quant_format_var.get(),
858
  custom_quant_map=self._custom_map_for_plan() or {},
859
  enable_q4_fallback=bool(self.q4_fallback_var.get()) and self.quant_format_var.get() not in ("q2_k", "q3_k"),
 
913
  cmd += ["--segment-k", str(int(self.segment_k_var.get()))]
914
  if int(self.segment_iters_var.get()) != 4:
915
  cmd += ["--segment-iters", str(int(self.segment_iters_var.get()))]
916
+ if int(self.segment_table_var.get()) != 256:
917
+ cmd += ["--segment-table", str(int(self.segment_table_var.get()))]
918
  cmd.append("--verbose")
919
  return " ".join(shlex.quote(part) for part in cmd)
920
 
 
978
  self._refresh_plan_view()
979
  self._update_cli_preview()
980
 
981
+ def _estimate_seg_bytes(self, ti, tgt: str | None) -> int:
982
+ if not bool(self.segment_var.get()) or tgt not in ("Q2_K", "Q3_K"):
983
+ return 0
984
+ try:
985
+ seg_k = max(2, int(self.segment_k_var.get()))
986
+ except Exception:
987
+ seg_k = 16
988
+ try:
989
+ seg_tbl = max(0, int(self.segment_table_var.get()))
990
+ except Exception:
991
+ seg_tbl = 256
992
+ content_len = label_channel_bytes(ti.n_elements, seg_k, seg_tbl)
993
+ key_len = len((SEG_KV_PREFIX + ti.name).encode("utf-8"))
994
+ return 8 + key_len + 4 + 12 + content_len
995
+
996
+ def _format_size_summary(self, planned_bytes: int, total_seg_bytes: int) -> str:
997
+ if planned_bytes >= 1024 ** 3:
998
+ planned_text = f"{planned_bytes / (1024 ** 3):.2f} GiB"
999
+ elif planned_bytes >= 1024 ** 2:
1000
+ planned_text = f"{planned_bytes / (1024 ** 2):.1f} MiB"
1001
+ else:
1002
+ planned_text = f"{planned_bytes / 1024:.1f} KiB"
1003
+ if total_seg_bytes > 0:
1004
+ try:
1005
+ seg_tbl = max(0, int(self.segment_table_var.get()))
1006
+ except Exception:
1007
+ seg_tbl = 256
1008
+ if total_seg_bytes >= 1024 ** 3:
1009
+ seg_txt = f"+{total_seg_bytes / (1024 ** 3):.2f} GiB"
1010
+ elif total_seg_bytes >= 1024 ** 2:
1011
+ seg_txt = f"+{total_seg_bytes / (1024 ** 2):.1f} MiB"
1012
+ else:
1013
+ seg_txt = f"+{total_seg_bytes / 1024:.1f} KiB"
1014
+ lut_lbl = f"LUT={seg_tbl}" if seg_tbl > 0 else "full"
1015
+ planned_text = f"{planned_text} [seg {seg_txt}, {lut_lbl}]"
1016
+ return planned_text
1017
+
1018
  def _refresh_plan_view(self) -> None:
1019
  if self._reader is None:
1020
  return
 
1022
  scope_selected = self.mode_var.get() == "selected"
1023
  n_q2, n_q3, n_q4, n_copy = 0, 0, 0, 0
1024
  n_q1 = 0
1025
+ planned_bytes = int(getattr(self._reader, "data_start", 0))
1026
+ total_seg_bytes = 0
1027
  for ti in self._reader.tensors:
1028
  tgt = self._plan_target(ti, chosen, scope_selected)
1029
  if tgt == "Q2_K":
 
1047
  planned_bytes += (ti.n_elements // QK1_0) * Q1_0_BLOCK_BYTES
1048
  else:
1049
  planned_bytes += ti.n_bytes
1050
+ seg_b = self._estimate_seg_bytes(ti, tgt)
1051
+ planned_bytes += seg_b
1052
+ total_seg_bytes += seg_b
1053
 
1054
  if not self.tree.exists(ti.name):
1055
  continue
 
1067
  vals[6] = state_str
1068
  self.tree.item(ti.name, tags=(tgt_label,), values=vals)
1069
 
1070
+ planned_text = self._format_size_summary(planned_bytes, total_seg_bytes)
 
 
 
 
 
1071
  scope_text = "selected" if scope_selected else "all"
1072
  self.status_var.set(
1073
  f"Plan ({self.quant_format_var.get().upper()}, {scope_text}): {n_q1} Q1_0 | {n_q2} Q2_K | {n_q3} Q3_K | {n_q4} Q4_0 | {n_copy} Copy "
 
1084
  plan_filt = self.plan_filter_var.get()
1085
  n_q2, n_q3, n_q4, n_copy = 0, 0, 0, 0
1086
  n_q1 = 0
1087
+ planned_bytes = int(getattr(self._reader, "data_start", 0))
1088
+ total_seg_bytes = 0
1089
  for ti in self._reader.tensors:
1090
  tgt = self._plan_target(ti, chosen, scope_selected)
1091
  if tgt == "Q2_K":
 
1109
  planned_bytes += (ti.n_elements // QK1_0) * Q1_0_BLOCK_BYTES
1110
  else:
1111
  planned_bytes += ti.n_bytes
1112
+ seg_b = self._estimate_seg_bytes(ti, tgt)
1113
+ planned_bytes += seg_b
1114
+ total_seg_bytes += seg_b
1115
 
1116
  tgt_label = tgt if tgt else "Copy"
1117
  if plan_filt != "ALL" and tgt_label.upper() != plan_filt:
 
1153
  keep_selected = [n for n in chosen if self.tree.exists(n)]
1154
  if keep_selected:
1155
  self.tree.selection_set(keep_selected)
1156
+ planned_text = self._format_size_summary(planned_bytes, total_seg_bytes)
 
 
 
 
 
1157
  scope_text = "selected" if scope_selected else "all"
1158
  self.status_var.set(
1159
  f"Plan ({self.quant_format_var.get().upper()}, {scope_text}): {n_q1} Q1_0 | {n_q2} Q2_K | {n_q3} Q3_K | {n_q4} Q4_0 | {n_copy} Copy "
patch_llamacpp_hpcq_seg.py CHANGED
@@ -78,13 +78,17 @@ struct hpcq_seg_entry {
78
  uintptr_t end; // payload end (exclusive)
79
  int64_t block_bytes; // 84 for Q2_K, 110 for Q3_K
80
  int64_t n_blocks; // number of 256-weight superblocks
81
- uint8_t enc; // 0 = 4-bit nibbles, 1 = 8-bit, 2 = 256-entry static table + 1B/blk
82
  uint8_t * packed; // compact packed label stream
83
  void * dev_packed; // lazily allocated device copy for CUDA (or NULL)
84
  };
85
 
86
- static inline HPCQ_HD size_t hpcq_seg_packed_bytes(int64_t n_blocks, uint8_t enc) {
87
- if (enc == 2) return 32768u + (size_t) n_blocks;
 
 
 
 
88
  if (enc == 0) return (size_t) n_blocks * 128u;
89
  return (size_t) n_blocks * 256u;
90
  }
@@ -139,9 +143,17 @@ static inline HPCQ_HD void hpcq_seg_make_order_block(
139
 
140
  static inline HPCQ_HD void hpcq_seg_make_order_at(
141
  const uint8_t * packed_base, uint8_t enc, int64_t blk_idx, uint16_t * order) {
142
- if (enc == 2) {
143
- const uint8_t m = packed_base[32768u + (size_t) blk_idx];
144
- hpcq_seg_make_order_block(packed_base + (size_t) m * 128u, 0, order);
 
 
 
 
 
 
 
 
145
  } else if (enc == 0) {
146
  hpcq_seg_make_order_block(packed_base + (size_t) blk_idx * 128u, 0, order);
147
  } else {
@@ -180,7 +192,8 @@ void ggml_hpcq_seg_bind_raw(
180
  if (n_labels == 0 || (n_labels % 256) != 0) return;
181
  const int64_t n_blocks = (int64_t) (n_labels / 256);
182
  if (n_blocks * block_bytes != n_bytes) return;
183
- const size_t need = hpcq_seg_packed_bytes(n_blocks, enc);
 
184
  if (hqs1_len < 10u + need) return;
185
 
186
  const uintptr_t base = (uintptr_t) payload;
@@ -550,7 +563,7 @@ static inline const uint8_t * hpcq_cuda_get_dev_packed(
550
  if (!seg) return nullptr;
551
  *out_enc = seg->enc;
552
  if (seg->dev_packed == nullptr) {{
553
- const size_t nbytes = hpcq_seg_packed_bytes(seg->n_blocks, seg->enc);
554
  void * d_ptr = nullptr;
555
  if (cudaMalloc(&d_ptr, nbytes) == cudaSuccess && d_ptr != nullptr) {{
556
  if (cudaMemcpy(d_ptr, seg->packed, nbytes, cudaMemcpyHostToDevice) == cudaSuccess) {{
 
78
  uintptr_t end; // payload end (exclusive)
79
  int64_t block_bytes; // 84 for Q2_K, 110 for Q3_K
80
  int64_t n_blocks; // number of 256-weight superblocks
81
+ uint8_t enc; // 0=nibbles, 1=bytes, 2=LUT+1B/blk, 3=LUT+2B/blk
82
  uint8_t * packed; // compact packed label stream
83
  void * dev_packed; // lazily allocated device copy for CUDA (or NULL)
84
  };
85
 
86
+ static inline HPCQ_HD size_t hpcq_seg_packed_bytes(int64_t n_blocks, uint8_t enc, const uint8_t * data) {
87
+ if (enc == 2 || enc == 3) {
88
+ const uint16_t tsize = (uint16_t) data[0] | ((uint16_t) data[1] << 8);
89
+ const size_t idx_bytes = (enc == 2) ? 1u : 2u;
90
+ return 2u + (size_t) tsize * 128u + (size_t) n_blocks * idx_bytes;
91
+ }
92
  if (enc == 0) return (size_t) n_blocks * 128u;
93
  return (size_t) n_blocks * 256u;
94
  }
 
143
 
144
  static inline HPCQ_HD void hpcq_seg_make_order_at(
145
  const uint8_t * packed_base, uint8_t enc, int64_t blk_idx, uint16_t * order) {
146
+ if (enc == 2 || enc == 3) {
147
+ const uint16_t tsize = (uint16_t) packed_base[0] | ((uint16_t) packed_base[1] << 8);
148
+ const size_t idx_off = 2u + (size_t) tsize * 128u;
149
+ uint16_t m = 0;
150
+ if (enc == 2) {
151
+ m = packed_base[idx_off + (size_t) blk_idx];
152
+ } else {
153
+ const uint8_t * p = packed_base + idx_off + (size_t) blk_idx * 2u;
154
+ m = (uint16_t) p[0] | ((uint16_t) p[1] << 8);
155
+ }
156
+ hpcq_seg_make_order_block(packed_base + 2u + (size_t) m * 128u, 0, order);
157
  } else if (enc == 0) {
158
  hpcq_seg_make_order_block(packed_base + (size_t) blk_idx * 128u, 0, order);
159
  } else {
 
192
  if (n_labels == 0 || (n_labels % 256) != 0) return;
193
  const int64_t n_blocks = (int64_t) (n_labels / 256);
194
  if (n_blocks * block_bytes != n_bytes) return;
195
+ if ((enc == 2 || enc == 3) && hqs1_len < 12u) return;
196
+ const size_t need = hpcq_seg_packed_bytes(n_blocks, enc, hqs1 + 10);
197
  if (hqs1_len < 10u + need) return;
198
 
199
  const uintptr_t base = (uintptr_t) payload;
 
563
  if (!seg) return nullptr;
564
  *out_enc = seg->enc;
565
  if (seg->dev_packed == nullptr) {{
566
+ const size_t nbytes = hpcq_seg_packed_bytes(seg->n_blocks, seg->enc, seg->packed);
567
  void * d_ptr = nullptr;
568
  if (cudaMalloc(&d_ptr, nbytes) == cudaSuccess && d_ptr != nullptr) {{
569
  if (cudaMemcpy(d_ptr, seg->packed, nbytes, cudaMemcpyHostToDevice) == cudaSuccess) {{
segmentation.py CHANGED
@@ -114,11 +114,12 @@ def scatter_to_original(xh_slots: np.ndarray, order: np.ndarray) -> np.ndarray:
114
  #
115
  # Layout: magic "HQS1" | uint8 k | uint8 enc | uint32 n_labels (little-endian)
116
  # | payload.
117
- # enc=0: packs labels 0..15 as nibbles (2 labels per byte; used for n <= 65536)
118
  # enc=1: stores raw uint8 labels (k > 16 only)
119
- # enc=2: generalized 256-entry permutation table (32 KB nibble table + 1 byte
120
- # index per 256-weight superblock = +0.03125 bpw; used for n > 65536,
121
- # k <= 16, n % 256 == 0).
 
122
 
123
  import struct as _struct
124
 
@@ -142,24 +143,29 @@ def _unpack_nibbles_1d(data: bytes, n: int) -> np.ndarray:
142
  return out[:n]
143
 
144
 
145
- def label_channel_bytes(n: int, k: int) -> int:
146
  """Exact content size of the packed label channel."""
 
147
  if int(k) <= 16:
148
- if n > SEG_TABLE_SIZE * 256 and n % 256 == 0:
149
- return 10 + SEG_TABLE_SIZE * 128 + (n // 256)
 
 
150
  return 10 + (n + 1) // 2
151
  return 10 + n
152
 
153
 
154
- def pack_label_channel(labels: np.ndarray, k: int) -> bytes:
155
  flat = np.ascontiguousarray(labels, dtype=np.int64).ravel()
156
  n = int(flat.size)
157
  k = int(k)
 
158
  if k <= 16:
159
  if flat.size and int(flat.max()) > 15:
160
  raise ValueError(f"label {int(flat.max())} does not fit nibble packing (k={k})")
161
- if n > SEG_TABLE_SIZE * 256 and n % 256 == 0:
162
- enc = 2
 
163
  n_blocks = n // 256
164
  rows = flat.reshape(n_blocks, 256).astype(np.uint8)
165
  lo = rows[:, 0::2]
@@ -171,17 +177,17 @@ def pack_label_channel(labels: np.ndarray, k: int) -> bytes:
171
  row_views, return_inverse=True, return_counts=True
172
  )
173
  uniq_rows = uniq_views.view(np.uint8).reshape(-1, 128)
174
- table = np.zeros((SEG_TABLE_SIZE, 128), dtype=np.uint8)
175
- if uniq_rows.shape[0] <= SEG_TABLE_SIZE:
176
  table[: uniq_rows.shape[0]] = uniq_rows
177
- block_idx = inv_idx.astype(np.uint8)
178
  else:
179
- top_u = np.argsort(-counts)[:SEG_TABLE_SIZE]
180
  table[:] = uniq_rows[top_u]
181
- remap = np.zeros(uniq_rows.shape[0], dtype=np.uint8)
182
- remap[top_u] = np.arange(SEG_TABLE_SIZE, dtype=np.uint8)
183
  block_idx = remap[inv_idx]
184
- data = table.tobytes() + block_idx.tobytes()
185
  else:
186
  enc = 0
187
  data = _pack_nibbles_1d(flat)
@@ -202,13 +208,19 @@ def unpack_label_channel(content: bytes) -> tuple[int, np.ndarray]:
202
  data = content[10:]
203
  if enc == 0:
204
  out = _unpack_nibbles_1d(data, n)
205
- elif enc == 2:
206
  n_blocks = n // 256
207
- tbytes = SEG_TABLE_SIZE * 128
208
- table = _unpack_nibbles_1d(data[:tbytes], SEG_TABLE_SIZE * 256).reshape(
209
- SEG_TABLE_SIZE, 256
 
210
  )
211
- idx = np.frombuffer(data[tbytes : tbytes + n_blocks], dtype=np.uint8)
 
 
 
 
 
212
  out = table[idx].reshape(-1)
213
  else:
214
  out = np.frombuffer(data, dtype=np.uint8).copy()
@@ -268,10 +280,10 @@ def segment_stage(
268
  Wo = None if W is None else np.asarray(W, dtype=np.float64)
269
  k = max(1, int(k))
270
  n_blocks, B = Xo.shape
271
- if n_blocks <= SEG_TABLE_SIZE or B != 256 or k > 16:
 
272
  return _segment_stage_dense(Xo, Wo, cfg, proxy, k)
273
 
274
- M = SEG_TABLE_SIZE
275
  table_labels = np.empty((M, 256), dtype=np.int64)
276
  # Entry 0: natural order (monotone labels so make_order is identity [0..255])
277
  table_labels[0] = _order_to_canonical_labels(np.arange(256, dtype=np.int64), k)
 
114
  #
115
  # Layout: magic "HQS1" | uint8 k | uint8 enc | uint32 n_labels (little-endian)
116
  # | payload.
117
+ # enc=0: packs labels 0..15 as nibbles (2 labels per byte; full per-block or small tensors)
118
  # enc=1: stores raw uint8 labels (k > 16 only)
119
+ # enc=2: generalized permutation lookup table with 1-byte block indices (table_size in 4..256):
120
+ # uint16 table_size | uint8 table[table_size * 128] | uint8 block_idx[n_blocks]
121
+ # enc=3: generalized permutation lookup table with 2-byte block indices (table_size in 257..65535):
122
+ # uint16 table_size | uint8 table[table_size * 128] | uint16 block_idx[n_blocks]
123
 
124
  import struct as _struct
125
 
 
143
  return out[:n]
144
 
145
 
146
+ def label_channel_bytes(n: int, k: int, table_size: int = SEG_TABLE_SIZE) -> int:
147
  """Exact content size of the packed label channel."""
148
+ table_size = int(table_size)
149
  if int(k) <= 16:
150
+ if table_size > 0 and n > table_size * 256 and n % 256 == 0:
151
+ n_blocks = n // 256
152
+ idx_bytes = 1 if table_size <= 256 else 2
153
+ return 10 + 2 + table_size * 128 + n_blocks * idx_bytes
154
  return 10 + (n + 1) // 2
155
  return 10 + n
156
 
157
 
158
+ def pack_label_channel(labels: np.ndarray, k: int, table_size: int = SEG_TABLE_SIZE) -> bytes:
159
  flat = np.ascontiguousarray(labels, dtype=np.int64).ravel()
160
  n = int(flat.size)
161
  k = int(k)
162
+ table_size = int(table_size)
163
  if k <= 16:
164
  if flat.size and int(flat.max()) > 15:
165
  raise ValueError(f"label {int(flat.max())} does not fit nibble packing (k={k})")
166
+ if table_size > 0 and n > table_size * 256 and n % 256 == 0:
167
+ enc = 2 if table_size <= 256 else 3
168
+ idx_dt = np.uint8 if enc == 2 else np.dtype("<u2")
169
  n_blocks = n // 256
170
  rows = flat.reshape(n_blocks, 256).astype(np.uint8)
171
  lo = rows[:, 0::2]
 
177
  row_views, return_inverse=True, return_counts=True
178
  )
179
  uniq_rows = uniq_views.view(np.uint8).reshape(-1, 128)
180
+ table = np.zeros((table_size, 128), dtype=np.uint8)
181
+ if uniq_rows.shape[0] <= table_size:
182
  table[: uniq_rows.shape[0]] = uniq_rows
183
+ block_idx = inv_idx.astype(idx_dt)
184
  else:
185
+ top_u = np.argsort(-counts)[:table_size]
186
  table[:] = uniq_rows[top_u]
187
+ remap = np.zeros(uniq_rows.shape[0], dtype=idx_dt)
188
+ remap[top_u] = np.arange(table_size, dtype=idx_dt)
189
  block_idx = remap[inv_idx]
190
+ data = _struct.pack("<H", table_size) + table.tobytes() + block_idx.tobytes()
191
  else:
192
  enc = 0
193
  data = _pack_nibbles_1d(flat)
 
208
  data = content[10:]
209
  if enc == 0:
210
  out = _unpack_nibbles_1d(data, n)
211
+ elif enc in (2, 3):
212
  n_blocks = n // 256
213
+ tsize = _struct.unpack("<H", data[:2])[0]
214
+ tbytes = tsize * 128
215
+ table = _unpack_nibbles_1d(data[2 : 2 + tbytes], tsize * 256).reshape(
216
+ tsize, 256
217
  )
218
+ if enc == 2:
219
+ idx = np.frombuffer(data[2 + tbytes : 2 + tbytes + n_blocks], dtype=np.uint8)
220
+ else:
221
+ idx = np.frombuffer(
222
+ data[2 + tbytes : 2 + tbytes + 2 * n_blocks], dtype=np.dtype("<u2")
223
+ )
224
  out = table[idx].reshape(-1)
225
  else:
226
  out = np.frombuffer(data, dtype=np.uint8).copy()
 
280
  Wo = None if W is None else np.asarray(W, dtype=np.float64)
281
  k = max(1, int(k))
282
  n_blocks, B = Xo.shape
283
+ M = int(getattr(cfg, "segment_table_size", SEG_TABLE_SIZE))
284
+ if M <= 0 or n_blocks <= M or B != 256 or k > 16:
285
  return _segment_stage_dense(Xo, Wo, cfg, proxy, k)
286
 
 
287
  table_labels = np.empty((M, 256), dtype=np.int64)
288
  # Entry 0: natural order (monotone labels so make_order is identity [0..255])
289
  table_labels[0] = _order_to_canonical_labels(np.arange(256, dtype=np.int64), k)