31827c3586
Every instruction family of the live hash is bit-exact on Intel's compiler (the family probe, runs a and b) and the hash is wrong on every lane in both exchange modes: the trace of lane 0's registers after every instruction, GPU against the CPU interpreter, names the first diverging instruction. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
74 lines
3.8 KiB
Python
74 lines
3.8 KiB
Python
#!/usr/bin/env python3
|
|
"""The register-trace pack for the Arc B580 bisect (7 October 2026, branch intel-arc, docs/plans/intel-arc.md section 7).
|
|
|
|
Copies a pack directory and appends to its kernel_bound.cl a kernel `igneum_hash_trace`: the text of
|
|
`igneum_hash_bound` (the kernel the worker's self-test runs on the vectors, init words from `initw`) with one more
|
|
argument, `__global uint* trace`, and a store of lane 0's eight registers (the work-item with get_global_id(0) == 0,
|
|
nonce = baseNonce) after the init block and after EVERY program instruction (`// N family`) and shadow instruction
|
|
(`// sN family`), in execution order: the same order igneum-pow's `trace-regs` prints from the CPU interpreter. A
|
|
labels file (trace_steps.txt: one line per snapshot, "k it n family") goes beside the pack; the host reads its line
|
|
count for the buffer size and prints the labels with the registers.
|
|
|
|
tools/intel-arc/make-trace-pack.py <pack dir> <out dir>
|
|
"""
|
|
import re
|
|
import shutil
|
|
import sys
|
|
from pathlib import Path
|
|
|
|
src = Path(sys.argv[1])
|
|
out = Path(sys.argv[2])
|
|
if out.exists():
|
|
shutil.rmtree(out)
|
|
shutil.copytree(src, out)
|
|
text = (out / "kernel_bound.cl").read_text()
|
|
lines = text.split("\n")
|
|
start = next(i for i, l in enumerate(lines) if l.startswith("IGNEUM_KERNEL_HASH void igneum_hash_bound("))
|
|
end = next(i for i in range(start, len(lines)) if lines[i] == "}")
|
|
body = lines[start : end + 1]
|
|
head = body[0]
|
|
assert ", __global const uint* initw" in head, head
|
|
# the trace argument LAST, so the self-test's argument order (ds, out, baseNonce, mask, initw[, hot][, scratch]) holds
|
|
head = head.replace("igneum_hash_bound(", "igneum_hash_trace(").replace(") {", ", __global uint* trace) {")
|
|
out_lines = [head]
|
|
labels = []
|
|
k = 0
|
|
it_depth = 0
|
|
STORE = "if (gid == 0u) { trace[tc * 8u + 0u] = r0; trace[tc * 8u + 1u] = r1; trace[tc * 8u + 2u] = r2; trace[tc * 8u + 3u] = r3; trace[tc * 8u + 4u] = r4; trace[tc * 8u + 5u] = r5; trace[tc * 8u + 6u] = r6; trace[tc * 8u + 7u] = r7; } tc += 1u;"
|
|
inst = re.compile(r"//\s+(s?)(\d+)\s+(\w+)\s*$")
|
|
seen_loop = False
|
|
for l in body[1:]:
|
|
out_lines.append(l)
|
|
if l.startswith(" uint r0, r1,"):
|
|
out_lines.append(" uint tc = 0u;")
|
|
continue
|
|
if l.startswith(" for (uint it = 0u;") and not seen_loop:
|
|
# the init snapshot sits before the loop: the registers after the eight init blocks
|
|
out_lines.insert(len(out_lines) - 1, " " + STORE)
|
|
labels.append("init")
|
|
seen_loop = True
|
|
continue
|
|
m = inst.search(l)
|
|
if m and seen_loop:
|
|
indent = l[: len(l) - len(l.lstrip())]
|
|
out_lines.append(indent + STORE)
|
|
labels.append(("s" if m.group(1) else "") + m.group(2) + " " + m.group(3))
|
|
# the snapshot count: init + 8 x (64 main + shadow reps x shadow instrs). The kernel loop carries the count itself;
|
|
# the labels file lists one iteration's labels and the host multiplies, so write the per-iteration list and the totals.
|
|
main_labels = [x for x in labels if x != "init" and not x.startswith("s")]
|
|
shadow_labels = [x for x in labels if x.startswith("s")]
|
|
ktext = "\n".join(lines[: end + 1] + ["", "// register trace of igneum_hash_bound for lane 0 (tools/intel-arc/make-trace-pack.py)"] + out_lines + lines[end + 1 :])
|
|
(out / "kernel_bound.cl").write_text(ktext)
|
|
reps = 0
|
|
for l in body:
|
|
m2 = re.search(r"for \(uint sh = 0u; sh < (\d+)u; \+\+sh\)", l)
|
|
if m2:
|
|
reps = int(m2.group(1))
|
|
total = 1 + 8 * (len(main_labels) + reps * len(shadow_labels))
|
|
with (out / "trace_steps.txt").open("w") as f:
|
|
f.write(f"total {total}\nmain {len(main_labels)}\nshadow {len(shadow_labels)}\nreps {reps}\n")
|
|
for x in main_labels:
|
|
f.write(f"m {x}\n")
|
|
for x in shadow_labels:
|
|
f.write(f"{x}\n")
|
|
print(f"trace pack {out}: {len(main_labels)} main, {len(shadow_labels)} shadow x {reps} reps, {total} snapshots of 8 registers ({total * 32} bytes)")
|