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);