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 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_015CeEGgpKdfkCTJb3YNMJcu

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;

image_analysis/bragg_integration/BraggIntegrationEngineMetal.mm

@@ -107,7 +107,6 @@ struct BraggIntegrationEngineMetal::Device {
     id<MTLCommandBuffer> unmark = nil;
     bool clear = false;
     id<MTLBuffer> rad_sum = nil, rad_cnt = nil;   // the radial background curve, as the clip kernel adds it up
-    id<MTLBuffer> curve = nil;                    // the curve handed over, which the clip follows
     // Per reflection, grown with the frames.
     size_t capacity = 0;
     id<MTLBuffer> xy = nil, panels = nil, level = nil, sums = nil, clips = nil, thr = nil, clip_out = nil;
@@ -221,16 +220,10 @@ std::vector<Reflection> 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<int32_t>(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<float *>(dev->curve.contents));
-        }
         f.static_mask_words = static_cast<uint32_t>(static_mask.size());
         f.n = static_cast<uint32_t>(npredicted);
 
@@ -319,7 +312,6 @@ std::vector<Reflection> 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);