fix(sycl): score 4:0:0 input in psnr_hvs_sycl and psnr_hvs_hip

CHANGELOG.md

@@ -1240,6 +1240,19 @@
   [the psnr_hvs page](docs/metrics/psnr-hvs.md#agreement-with-the-cpu-extractor).
 
 
+- **SYCL and HIP `psnr_hvs` score 4:0:0 input, and the HIP twin takes
+  `enable_chroma`** (`T-SYCL-HIP-PSNR-HVS-YUV400-REFUSED-2026-10-01`).
+  `psnr_hvs_sycl` and `psnr_hvs_hip` refused 4:0:0 pictures at `init()`
+  (`YUV400P unsupported`), where the CPU extractor and `psnr_hvs_cuda` score the
+  luma plane and emit `psnr_hvs_y` and `psnr_hvs`. Both twins now do the same.
+  `psnr_hvs_hip` also gains the CPU extractor's `enable_chroma` option (default
+  `true`): with `false` it uploads, dispatches and scores luma only, like the
+  CUDA and SYCL twins. Scores for 4:0:0 and for `enable_chroma=false` are
+  identical to the CPU's at `--precision max`, measured on an Arc A380, a
+  gfx1036 and an RTX 4090. See
+  [the psnr_hvs page](docs/metrics/psnr-hvs.md#sample-conversion).
+
+
 - **SYCL: kernels stay out of scratch memory, and a wrong-result driver is
   reported.** On an Arc A380 under the Linux xe driver, a SYCL kernel that keeps
   a private array in memory or spills registers returns wrong values, with no

changelog.d/fixed/sycl-hip-psnr-hvs-yuv400.md

@@ -0,0 +1,11 @@
+- **SYCL and HIP `psnr_hvs` score 4:0:0 input, and the HIP twin takes
+  `enable_chroma`** (`T-SYCL-HIP-PSNR-HVS-YUV400-REFUSED-2026-10-01`).
+  `psnr_hvs_sycl` and `psnr_hvs_hip` refused 4:0:0 pictures at `init()`
+  (`YUV400P unsupported`), where the CPU extractor and `psnr_hvs_cuda` score the
+  luma plane and emit `psnr_hvs_y` and `psnr_hvs`. Both twins now do the same.
+  `psnr_hvs_hip` also gains the CPU extractor's `enable_chroma` option (default
+  `true`): with `false` it uploads, dispatches and scores luma only, like the
+  CUDA and SYCL twins. Scores for 4:0:0 and for `enable_chroma=false` are
+  identical to the CPU's at `--precision max`, measured on an Arc A380, a
+  gfx1036 and an RTX 4090. See
+  [the psnr_hvs page](docs/metrics/psnr-hvs.md#sample-conversion).

core/src/feature/hip/AGENTS.md

@@ -249,8 +249,14 @@ Load-bearing, each one breaks bit-identity on its own:
 - no sum of terms in kernel or host TU (per-block partials round differently:
   1e-2 dB at 3840x2160).
 
+Plane count = `psnr_hvs_plane_count(s)` (`s->n_planes` clamped to 3): 1 for
+`enable_chroma=false` or 4:0:0 (CPU `psnr_hvs.c::init` rule), else 3. Staging,
+uploads, kernel `args.n_planes`, scores and emitted features all loop to it; a
+fixed `PSNR_HVS_NUM_PLANES` loop there re-breaks 4:0:0 (reads `data[1]` of a
+luma-only picture).
+
 Guards: `test_hip_psnr_hvs_parity{,_large}` (device, `==` on all four outputs,
-3840x2160 included), `test_psnr_hvs_twin_exact_sum_contract.py` and
+3840x2160, 4:0:0 and `enable_chroma=false` included), `test_psnr_hvs_twin_exact_sum_contract.py` and
 `test_psnr_hvs_score` (device-free). Gate cell = tolerance 0 (`EXACT_TWINS`,
 `scripts/ci/cross_backend_calibration.py`). Upstream change to
 `calc_psnrhvs()` arithmetic or order -> mirror it here in the same PR.

core/src/feature/hip/integer_psnr_hvs_hip.c

@@ -16,7 +16,10 @@
  *    not wait for the device.
  *  - collect() waits once and emits `psnr_hvs_y/cb/cr` and
  *    `psnr_hvs = 0.8*Y + 0.1*(Cb + Cr)`.
- *  - Rejects YUV400P (no chroma) and bpc > 12.
+ *  - With `enable_chroma=false`, or for 4:0:0 input, only the luma plane is
+ *    staged, dispatched and scored, and `psnr_hvs` is the luma score, as in
+ *    the CPU extractor (third_party/xiph/psnr_hvs.c::init).
+ *  - Rejects bpc > 12.
  *
  *  ADR-1397: the scores are the CPU extractor's bit for bit. The kernel
  *  stores the 64 masked coefficient errors of every block, computed in the
@@ -33,6 +36,7 @@
  */
 
 #include <errno.h>
+#include <stdbool.h>
 #include <stddef.h>
 #include <stdint.h>
 #include <string.h>
@@ -82,6 +86,12 @@ typedef struct PsnrHvsStateHip {
     unsigned total_blocks;
     size_t row_bytes[PSNR_HVS_NUM_PLANES];
     unsigned bpc;
+    /* enable_chroma: when false, only the luma plane is scored. Default true,
+     * as in the CPU extractor. */
+    bool enable_chroma;
+    /* Planes staged, dispatched and scored: 1 when enable_chroma is false or
+     * the input is 4:0:0, else 3. */
+    unsigned n_planes;
 
 #ifdef HAVE_HIPCC
     hipModule_t module;
@@ -103,7 +113,25 @@ typedef struct PsnrHvsStateHip {
     VmafDictionary *feature_name_dict;
 } PsnrHvsStateHip;
 
-static const VmafOption options[] = {{0}};
+/* Planes the twin stages, dispatches and scores: s->n_planes, never more than
+ * the three the state's arrays hold. */
+static unsigned psnr_hvs_plane_count(const PsnrHvsStateHip *s)
+{
+    return s->n_planes < (unsigned)PSNR_HVS_NUM_PLANES ? s->n_planes :
+                                                         (unsigned)PSNR_HVS_NUM_PLANES;
+}
+
+static const VmafOption options[] = {
+    {
+        .name = "enable_chroma",
+        .help = "enable psnr_hvs calculation for chroma channels (Cb and Cr); "
+                "when false only the luma plane is scored and psnr_hvs equals "
+                "psnr_hvs_y",
+        .offset = offsetof(PsnrHvsStateHip, enable_chroma),
+        .type = VMAF_OPT_TYPE_BOOL,
+        .default_val.b = true,
+    },
+    {0}};
 
 #ifdef HAVE_HIPCC
 static int psnr_hvs_hip_rc(hipError_t rc)
@@ -197,7 +225,7 @@ static int psnr_hvs_alloc_pair(void **device, void **pinned, size_t bytes)
 static int psnr_hvs_alloc_buffers(PsnrHvsStateHip *s)
 {
     int err = 0;
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES && !err; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s) && !err; p++) {
         const size_t bytes = s->row_bytes[p] * s->height[p];
         err = psnr_hvs_alloc_pair(&s->d_ref[p], &s->h_ref[p], bytes);
         if (!err)
@@ -242,13 +270,22 @@ static int psnr_hvs_unwind_module(PsnrHvsStateHip *s, int err)
 }
 #endif /* HAVE_HIPCC */
 
-/* Per-plane dimensions for `pix_fmt`, by picture.c's ceil rule. */
+/* Per-plane dimensions for `pix_fmt`, by picture.c's ceil rule, and the number
+ * of planes to score. */
 static int psnr_hvs_set_plane_dims(PsnrHvsStateHip *s, enum VmafPixelFormat pix_fmt, unsigned w,
                                    unsigned h)
 {
     s->width[0] = w;
     s->height[0] = h;
+    /* 4:0:0 has no chroma planes: luma only whatever enable_chroma says, as in
+     * the CPU extractor (third_party/xiph/psnr_hvs.c::init). */
+    s->n_planes =
+        (s->enable_chroma && pix_fmt != VMAF_PIX_FMT_YUV400P) ? (unsigned)PSNR_HVS_NUM_PLANES : 1u;
     switch (pix_fmt) {
+    case VMAF_PIX_FMT_YUV400P:
+        s->width[1] = s->width[2] = 0u;
+        s->height[1] = s->height[2] = 0u;
+        break;
     case VMAF_PIX_FMT_YUV420P:
         s->width[1] = s->width[2] = (w + 1u) >> 1;
         s->height[1] = s->height[2] = (h + 1u) >> 1;
@@ -268,8 +305,8 @@ static int psnr_hvs_set_plane_dims(PsnrHvsStateHip *s, enum VmafPixelFormat pix_
     return 0;
 }
 
-/* Block grid of every plane and the planes' offsets in the one term buffer
- * the single dispatch writes. */
+/* Block grid of every active plane and the planes' offsets in the one term
+ * buffer the single dispatch writes. */
 static int psnr_hvs_set_plane_geometry(PsnrHvsStateHip *s, enum VmafPixelFormat pix_fmt, unsigned w,
                                        unsigned h)
 {
@@ -279,10 +316,10 @@ static int psnr_hvs_set_plane_geometry(PsnrHvsStateHip *s, enum VmafPixelFormat
 
     const size_t bytes_per_sample = (s->bpc > 8u) ? 2u : 1u;
     s->total_blocks = 0u;
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         if (s->width[p] < (unsigned)PSNR_HVS_BLOCK || s->height[p] < (unsigned)PSNR_HVS_BLOCK) {
             vmaf_log(VMAF_LOG_LEVEL_ERROR,
-                     "psnr_hvs_hip: plane %d dims %ux%u smaller than 8x8 block\n", p, s->width[p],
+                     "psnr_hvs_hip: plane %u dims %ux%u smaller than 8x8 block\n", p, s->width[p],
                      s->height[p]);
             return -EINVAL;
         }
@@ -296,19 +333,13 @@ static int psnr_hvs_set_plane_geometry(PsnrHvsStateHip *s, enum VmafPixelFormat
     return 0;
 }
 
-static int psnr_hvs_validate_input(enum VmafPixelFormat pix_fmt, unsigned bpc, unsigned w,
-                                   unsigned h)
+static int psnr_hvs_validate_input(unsigned bpc, unsigned w, unsigned h)
 {
     if (bpc > 12u) {
         vmaf_log(VMAF_LOG_LEVEL_ERROR, "psnr_hvs_hip: invalid bitdepth (%u); bpc must be <= 12\n",
                  bpc);
         return -EINVAL;
     }
-    if (pix_fmt == VMAF_PIX_FMT_YUV400P) {
-        vmaf_log(VMAF_LOG_LEVEL_ERROR,
-                 "psnr_hvs_hip: YUV400P unsupported (psnr_hvs needs all 3 planes)\n");
-        return -EINVAL;
-    }
     if (w < (unsigned)PSNR_HVS_BLOCK || h < (unsigned)PSNR_HVS_BLOCK) {
         vmaf_log(VMAF_LOG_LEVEL_ERROR, "psnr_hvs_hip: input %ux%u smaller than 8x8 block\n", w, h);
         return -EINVAL;
@@ -321,7 +352,7 @@ static int init_fex_hip(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt,
 {
     PsnrHvsStateHip *s = fex->priv;
 
-    int err = psnr_hvs_validate_input(pix_fmt, bpc, w, h);
+    int err = psnr_hvs_validate_input(bpc, w, h);
     if (err != 0)
         return err;
 
@@ -366,7 +397,7 @@ static int psnr_hvs_picture_matches(const PsnrHvsStateHip *s, const VmafPicture
 {
     if (pic == NULL || pic->bpc != s->bpc)
         return 0;
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         if (pic->data[p] == NULL || pic->w[p] != s->width[p] || pic->h[p] != s->height[p])
             return 0;
     }
@@ -377,7 +408,8 @@ static int psnr_hvs_picture_matches(const PsnrHvsStateHip *s, const VmafPicture
  * caller when submit() returns, so the device copy must not read it
  * (core/src/feature/hip/AGENTS.md, "Picture uploads"); the staging buffer is
  * the extractor's until the next submit(), which follows collect(). */
-static void psnr_hvs_stage_plane(const PsnrHvsStateHip *s, const VmafPicture *pic, int p, void *dst)
+static void psnr_hvs_stage_plane(const PsnrHvsStateHip *s, const VmafPicture *pic, unsigned p,
+                                 void *dst)
 {
     const uint8_t *src = (const uint8_t *)pic->data[p];
     uint8_t *out = (uint8_t *)dst;
@@ -393,7 +425,7 @@ static void psnr_hvs_stage_plane(const PsnrHvsStateHip *s, const VmafPicture *pi
 
 static int psnr_hvs_enqueue_uploads(const PsnrHvsStateHip *s, hipStream_t str)
 {
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         const size_t bytes = s->row_bytes[p] * s->height[p];
         hipError_t rc = hipMemcpyAsync(s->d_ref[p], s->h_ref[p], bytes, hipMemcpyHostToDevice, str);
         if (rc == hipSuccess)
@@ -415,15 +447,15 @@ static int psnr_hvs_enqueue_frame(const PsnrHvsStateHip *s)
 
     struct PsnrHvsHipKernelArgs args;
     (void)memset(&args, 0, sizeof(args));
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         args.plane[p].ref = s->d_ref[p];
         args.plane[p].dist = s->d_dist[p];
         args.plane[p].width = s->width[p];
         args.plane[p].blocks_x = s->num_blocks_x[p];
         args.plane[p].first_block = s->first_block[p];
     }
     args.terms = s->d_terms;
-    args.n_planes = PSNR_HVS_NUM_PLANES;
+    args.n_planes = psnr_hvs_plane_count(s);
     args.total_blocks = s->total_blocks;
     args.wide = (s->bpc > 8u) ? 1u : 0u;
     void *params[] = {&args};
@@ -450,7 +482,7 @@ static int submit_fex_hip(VmafFeatureExtractor *fex, VmafPicture *ref_pic, VmafP
 #ifdef HAVE_HIPCC
     if (!psnr_hvs_picture_matches(s, ref_pic) || !psnr_hvs_picture_matches(s, dist_pic))
         return -EINVAL;
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         psnr_hvs_stage_plane(s, ref_pic, p, s->h_ref[p]);
         psnr_hvs_stage_plane(s, dist_pic, p, s->h_dist[p]);
     }
@@ -494,25 +526,25 @@ static void psnr_hvs_plane_scores(const PsnrHvsStateHip *s, double plane_score[]
     // SAFETY: s->h_terms holds PSNR_HVS_HIP_TERMS floats for each of the
     // s->total_blocks blocks, and s->first_block[p] + s->num_blocks[p] <=
     // s->total_blocks holds by construction (psnr_hvs_set_plane_geometry).
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         const float *plane_terms =
             s->h_terms + ((size_t)s->first_block[p] * (size_t)PSNR_HVS_HIP_TERMS);
         plane_score[p] = vmaf_psnr_hvs_plane_score(plane_terms, s->num_blocks[p], s->bpc);
     }
 }
 
 /* The CPU extract() expressions: dB per plane, then the weighted score. */
-static int psnr_hvs_append_scores(VmafFeatureCollector *feature_collector,
+static int psnr_hvs_append_scores(VmafFeatureCollector *feature_collector, const PsnrHvsStateHip *s,
                                   const double plane_score[], unsigned index)
 {
     static const char *plane_features[PSNR_HVS_NUM_PLANES] = {"psnr_hvs_y", "psnr_hvs_cb",
                                                               "psnr_hvs_cr"};
     int err = 0;
-    for (int p = 0; p < PSNR_HVS_NUM_PLANES; p++) {
+    for (unsigned p = 0; p < psnr_hvs_plane_count(s); p++) {
         err |= vmaf_feature_collector_append(feature_collector, plane_features[p],
                                              vmaf_psnr_hvs_score_db(plane_score[p]), index);
     }
-    const double combined = vmaf_psnr_hvs_combined_score(plane_score, PSNR_HVS_NUM_PLANES);
+    const double combined = vmaf_psnr_hvs_combined_score(plane_score, psnr_hvs_plane_count(s));
     err |= vmaf_feature_collector_append(feature_collector, "psnr_hvs",
                                          vmaf_psnr_hvs_score_db(combined), index);
     return err;
@@ -529,9 +561,9 @@ static int collect_fex_hip(VmafFeatureExtractor *fex, unsigned index,
         return wait_err;
 
 #ifdef HAVE_HIPCC
-    double plane_score[PSNR_HVS_NUM_PLANES];
+    double plane_score[PSNR_HVS_NUM_PLANES] = {0.0, 0.0, 0.0};
     psnr_hvs_plane_scores(s, plane_score);
-    return psnr_hvs_append_scores(feature_collector, plane_score, index);
+    return psnr_hvs_append_scores(feature_collector, s, plane_score, index);
 #else
     (void)feature_collector;
     (void)index;

core/src/feature/sycl/AGENTS.md

@@ -357,7 +357,9 @@ HIP / Metal motion twins listed in Twin-update table above — same PR.
   Kernel stays scratch-free (ADR-1395): `private_mem_size` and
   `spill_memory_size` 0 on the A380 at SIMD16 and forced SIMD32; terms go
   straight to USM, no private `float[64]`; `test_sycl_kernel_scratch` fails
-  if the kernel gains any. Guards:
+  if the kernel gains any. `n_active_planes` = 1 for `enable_chroma=false`
+  or 4:0:0 (CPU `psnr_hvs.c::init` rule, set in `configure_hvs_geometry()`),
+  else 3; 4:0:0 must not reach the shared chroma planes. Guards:
   `test_sycl_psnr_hvs_parity{,_simd32,_large}` (device, `==` on all four
   outputs, 3840x2160 included), `test_sycl_fp_arith_contract`
   (`sqrt_prod_rn` on the device vs the host's fp64 expression),

core/src/feature/sycl/integer_psnr_hvs_sycl.cpp

@@ -659,19 +659,13 @@ static const VmafOption options_psnr_hvs_sycl[] = {
 namespace
 {
 
-static int validate_hvs_input(enum VmafPixelFormat format, unsigned bpc, unsigned width,
-                              unsigned height)
+static int validate_hvs_input(unsigned bpc, unsigned width, unsigned height)
 {
     if (bpc > 12) {
         vmaf_log(VMAF_LOG_LEVEL_ERROR, "psnr_hvs_sycl: invalid bitdepth (%u); bpc must be ≤ 12\n",
                  bpc);
         return -EINVAL;
     }
-    if (format == VMAF_PIX_FMT_YUV400P) {
-        vmaf_log(VMAF_LOG_LEVEL_ERROR,
-                 "psnr_hvs_sycl: YUV400P unsupported (psnr_hvs needs all 3 planes)\n");
-        return -EINVAL;
-    }
     if (width < (unsigned)PSNR_HVS_BLOCK || height < (unsigned)PSNR_HVS_BLOCK) {
         vmaf_log(VMAF_LOG_LEVEL_ERROR, "psnr_hvs_sycl: input %ux%u smaller than 8×8 block\n", width,
                  height);
@@ -690,7 +684,15 @@ static int configure_hvs_geometry(PsnrHvsStateSycl *s, enum VmafPixelFormat form
 {
     s->width[0] = width;
     s->height[0] = height;
+    /* 4:0:0 has no chroma planes: luma only whatever enable_chroma says, as in
+     * the CPU extractor (third_party/xiph/psnr_hvs.c::init). */
+    s->n_active_planes =
+        (s->enable_chroma && format != VMAF_PIX_FMT_YUV400P) ? (unsigned)PSNR_HVS_NUM_PLANES : 1U;
     switch (format) {
+    case VMAF_PIX_FMT_YUV400P:
+        s->width[1] = s->width[2] = 0U;
+        s->height[1] = s->height[2] = 0U;
+        break;
     case VMAF_PIX_FMT_YUV420P:
         s->width[1] = s->width[2] = (width + 1u) >> 1;
         s->height[1] = s->height[2] = (height + 1u) >> 1;
@@ -717,7 +719,6 @@ namespace
 
 static int configure_hvs_blocks(PsnrHvsStateSycl *s)
 {
-    s->n_active_planes = s->enable_chroma ? (unsigned)PSNR_HVS_NUM_PLANES : 1U;
     s->total_blocks = 0U;
     for (int plane = 0; std::cmp_less(plane, s->n_active_planes); plane++) {
         if (s->width[plane] < (unsigned)PSNR_HVS_BLOCK ||
@@ -786,7 +787,7 @@ static int init_fex_sycl(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt
 {
     auto *s = static_cast<PsnrHvsStateSycl *>(fex->priv);
 
-    const int input_err = validate_hvs_input(pix_fmt, bpc, w, h);
+    const int input_err = validate_hvs_input(bpc, w, h);
     if (input_err) {
         return input_err;
     }