lusoris · Sep 29, 2026 · Sep 29, 2026 · Sep 29, 2026
@@ -183,6 +183,21 @@
   does not have; it lists `enable_db` and `clip_db`
   ([SSIM](docs/metrics/ssim.md#options)).
 
+
+- **`motion_sycl` matches the CPU `motion` exactly.** The SYCL twin blurred
+  each frame and differenced the blurred frames, while the CPU (since the
+  upstream pipelined-motion port) blurs the frame difference and rounds after
+  each filter pass. The two orders round differently, so `motion2` was up to
+  2.0e-4 off on 17x17 frames, 1.3e-5 on the Netflix 576x324 pair and 5.6e-6
+  at 4K. `motion_sycl` and `motion_v2_sycl` now run one kernel with the CPU's
+  arithmetic and agree with the CPU bit for bit at every size and bit depth
+  tested, on an Arc B580 and a UHD 770 (ADR-1371). The 4K motion step costs
+  about 11% more device time on both GPUs. With `motion_add_uv=true`,
+  `motion_sycl` no longer waits on the device inside `submit()`: the U and V
+  planes are staged in pinned memory and uploaded on the compute queue, which
+  cuts host time per 4K frame from 5.4 to 0.6 ms on a UHD 770
+  ([SYCL backend](docs/backends/sycl/overview.md#motion_sycl-matches-the-cpu-motion-exactly-2026-09-29)).
+
 ## [1.0.0-rc.2] - 2026-09-28
 ### Changed
 

@@ -0,0 +1,13 @@
+- **`motion_sycl` matches the CPU `motion` exactly.** The SYCL twin blurred
+  each frame and differenced the blurred frames, while the CPU (since the
+  upstream pipelined-motion port) blurs the frame difference and rounds after
+  each filter pass. The two orders round differently, so `motion2` was up to
+  2.0e-4 off on 17x17 frames, 1.3e-5 on the Netflix 576x324 pair and 5.6e-6
+  at 4K. `motion_sycl` and `motion_v2_sycl` now run one kernel with the CPU's
+  arithmetic and agree with the CPU bit for bit at every size and bit depth
+  tested, on an Arc B580 and a UHD 770 (ADR-1371). The 4K motion step costs
+  about 11% more device time on both GPUs. With `motion_add_uv=true`,
+  `motion_sycl` no longer waits on the device inside `submit()`: the U and V
+  planes are staged in pinned memory and uploaded on the compute queue, which
+  cuts host time per 4K frame from 5.4 to 0.6 ms on a UHD 770
+  ([SYCL backend](docs/backends/sycl/overview.md#motion_sycl-matches-the-cpu-motion-exactly-2026-09-29)).
@@ -269,7 +269,7 @@ lookups back into the runtime test. See Research-2091.
   `integer_motion_v2.metal::mv2_mirror` uses iterated reflect-101
   `idx = (idx < 0) ? -idx : 2 * (sup - 1) - idx`, bit-identical to CPU
   `integer_motion_v2.c::mirror`, CUDA `motion_v2_score.cu::mv2_mirror`,
-  SYCL `integer_motion_v2_sycl.cpp::dev_mirror_mv2`, and HIP
+  SYCL `integer_motion_pipeline_sycl.cpp::reflect_101`, and HIP
   `motion_v2_score.hip::mv2_mirror`. Never revert to single-bounce or
   `- 1` edge-replicating form.
 

@@ -20,7 +20,7 @@
  *
  *  Mirror padding: reflect-101 (`2 * (sup - 1) - idx`, iterated),
  *  identical to CPU integer_motion_v2.c::mirror, CUDA mv2_mirror,
- *  SYCL dev_mirror_mv2 and HIP mv2_mirror; see mv2_mirror below.
+ *  SYCL reflect_101 and HIP mv2_mirror; see mv2_mirror below.
  *
  *  Threadgroup layout: 16 × 16 threads per group, +2 pixel halo on
  *  each side → 20 × 20 shared tile. Inner pitch padded to 21 to
@@ -55,7 +55,7 @@ constant int MV2_FILTER[5] = {3571, 16004, 26386, 16004, 3571};
  *     Reflect-101 skips it: `2 * (sup - 1) - idx`. Every other backend
  *     already carries the corrected form -- CPU
  *     (integer_motion_v2.c::mirror), CUDA (fixed in PR #120 / T7-15),
- *     SYCL (integer_motion_v2_sycl.cpp::dev_mirror_mv2) and HIP
+ *     SYCL (integer_motion_pipeline_sycl.cpp::reflect_101) and HIP
  *     (integer_motion_v2/motion_v2_score.hip) -- and the SYCL fix records
  *     the measured impact of the `- 1` form as a systematic ~2.6e-3 motion
  *     drift vs CPU on every frame after the first. Metal was the last

@@ -135,12 +135,31 @@ HIP / Metal motion twins listed in Twin-update table above — same PR.
   `motion_five_frame_window=true` returns `-ENOTSUP` at `init()` with
   `WARNING` log. See [../../AGENTS.md §"motion3_score GPU contract"](../../AGENTS.md).
 
+- **Motion SAD = one shared kernel, difference first
+  (T-SYCL-MOTION-TINY-FRAME-PARITY-2026-09-29).** `motion_sycl` and
+  `motion_v2_sycl` both call `motion_sycl_pipeline::enqueue_sad()`
+  (`integer_motion_pipeline_sycl.{h,cpp}`): `sum |blur(prev - cur)|`,
+  vertical pass rounded `>> bpc`, horizontal `>> 16`, reflect-101 borders =
+  CPU `integer_motion.c` / `integer_motion_v2.c` `motion_score_pipeline_*`
+  since the Netflix a4a1492d port (PR #532). Bit-exact; gate
+  `test_sycl_motion_tiny_frames` compares with `==` (3x3 .. 1283x723, 8/10/16
+  bit). `blur(cur) - blur(prev)` rounds twice per pixel -> 2e-4 off at 17x17;
+  never reintroduce it. `prev - cur` order load-bearing (arithmetic shift
+  floors negatives). Vertical sum int32 up to 15 bpc, int64 at 16 (host picks
+  the `submit_sad<Acc>` instance). Kernel lives only in the pipeline TU
+  (Research-2090 name collision); extractor TUs hold none. `motion_sycl` keeps
+  the raw luma of the previous frame in `d_raw_y[2]` (device memcpy from the
+  shared frame after the kernel), because the shared frame buffers are
+  overwritten by the next upload. Measured cost vs the old per-frame blur
+  (4K micro-benchmark, kernel + copy): about +11% on B580 and UHD 770.
+
 - **`integer_motion_sycl.cpp::motion_add_uv` GPU contract** (ADR-0989).
-  When `motion_add_uv=true`, `submit_fex_sycl` uploads U and V plane data
-  H2D to `d_ref_u[cur_blur]` / `d_ref_v[cur_blur]` before calling
-  `vmaf_sycl_graph_submit`. `enqueue_motion_work` launches additional
-  `launch_blur_sad_fused` kernels for U and V, each writing to
-  `d_blur_u/v[cur]`, accumulating into `d_sad_u` / `d_sad_v`.
+  When `motion_add_uv=true`, `submit_fex_sycl` packs the reference U and V
+  planes into pinned host staging (`h_stage_u` / `h_stage_v`);
+  `motion_pre_graph` copies them H2D on the combined queue into
+  `d_ref_u[cur_slot]` / `d_ref_v[cur_slot]`; `enqueue_motion_work` runs the
+  shared SAD kernel on `d_ref_*[1 - cur_slot]` - `d_ref_*[cur_slot]`,
+  accumulating into `d_sad_u` / `d_sad_v`.
   `collect_fex_sycl` sums Y + U + V contributions, each normalized by
   respective plane area (`chroma_w × chroma_h` for UV in YUV420P). The
   numerical gate is the scalar fixed-point oracle in
@@ -152,14 +171,15 @@ HIP / Metal motion twins listed in Twin-update table above — same PR.
   `-ENOTSUP` with `WARNING` until their kernel ports land. On rebase:
   if upstream Netflix adds `motion_add_uv` to `integer_motion.c`, verify
   per-plane normalization formula stays consistent.
-  **Queue-sync invariant (ADR-1034)**: `vmaf_sycl_memcpy_h2d_async` submits
-  UV H2D copies to `state->queue` (primary queue), NOT same as
-  `copy_queue` (DMA engine used for Y-plane uploads). `vmaf_sycl_graph_submit`
-  barriers `combined_queue` only on `last_upload_event` from `copy_queue`.
-  So `submit_fex_sycl` calls `vmaf_sycl_queue_wait(state)` after UV copies
-  to flush primary queue before graph submission. If future PR routes UV H2D
-  through `copy_queue` and updates `last_upload_event`, `vmaf_sycl_queue_wait`
-  call can be removed in favour of GPU-side barrier — update this note then.
+  **Queue-sync invariant (T-SYCL-MOTION-ADD-UV-SUBMIT-WAIT-2026-09-29,
+  supersedes the ADR-1034 primary-queue wait)**: no host wait in `submit()`.
+  UV H2D rides the in-order combined queue in `pre_fn`, ahead of the kernels
+  (graph replay is fenced by `ext_oneapi_submit_barrier()`). Staging is
+  safe to refill in the next `submit()`: the graph fires on the last
+  extractor's submit, after staging, and this extractor's `collect()` of the
+  previous frame (`vmaf_sycl_graph_wait`) drained the copy that read it. Do
+  not move the UV copies back to `vmaf_sycl_memcpy_h2d_async` (primary queue)
+  — that needs the host wait again.
 
 - **`integer_vif_sycl.cpp` rd_stride uses ceiling division for odd widths** (ADR-1034).
   Both `launch_vif_hori_impl` (scalar/SIMD-32) and `launch_vif_fused_impl` (SIMD-16)
@@ -272,8 +292,8 @@ HIP / Metal motion twins listed in Twin-update table above — same PR.
   the USM buffer -> `UR_RESULT_ERROR_DEVICE_LOST` when the page is unmapped.
   Every single-reflection loader wraps the reflected index in
   `vmaf_sycl_tile_index()`: `integer_adm` `launch_dwt_vert_pair`, `integer_vif`
-  `dev_vert_load_tile` + `dev_fused_load_tile`, `integer_motion`
-  `motion_load_tile`, `integer_motion_v2` `mv2_load_diff`, `float_motion`
+  `dev_vert_load_tile` + `dev_fused_load_tile`, `integer_motion_pipeline`
+  `load_diff` (motion + motion_v2), `float_motion`
   `fm_load_tile`, `float_vif` `load_vif_tile`. Identity for consumed samples ->
   no score change. New tiled kernel = same wrap. Per-output reflections
   (`dev_hori_convolve_border`, float VIF decimate) are consumed-only and stay
@@ -516,6 +536,7 @@ ADR-0884 / ADR-0946 backlog must update in same PR.
 |---|---|---|---|
 | `integer_cambi_sycl.cpp` | `cambi.c` | `test_sycl_cambi_parity.c` (bit-exact, 4 frames), `test_integer_cambi_sycl.c` (smoke) | [ADR-1357](../../../../docs/adr/1357-sycl-cambi-device-resident.md) |
 | `integer_motion_sycl.cpp` (motion3) | `integer_motion.c` | `test_sycl_motion3_parity.c` | ADR-0219 |
+| `integer_motion_pipeline_sycl.cpp` (motion + motion_v2 SAD) | `integer_motion.c`, `integer_motion_v2.c` | `test_sycl_motion_tiny_frames.c` (bit-exact, 3x3 .. 1283x723) | T-SYCL-MOTION-TINY-FRAME-PARITY-2026-09-29 |
 | `integer_motion_sycl.cpp` (motion_add_uv) | `float_motion.c` | `test_sycl_motion_add_uv_parity.c` | ADR-0989 |
 | `integer_psnr_sycl.cpp` | `integer_psnr.c` | `test_sycl_psnr_parity.c` | ADR-0868 (round 1) |
 | `integer_vif_sycl.cpp` | `integer_vif.c` | `test_sycl_vif_parity.c` | ADR-0868 (round 1) |

@@ -0,0 +1,198 @@
+/**
+ *  Copyright 2016-2026 Netflix, Inc.
+ *  Copyright 2026 Lusoris
+ *  SPDX-License-Identifier: BSD-2-Clause-Patent
+ *
+ *  Motion SAD kernel shared by motion_sycl and motion_v2_sycl; see
+ *  integer_motion_pipeline_sycl.h for the arithmetic contract.
+ *
+ *  One work-item per output pixel. A work-group stages the (prev - cur)
+ *  difference of its 8x32 outputs plus the two-sample filter halo in local
+ *  memory, filters it vertically at the five columns the horizontal tap
+ *  needs, filters those horizontally, and sums |h| through a sub-group
+ *  reduction and one device atomic per work-group. Integer arithmetic
+ *  throughout, so the order of the sums cannot change the result. The copy
+ *  of `cur` the caller may ask for is a device memcpy after the kernel.
+ */
+
+#include "integer_motion_pipeline_sycl.h"
+
+#include <sycl/sycl.hpp>
+
+#include "sycl_compat.h"
+#include "sycl_tile_index.h"
+
+#include <cstddef>
+#include <cstdint>
+#include <utility>
+
+namespace motion_sycl_pipeline
+{
+namespace
+{
+
+constexpr int32_t kFilter[5] = {3571, 16004, 26386, 16004, 3571}; /* sum 65536 */
+constexpr int kWgX = 32;
+constexpr int kWgY = 8;
+constexpr int kHalf = 2;
+constexpr int kTileW = kWgX + 2 * kHalf; /* 36 */
+constexpr int kTileH = kWgY + 2 * kHalf; /* 12 */
+constexpr unsigned kGroupSize = kWgX * kWgY;
+constexpr int kMaxSubgroups = 32;
+
+/* The vertical pass sums |filter[k] * diff| <= 65536 * (2^bpc - 1), plus a
+ * rounding term of 2^(bpc - 1). That fits int32 up to 15 bits per sample;
+ * 16-bit input takes the int64 path, as the CPU's 16-bit pipeline does. */
+constexpr unsigned kInt32VerticalMaxBpc = 15;
+
+/* Reflect-101, the CPU's mirror(): -1 -> 1, extent -> extent - 2. */
+inline int reflect_101(int idx, int extent)
+{
+    if (idx < 0) {
+        return -idx;
+    }
+    if (idx >= extent) {
+        return (2 * extent) - idx - 2;
+    }
+    return idx;
+}
+
+inline int32_t read_sample(const void *plane, size_t offset, unsigned bpc)
+{
+    if (bpc <= 8) {
+        return static_cast<const uint8_t *>(plane)[offset];
+    }
+    return static_cast<const uint16_t *>(plane)[offset];
+}
+
+/* Stage prev - cur for the work-group's outputs and halo. The CPU differences
+ * in that order, and the order matters: the arithmetic shift rounds a
+ * negative sum towards minus infinity. */
+inline void load_diff(sycl::nd_item<2> item, const sycl::local_accessor<int32_t, 2> &diff,
+                      const SadArgs &args)
+{
+    const int tile_y = (int)(item.get_group(0) * kWgY) - kHalf;
+    const int tile_x = (int)(item.get_group(1) * kWgX) - kHalf;
+    const bool interior = (tile_y >= 0) && (tile_y + kTileH <= (int)args.height) && (tile_x >= 0) &&
+                          (tile_x + kTileW <= (int)args.width);
+    constexpr unsigned tile_elems = kTileH * kTileW;
+    for (unsigned i = item.get_local_linear_id(); i < tile_elems; i += kGroupSize) {
+        const unsigned row = i / kTileW;
+        const unsigned col = i % kTileW;
+        int y = tile_y + (int)row;
+        int x = tile_x + (int)col;
+        if (!interior) {
+            y = vmaf_sycl_tile_index(reflect_101(y, (int)args.height), (int)args.height);
+            x = vmaf_sycl_tile_index(reflect_101(x, (int)args.width), (int)args.width);
+        }
+        const size_t offset = ((size_t)y * args.width) + (size_t)x;
+        diff[row][col] =
+            read_sample(args.prev, offset, args.bpc) - read_sample(args.cur, offset, args.bpc);
+    }
+}
+
+/* Vertical 5-tap filter of tile column `col` around local row `ly`. */
+template <typename Acc>
+inline int32_t vertical_tap(const sycl::local_accessor<int32_t, 2> &diff, unsigned ly, unsigned col,
+                            unsigned bpc)
+{
+    const Acc sum = ((Acc)kFilter[0] * (diff[ly][col] + diff[ly + 4][col])) +
+                    ((Acc)kFilter[1] * (diff[ly + 1][col] + diff[ly + 3][col])) +
+                    ((Acc)kFilter[2] * diff[ly + 2][col]);
+    const Acc round = (Acc)1 << (bpc - 1);
+    return (int32_t)((sum + round) >> bpc);
+}
+
+/* |h| at this work-item's pixel, 0 for the padding work-items. `Acc` is the
+ * vertical accumulator; the host picks it per launch, so the kernel carries
+ * no per-pixel width test. */
+template <typename Acc>
+inline int64_t filtered_abs(sycl::nd_item<2> item, const sycl::local_accessor<int32_t, 2> &diff,
+                            const SadArgs &args)
+{
+    const int x = (int)item.get_global_id(1);
+    const int y = (int)item.get_global_id(0);
+    if (!std::cmp_less(x, args.width) || !std::cmp_less(y, args.height)) {
+        return 0;
+    }
+    const unsigned lx = item.get_local_id(1);
+    const unsigned ly = item.get_local_id(0);
+    int32_t v[5];
+#pragma unroll
+    for (unsigned hx = 0; hx < 5; hx++) {
+        v[hx] = vertical_tap<Acc>(diff, ly, lx + hx, args.bpc);
+    }
+    const int64_t sum = ((int64_t)kFilter[0] * (v[0] + v[4])) +
+                        ((int64_t)kFilter[1] * (v[1] + v[3])) + ((int64_t)kFilter[2] * v[2]);
+    const int64_t h = (sum + 32768) >> 16;
+    return (h < 0) ? -h : h;
+}
+
+inline void reduce_sad(sycl::nd_item<2> item, const sycl::local_accessor<int64_t, 1> &scratch,
+                       int64_t value, int64_t *sad)
+{
+    const sycl::sub_group subgroup = item.get_sub_group();
+    const int64_t subgroup_sum = sycl::reduce_over_group(subgroup, value, sycl::plus<int64_t>{});
+    if (subgroup.get_local_linear_id() == 0) {
+        scratch[subgroup.get_group_linear_id()] = subgroup_sum;
+    }
+    item.barrier(sycl::access::fence_space::local_space);
+    if (item.get_local_linear_id() == 0) {
+        int64_t total = 0;
+        const uint32_t subgroups = subgroup.get_group_linear_range();
+        for (uint32_t i = 0; i < subgroups; i++) {
+            total += scratch[i];
+        }
+        const sycl::atomic_ref<int64_t, sycl::memory_order::relaxed, sycl::memory_scope::device,
+                               sycl::access::address_space::global_space>
+            output(*sad);
+        output.fetch_add(total);
+    }
+}
+
+size_t plane_bytes(const SadArgs &args)
+{
+    return (size_t)args.width * args.height * ((args.bpc <= 8) ? 1U : 2U);
+}
+
+template <typename Acc> void submit_sad(sycl::queue &queue, const SadArgs &args)
+{
+    const size_t global_h = (((size_t)args.height + kWgY - 1) / kWgY) * kWgY;
+    const size_t global_w = (((size_t)args.width + kWgX - 1) / kWgX) * kWgX;
+    const sycl::nd_range<2> range({global_h, global_w}, {(size_t)kWgY, (size_t)kWgX});
+
+    queue.submit([&](sycl::handler &cgh) {
+        const sycl::local_accessor<int32_t, 2> diff(sycl::range<2>(kTileH, kTileW), cgh);
+        const sycl::local_accessor<int64_t, 1> scratch(sycl::range<1>(kMaxSubgroups), cgh);
+        const SadArgs kernel_args = args;
+        cgh.parallel_for(range, [=](sycl::nd_item<2> item) VMAF_SYCL_REQD_SG_SIZE(32) {
+            load_diff(item, diff, kernel_args);
+            item.barrier(sycl::access::fence_space::local_space);
+            const int64_t value = filtered_abs<Acc>(item, diff, kernel_args);
+            reduce_sad(item, scratch, value, kernel_args.sad);
+        });
+    });
+}
+
+} // namespace
+
+void enqueue_sad(sycl::queue &queue, const SadArgs &args)
+{
+    if (args.bpc <= kInt32VerticalMaxBpc) {
+        submit_sad<int32_t>(queue, args);
+    } else {
+        submit_sad<int64_t>(queue, args);
+    }
+    /* A plain copy, not a store in the kernel: one more message per pixel
+     * there measured no faster on a UHD 770 or an Arc B580. */
+    enqueue_copy(queue, args);
+}
+
+void enqueue_copy(sycl::queue &queue, const SadArgs &args)
+{
+    if (args.cur_copy != nullptr) {
+        queue.memcpy(args.cur_copy, args.cur, plane_bytes(args));
+    }
+}
+
+} // namespace motion_sycl_pipeline