Native Metal SSB redraw precision at the lowest scan frequencies, 2026-10-25#
Question#
Is the native interactive SSB redraw (ssb-redraw, method ssb, 1–4×) within
twice the CUDA reference’s own float32 error, measured against a float64
evaluation of the same definition, on every measured acquisition, fitted optics
set and output sampling, at the same drag speed?
Workload and boundary#
Five measured 512 × 512 × 192 × 192 acquisitions at 300 kV and 30 mrad, with 8,800–9,300 BF pixels: three at a 0.373 Å scan step (Auto 3×) and two at 0.525 Å (Auto 4×). Eight fitted optics sets (C10 … C56, sample tilt, depth spread), each reconstructed at 1×, 2× and Auto: 24 images.
References: the CUDA SSB (float32 corrections, error-free float32 sums met in float64, complex128 inverse FFT) and an independent float64 evaluation of the definition (float64 FFT of the counts, exact chi, complex128 sums and inverse FFT). e_C = rms(CUDA − float64), per image, 3.2e-6 to 8.7e-6 rad.
Runner:
metal-ssb-fit-benchmark --fixture-parity CASE --oracle-dir DIR --methods ssb(rms and maximum of the wrapped phase difference); timing withmetal-ssb-redraw-benchmarkas in the 2026-10-24 redraw report (synthetic counts on the 0.373 Å geometry, every BF pixel, all 14 terms, tilt, depth; 20 warm calls with a changed C10). Apple M5 Max, macOS 26.4. Before = the same runners built fromf829cab4with the same release flags; before and after alternated run by run under the shared GPU lock. Raw outputs are in the private evidence archive.
Finding#
v0.2 was 2.28 and 2.45 e_C from the float64 definition at 1× on the two 0.525 Å acquisitions (1.36–1.51 e_C at 1× on two others). The whole excess sat at the first scan frequencies (frequency index ±1): 8.1e-6 rad of error power there against CUDA’s 2.5e-6, every other frequency band equal to CUDA’s.
There the two sidebands nearly coincide. With
δ− = chi(k − q) − chi(k), δ+ = chi(k + q) − chi(k), μ = (δ+ − δ−)/2 and
ν = −(δ− + δ+)/2, the sideband difference is
gamma = exp(iμ) ((a − b) cos ν + i (a + b) sin ν), so the correction’s
direction is set by μ and by the sign of the second difference
chi(k + q) + chi(k − q) − 2 chi(k). At index 1 of a 0.525 Å scan
(q = 0.0037 / Å) that second difference is a few 1e-5 rad when C10 and C12
partly cancel along the scan direction, while chi reaches tens of radians.
v0.2 formed δ± from separately rounded float32 chi values (a row-collapsed
absolute polynomial minus chi(k) from the probe kernel), errors of several
1e-6 rad, so the sign flipped for a sizable fraction of BF pixels, each flip
adding 2 |G(q, k)|. The 0.373 Å scans have twice the q² at index 1 and cancel
less. Trigonometry was not the cause: on the M5 Max metal::fast::sincos and
metal::precise::sincos return identical values (within 1.3e-7 of double for
|x| < 100), and swapping them changed no output bit. Accumulation order and the
float32 inverse FFT contribute equally to both implementations at every other
frequency.
Change (v0.3)#
redraw_taylor(once per redraw, one thread per BF pixel) shifts the absolute 7 × 7 chi polynomial to each pixel: δ±(q) = Σ d_ij(k) (±q_row)^i (±q_col)^j, stored coefficient-major; it also writes the pixel’s probe geometry, replacing the probe dispatch. The threadgroup collapses the row power at ±q_row (plus and minus from one read) and Horner runs in ±q_col. Rounding is now proportional to δ± (about 1e-8 rad at index 1) instead of to chi, and chi(k) is no longer subtracted.Pairs with both sidebands possible form gamma from μ and ν as above, so the rounding of sin and cos is not amplified by 1/|gamma| where the sidebands nearly cancel. Single-sideband pairs, accumulation, mirror, origin and FFT are unchanged.
Results#
Worst over the optics sets of each acquisition; native − float64 / e_C and native − CUDA / e_C (rms). The rule is ≤ 2.
Acquisition |
Step |
1× v0.2 |
1× v0.3 |
2× v0.2 |
2× v0.3 |
Auto v0.2 |
Auto v0.3 |
|---|---|---|---|---|---|---|---|
1 (three optics sets) |
0.373 Å |
1.20 / 1.08 |
0.95 / 1.11 |
1.13 / 1.00 |
1.07 / 1.01 |
1.13 / 1.00 |
1.08 / 1.01 |
2 |
0.373 Å |
1.36 / 1.61 |
0.96 / 1.13 |
1.14 / 1.26 |
0.98 / 1.06 |
1.14 / 1.26 |
0.99 / 1.06 |
3 |
0.373 Å |
1.51 / 1.67 |
0.81 / 1.01 |
1.18 / 1.30 |
0.90 / 1.03 |
1.17 / 1.29 |
0.90 / 1.02 |
4 |
0.525 Å |
2.28 / 2.20 |
0.74 / 1.07 |
1.45 / 1.40 |
0.93 / 0.97 |
1.34 / 1.25 |
0.95 / 0.93 |
5 (two optics sets) |
0.525 Å |
2.45 / 2.53 |
0.62 / 1.11 |
1.57 / 1.62 |
0.82 / 1.00 |
1.40 / 1.44 |
0.85 / 0.97 |
All 24 v0.3 images: native − float64 at most 1.08 e_C, native − CUDA at most
1.13 e_C, maximum absolute differences at most 0.57 × 10 e_C. At index ±1 the
error power fell from 8.1e-6 to 1.5e-7 rad (acquisition 4) and from 9.5e-6 to
4.5e-7 (acquisition 5), 7–17 times below CUDA’s; the remaining error is the
high-frequency float32 floor, equal to CUDA’s. Synthetic check
(testRedrawResolvesNearlyCancellingSidebands, C10 and C12 balanced along the
scan rows, |gamma| about 2e-5 at index 1): v0.2 4.5e-3 rad rms against the
double-precision definition, v0.3 3.7e-8.
Redraw time per frame, median of the per-run medians / median p90, v0.2 → v0.3, runs alternating on the final builds (other agents’ GPU work ran between runs):
Scan |
Order |
Output |
Rounds |
v0.2 |
v0.3 |
Paired difference |
|---|---|---|---|---|---|---|
512² |
fifth |
3× |
12 |
30.5 / 31.3 |
30.8 / 31.4 |
+0.1 ms |
512² |
low |
3× |
8 |
40.3 (loaded; minimum 29.4) |
39.9 (minimum 29.8) |
−0.1 ms |
256² |
fifth |
3× |
8 |
10.1 / 10.3 |
10.3 / 10.7 |
+0.3 ms |
512² |
fifth |
1× |
6 |
23.9 |
23.7 |
+0.1 ms |
At 512² the difference is within the run-to-run noise; at 256², where one wave of threadgroups exposes the per-batch collapse from device memory (constant memory before), v0.3 costs about 3 %. The 512² fifth-order 3× drag stays within 33 ms per frame.
Conclusion#
The native SSB redraw now meets the 2× rule everywhere: below CUDA’s own error on 20 of 24 images and within 8 % of it on the rest, at the same drag speed at 512² (about 3 % slower at 256²). v0.3 replaces v0.2 with no fallback. The in-app headed drag timing is still to be recorded.