From ecd67ec44f5439e3fba28aa51eded601d887baa9 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Thu, 22 Oct 2026 12:53:37 +0200 Subject: [PATCH] Metal integration engine: back to its panel-s5 kernels; Mac stills keep the old clip, no saturation count The Metal kernel and host code are restored to their panel-s5 content, so neither the background clip that follows a handed-over radial curve nor the count of saturated pixels in pass A runs on the Metal path. This is an accepted, temporary platform difference, to be closed once the kernels can be checked on a Mac: on macOS a still's background clip stays the plain one (the radial correction itself still uses the handed-over curve, as on the CPU), and an overloaded observation kept as a lower bound counts its saturated pixels as 0 instead of at the saturation limit. One host-side change is kept, with no effect on any result: the radial sums are accumulated on the GPU only where no curve was handed over. StartPassA now sizes those sums for that case alone, so the old host code would write past them on every still with a curve. The shared structs need nothing on the Metal side: BraggPassA.h is compiled into the kernel source, the box-sum kernel zero-initialises BraggBoxSums (n_saturated = 0) and the host zero-initialises BraggPassAFrame (n_curve = 0). Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_015CeEGgpKdfkCTJb3YNMJcu --- .../BraggIntegrationEngineMetal.metal | 12 +----------- .../bragg_integration/BraggIntegrationEngineMetal.mm | 12 ++---------- 2 files changed, 3 insertions(+), 21 deletions(-) diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineMetal.metal b/image_analysis/bragg_integration/BraggIntegrationEngineMetal.metal index 2489bb5ba..46d12341f 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineMetal.metal +++ b/image_analysis/bragg_integration/BraggIntegrationEngineMetal.metal @@ -125,8 +125,6 @@ kernel void box_sums(device const float2 *xy [[buffer(0)]], && ((static_mask[(uint) idx / 32] >> ((uint) idx % 32)) & 1u); if (v == PX_SATURATED || (v == PX_MASKED && on && !masked)) s.overloaded = 1; - if (v == PX_SATURATED) - ++s.n_saturated; if (!valid_px(v)) continue; s.I_sum += v; s.I_sum_x += (int64_t) x * v; @@ -169,7 +167,6 @@ kernel void clip_sums(device const float2 *xy [[buffer(0)]], device atomic_uint *rad_sum [[buffer(8)]], device atomic_int *rad_cnt [[buffer(9)]], device const BraggPanel *panels [[buffer(10)]], - device const float *curve [[buffer(11)]], uint i [[thread_position_in_grid]]) { if (i >= f.n || !clips[i]) return; @@ -182,20 +179,13 @@ kernel void clip_sums(device const float2 *xy [[buffer(0)]], const int y0 = max(0, (int) floor(py - st.ey_out - 1.0f)); const int y1 = min(f.height - 1, (int) ceil(py + st.ey_out + 1.0f)); const int32_t t = thr[i]; - const int b0 = f.n_curve > 0 ? clamp((int) round(st.r0), 0, f.n_curve - 1) : 0; BraggClipSums c = {}; for (int y = y0; y <= y1; ++y) for (int x = x0; x <= x1; ++x) { const BraggStencilDist d = BraggStencilDistances(st, x - px, y - py); const int idx = y * f.width + x; const int32_t v = BraggPanelHolds(pn, x, y) ? img[idx] : PX_MASKED; - // Where the frame's radial curve is in use the clip follows it (RadialClipOffset). - int64_t lim = t; - if (f.n_curve > 0) { - const int b = clamp((int) round(st.r0 + d.rad), 0, f.n_curve - 1); - lim += (int64_t) round(curve[b] - curve[b0]); - } - if (!(d.inner >= f.r2_sq && d.outer < f.r3_sq && valid_px(v) && mask[idx] != MASK_FLUX && (int64_t) v <= lim)) + if (!(d.inner >= f.r2_sq && d.outer < f.r3_sq && valid_px(v) && mask[idx] != MASK_FLUX && v <= t)) continue; c.sum += v; ++c.n; diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineMetal.mm b/image_analysis/bragg_integration/BraggIntegrationEngineMetal.mm index e3c5d0ab4..c5cf46860 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineMetal.mm +++ b/image_analysis/bragg_integration/BraggIntegrationEngineMetal.mm @@ -107,7 +107,6 @@ struct BraggIntegrationEngineMetal::Device { id unmark = nil; bool clear = false; id rad_sum = nil, rad_cnt = nil; // the radial background curve, as the clip kernel adds it up - id curve = nil; // the curve handed over, which the clip follows // Per reflection, grown with the frames. size_t capacity = 0; id xy = nil, panels = nil, level = nil, sums = nil, clips = nil, thr = nil, clip_out = nil; @@ -221,16 +220,10 @@ std::vector BraggIntegrationEngineMetal::Run(const ImagePreprocessor f.inv_claim = inv_claim; f.use_owner = use_owner; f.exclude = overlap == OverlapMode::Exclude && mode != IntegratorMode::BoxSum; - // The curve is added up here only where no ring profile was handed over (RadialBackgroundFromRings). + // The curve is added up on the GPU only where none was handed over (RadialBackgroundFromRings): + // StartPassA sizes its sums for that case alone. The clip does not follow a handed-over curve here. const bool accumulate_radial = bkg_radial && radial_curve.empty(); f.n_rad = accumulate_radial ? a.n_rad : 0; - // The clip follows the handed-over curve where the correction uses it (RadialClipOffset). - f.n_curve = bkg_radial && !radial_curve.empty() ? static_cast(radial_curve.size()) : 0; - if (f.n_curve > 0) { - if (dev->curve == nil || dev->curve.length < radial_curve.size() * sizeof(float)) - dev->curve = buffer(radial_curve.size() * sizeof(float)); - std::copy(radial_curve.begin(), radial_curve.end(), static_cast(dev->curve.contents)); - } f.static_mask_words = static_cast(static_mask.size()); f.n = static_cast(npredicted); @@ -319,7 +312,6 @@ std::vector BraggIntegrationEngineMetal::Run(const ImagePreprocessor [enc setBuffer:(accumulate_radial ? dev->rad_sum : dev->clips) offset:0 atIndex:8]; // not written without the curve [enc setBuffer:(accumulate_radial ? dev->rad_cnt : dev->clips) offset:0 atIndex:9]; [enc setBuffer:dev->panels offset:0 atIndex:10]; - [enc setBuffer:(f.n_curve > 0 ? dev->curve : dev->thr) offset:0 atIndex:11]; // not read without a curve [enc dispatchThreads:grid threadsPerThreadgroup:group]; [enc endEncoding]; Wait(cb);