sync 6fdf6301e2bb

Browse files

Files changed (7) hide show

README.md +4 -2
build/webgpu/bench.json +56 -3
build/webgpu/grid-sample.wgsl.jinja +13 -49
build/webgpu/grid-sample3d.wgsl.jinja +51 -54
build/webgpu/manifest.json +12 -23
build/webgpu/metadata.json +8 -8
build/webgpu/test.json +309 -2

README.md CHANGED Viewed

@@ -49,6 +49,8 @@ Default values (overridable per request):
 One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
 - `ncdhw_rank5_channel_vector` — Shares volumetric coordinates, padding and interpolation across a vector of channels; uses two lanes for two channels and four lanes otherwise, with masked tails and device-capped workgroups.
 - `ncdhw_rank5` — Portable scalar volumetric sampling for all interpolation and padding modes.
@@ -57,14 +59,14 @@ One implementation is selected per call from the device capabilities, the reques
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`grid-sample.wgsl.jinja`](build/webgpu/grid-sample.wgsl.jinja)
 - [`grid-sample3d.wgsl.jinja`](build/webgpu/grid-sample3d.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

 One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `nchw_rank4_channel_vector` — Resolves the sampling coordinate once per pixel and reuses it across four channels, so one thread gathers from four channel planes instead of reading one coalesced. This amortizes the coordinate math across channels, trading a coalesced read for a scattered gather; adapters where scattered gathers are slow use the coalesced variant instead.
+- `nchw_rank4` — Portable scalar spatial sampling: one thread per output element, so reads stay coalesced within one channel plane. The route for every adapter whose gather throughput the channel vector would spend.
 - `ncdhw_rank5_channel_vector` — Shares volumetric coordinates, padding and interpolation across a vector of channels; uses two lanes for two channels and four lanes otherwise, with masked tails and device-capped workgroups.
 - `ncdhw_rank5` — Portable scalar volumetric sampling for all interpolation and padding modes.
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
 - [`grid-sample.wgsl.jinja`](build/webgpu/grid-sample.wgsl.jinja)
 - [`grid-sample3d.wgsl.jinja`](build/webgpu/grid-sample3d.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

build/webgpu/bench.json CHANGED Viewed

@@ -10,7 +10,7 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 256, 256] } }
     },
     {
-      "name": "nchw_1x32x256x256_linear_zeros_f32_healthy",
       "preset": "smoke",
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
       "inputs": {
@@ -63,7 +63,7 @@
       }
     },
     {
-      "name": "ncdhw_1x4x16x32x32_linear_zeros_f32_volumetric_unbenched_variant",
       "preset": "smoke",
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
       "inputs": {
@@ -412,6 +412,59 @@
       "bench": {
         "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
       }
     }
-  ]
 }

       "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 256, 256] } }
     },
     {
+      "name": "nchw_1x32x256x256_linear_zeros_f32_control",
       "preset": "smoke",
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
       "inputs": {
       }
     },
     {
+      "name": "ncdhw_1x4x16x32x32_linear_zeros_f32_volumetric",
       "preset": "smoke",
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
       "inputs": {
       "bench": {
         "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
       }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_float32_c1",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16304 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16404 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_float16_c1",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 1, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16305 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16405 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_float32_c1",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16316 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16416 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_float16_c1",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 1, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16317 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16417 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
     }
+  ],
+  "tunableSpace": { "RESOLVE_REFLECTION_TAPS": [0, 1] }
 }

build/webgpu/grid-sample.wgsl.jinja CHANGED Viewed

@@ -1,36 +1,8 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) {
-    return;
-  }
-{%- endif %}
-{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
@@ -72,8 +44,8 @@ fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   }
   return x0;
 }
 {% endif %}
 {% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
@@ -92,7 +64,7 @@ fn sample_floor(v: f32) -> i32 {
   }
   return i32(floor(v));
 }
-{%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
     {% if paddingMode == "reflection" %}
@@ -110,10 +82,9 @@ fn sample_round(v: f32) -> i32 {
   }
   return i32(round(v));
 }
-{%- endif %}
 {% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
@@ -128,9 +99,9 @@ fn cubic_one(x0: f32, a: f32) -> f32 {
   if (x0 < 2.0) { return a * x0 * x0 * x0 - 5.0 * a * x0 * x0 + 8.0 * a * x0 - 4.0 * a; }
   return 0.0;
 }
-{% endif -%}
 {% if channelWidthSpec is not defined %}
 fn pixel(base: u32, h: i32, w: i32) -> f32 {
   {% if paddingMode == "zeros" %}
   if (h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) { return 0.0; }
@@ -152,12 +123,11 @@ fn pixel(base: u32, h: i32, w: i32) -> f32 {
   {% endif %}
   return f32(x[base + hh * params.inW + ww]);
 }
 {% endif %}
 {% if channelWidthSpec is defined %}
 {% set channelVec = "vec" ~ channelWidthSpec ~ "<f32>" %}
 {% set components = ["x", "y", "z", "w"] %}
 // A block of NCHW channels shares one grid coordinate. Tail lanes alias the
 // final channel safely because main suppresses their stores.
 {% if paddingMode != "reflection" or modeSpec == "nearest" %}
@@ -188,12 +158,9 @@ fn pixel_channels(n: u32, c0: u32, h: i32, w: i32) -> {{ channelVec }} {
     f32(x[(nb + {% if channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
-{% endif %}
-{% if paddingMode == "reflection" %}
 // Reflection gives every tap in one row or column the same resolved coordinate.
 // Resolve each row and column once in main and reuse it for all channel loads.
-{% if modeSpec != "nearest" %}
 fn pixel_channels_resolved(n: u32, c0: u32, h: u32, w: u32) -> {{ channelVec }} {
   let plane = params.inH * params.inW;
   let offset = h * params.inW + w;
@@ -203,15 +170,12 @@ fn pixel_channels_resolved(n: u32, c0: u32, h: u32, w: u32) -> {{ channelVec }}
     f32(x[(nb + {% if channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
 {% endif %}
-{% else %}
 {% endif %}
-{% endif -%}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("i", "") }}
 {% if channelWidthSpec is defined %}
   let out_plane = params.outH * params.outW;
   let channel_blocks = (params.C + {{ channelWidthSpec - 1 }}u) / {{ channelWidthSpec }}u;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
   }
   return x0;
 }
 {% endif %}
 {% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
   }
   return i32(floor(v));
 }
+{% else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
     {% if paddingMode == "reflection" %}
   }
   return i32(round(v));
 }
+{% endif %}
 {% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
   if (x0 < 2.0) { return a * x0 * x0 * x0 - 5.0 * a * x0 * x0 + 8.0 * a * x0 - 4.0 * a; }
   return 0.0;
 }
+{% endif %}
 {% if channelWidthSpec is not defined %}
 fn pixel(base: u32, h: i32, w: i32) -> f32 {
   {% if paddingMode == "zeros" %}
   if (h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) { return 0.0; }
   {% endif %}
   return f32(x[base + hh * params.inW + ww]);
 }
 {% endif %}
 {% if channelWidthSpec is defined %}
 {% set channelVec = "vec" ~ channelWidthSpec ~ "<f32>" %}
 {% set components = ["x", "y", "z", "w"] %}
 // A block of NCHW channels shares one grid coordinate. Tail lanes alias the
 // final channel safely because main suppresses their stores.
 {% if paddingMode != "reflection" or modeSpec == "nearest" %}
     f32(x[(nb + {% if channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
+{% else %}
 // Reflection gives every tap in one row or column the same resolved coordinate.
 // Resolve each row and column once in main and reuse it for all channel loads.
 fn pixel_channels_resolved(n: u32, c0: u32, h: u32, w: u32) -> {{ channelVec }} {
   let plane = params.inH * params.inW;
   let offset = h * params.inW + w;
     f32(x[(nb + {% if channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
 {% endif %}
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "i", "") }}
 {% if channelWidthSpec is defined %}
   let out_plane = params.outH * params.outW;
   let channel_blocks = (params.C + {{ channelWidthSpec - 1 }}u) / {{ channelWidthSpec }}u;

build/webgpu/grid-sample3d.wgsl.jinja CHANGED Viewed

@@ -1,36 +1,9 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) {
-    return;
-  }
-{%- endif %}
-{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
@@ -72,8 +45,8 @@ fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   }
   return x0;
 }
 {% endif %}
 {% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
@@ -92,7 +65,7 @@ fn sample_floor(v: f32) -> i32 {
   }
   return i32(floor(v));
 }
-{%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
     {% if paddingMode == "reflection" %}
@@ -110,10 +83,9 @@ fn sample_round(v: f32) -> i32 {
   }
   return i32(round(v));
 }
-{%- endif %}
 {% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
@@ -128,12 +100,34 @@ fn cubic_one(x0: f32, a: f32) -> f32 {
   if (x0 < 2.0) { return a * x0 * x0 * x0 - 5.0 * a * x0 * x0 + 8.0 * a * x0 - 4.0 * a; }
   return 0.0;
 }
-{% endif -%}
 {% set volumeVector = channelWidthSpec is defined %}
 {% set volumeT = "vec" ~ channelWidthSpec ~ "<f32>" if volumeVector else "f32" %}
 {% set voxelArgs = "img_base, c, " if volumeVector and channelTail else "img_base, " %}
-fn voxel(base: u32,{% if volumeVector and channelTail %} c0: u32,{% endif %} d: i32, h: i32, w: i32) -> {{ volumeT }} {
   {% if paddingMode == "zeros" %}
   if (d < 0 || d >= i32(params.inD) || h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) {
     return {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
@@ -159,6 +153,7 @@ fn voxel(base: u32,{% if volumeVector and channelTail %} c0: u32,{% endif %} d:
   let hh = u32(clamp(rh, 0, i32(params.inH) - 1));
   let ww = u32(clamp(rw, 0, i32(params.inW) - 1));
   {% endif %}
 {% if volumeVector %}
   // All channels share the resolved voxel coordinates. Incomplete vectors
   // alias the final channel on loads; main masks their output stores.
@@ -173,7 +168,7 @@ fn voxel(base: u32,{% if volumeVector and channelTail %} c0: u32,{% endif %} d:
 @compute @workgroup_size({{ volumeWorkgroup if volumeVector else tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d(guardInline=true) }}
   let ow = i % params.outW;
   var t = i / params.outW;
@@ -200,22 +195,23 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let result = voxel({{ voxelArgs }}sample_round(sz), sample_round(sy), sample_round(sx));
   {% elif modeSpec == "cubic" %}
   // Tricubic applies separable four-tap Keys interpolation along each spatial
-  // axis; voxel() already resolves every padding mode, so
-  // the 64-tap window needs no boundary handling of its own.
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let z0 = sample_floor(sz) - 1;
   let cx = cubic_coeffs(sx - f32(x0 + 1));
   let cy = cubic_coeffs(sy - f32(y0 + 1));
   let cz = cubic_coeffs(sz - f32(z0 + 1));
   var result = {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
   for (var kz = 0i; kz < 4i; kz = kz + 1i) {
     var plane = {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
     for (var ky = 0i; ky < 4i; ky = ky + 1i) {
-      let row = cx.x * voxel({{ voxelArgs }}z0 + kz, y0 + ky, x0)
-        + cx.y * voxel({{ voxelArgs }}z0 + kz, y0 + ky, x0 + 1)
-        + cx.z * voxel({{ voxelArgs }}z0 + kz, y0 + ky, x0 + 2)
-        + cx.w * voxel({{ voxelArgs }}z0 + kz, y0 + ky, x0 + 3);
       plane = plane + cy[u32(ky)] * row;
     }
     result = result + cz[u32(kz)] * plane;
@@ -227,14 +223,15 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let wx = sx - f32(x0);
   let wy = sy - f32(y0);
   let wz = sz - f32(z0);
-  let v000 = voxel({{ voxelArgs }}z0, y0, x0);
-  let v001 = voxel({{ voxelArgs }}z0, y0, x0 + 1);
-  let v010 = voxel({{ voxelArgs }}z0, y0 + 1, x0);
-  let v011 = voxel({{ voxelArgs }}z0, y0 + 1, x0 + 1);
-  let v100 = voxel({{ voxelArgs }}z0 + 1, y0, x0);
-  let v101 = voxel({{ voxelArgs }}z0 + 1, y0, x0 + 1);
-  let v110 = voxel({{ voxelArgs }}z0 + 1, y0 + 1, x0);
-  let v111 = voxel({{ voxelArgs }}z0 + 1, y0 + 1, x0 + 1);
   let front = (1.0 - wy) * ((1.0 - wx) * v000 + wx * v001) + wy * ((1.0 - wx) * v010 + wx * v011);
   let back = (1.0 - wy) * ((1.0 - wx) * v100 + wx * v101) + wy * ((1.0 - wx) * v110 + wx * v111);
   let result = (1.0 - wz) * front + wz * back;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) { return; }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
   }
   return x0;
 }
 {% endif %}
 {% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
   }
   return i32(floor(v));
 }
+{% else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
     {% if paddingMode == "reflection" %}
   }
   return i32(round(v));
 }
+{% endif %}
 {% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
   if (x0 < 2.0) { return a * x0 * x0 * x0 - 5.0 * a * x0 * x0 + 8.0 * a * x0 - 4.0 * a; }
   return 0.0;
 }
+{% endif %}
 {% set volumeVector = channelWidthSpec is defined %}
 {% set volumeT = "vec" ~ channelWidthSpec ~ "<f32>" if volumeVector else "f32" %}
 {% set voxelArgs = "img_base, c, " if volumeVector and channelTail else "img_base, " %}
+{% if resolveReflectionTaps %}
+fn reflected_index(index: i32, length: u32) -> u32 {
+{% if alignCorners %}
+  let reflected = i32(reflect_coord(f32(index), 0.0, f32(length) - 1.0));
+{% else %}
+  let reflected = i32(reflect_coord(f32(index), -0.5, f32(length) - 0.5));
+{% endif %}
+  return u32(clamp(reflected, 0, i32(length) - 1));
+}
+{% endif %}
+{% macro reflected_axis_taps() %}
+{% if resolveReflectionTaps %}
+{% for axis, length in [["x", "inW"], ["y", "inH"], ["z", "inD"]] %}
+  let i{{ axis }} = vec{{ 4 if modeSpec == "cubic" else 2 }}<u32>(
+{% for tap in range(4 if modeSpec == "cubic" else 2) %}
+    reflected_index({{ axis }}0{% if tap %} + {{ tap }}{% endif %}, params.{{ length }}){% if loop.last %});{% else %},{% endif %}
+{% endfor %}
+{% endfor %}
+{% endif %}
+{% endmacro %}
+fn voxel(base: u32,{% if volumeVector and channelTail %} c0: u32,{% endif %} {% if resolveReflectionTaps %}dd: u32, hh: u32, ww: u32{% else %}d: i32, h: i32, w: i32{% endif %}) -> {{ volumeT }} {
+{% if not resolveReflectionTaps %}
   {% if paddingMode == "zeros" %}
   if (d < 0 || d >= i32(params.inD) || h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) {
     return {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
   let hh = u32(clamp(rh, 0, i32(params.inH) - 1));
   let ww = u32(clamp(rw, 0, i32(params.inW) - 1));
   {% endif %}
+{% endif %}
 {% if volumeVector %}
   // All channels share the resolved voxel coordinates. Incomplete vectors
   // alias the final channel on loads; main masks their output stores.
 @compute @workgroup_size({{ volumeWorkgroup if volumeVector else tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(volumeWorkgroup if volumeVector else tunables.WORKGROUP_SIZE, guardInline=true) }}
   let ow = i % params.outW;
   var t = i / params.outW;
   let result = voxel({{ voxelArgs }}sample_round(sz), sample_round(sy), sample_round(sx));
   {% elif modeSpec == "cubic" %}
   // Tricubic applies separable four-tap Keys interpolation along each spatial
+  // axis. Reflection taps are resolved once per axis when enabled; other
+  // padding modes are handled by voxel().
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let z0 = sample_floor(sz) - 1;
   let cx = cubic_coeffs(sx - f32(x0 + 1));
   let cy = cubic_coeffs(sy - f32(y0 + 1));
   let cz = cubic_coeffs(sz - f32(z0 + 1));
+{{ reflected_axis_taps() }}
   var result = {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
   for (var kz = 0i; kz < 4i; kz = kz + 1i) {
     var plane = {% if volumeVector %}{{ volumeT }}(0.0){% else %}0.0{% endif %};
     for (var ky = 0i; ky < 4i; ky = ky + 1i) {
+      let row = cx.x * voxel({{ voxelArgs }}{{ "iz[u32(kz)], iy[u32(ky)], ix.x" if resolveReflectionTaps else "z0 + kz, y0 + ky, x0" }})
+        + cx.y * voxel({{ voxelArgs }}{{ "iz[u32(kz)], iy[u32(ky)], ix.y" if resolveReflectionTaps else "z0 + kz, y0 + ky, x0 + 1" }})
+        + cx.z * voxel({{ voxelArgs }}{{ "iz[u32(kz)], iy[u32(ky)], ix.z" if resolveReflectionTaps else "z0 + kz, y0 + ky, x0 + 2" }})
+        + cx.w * voxel({{ voxelArgs }}{{ "iz[u32(kz)], iy[u32(ky)], ix.w" if resolveReflectionTaps else "z0 + kz, y0 + ky, x0 + 3" }});
       plane = plane + cy[u32(ky)] * row;
     }
     result = result + cz[u32(kz)] * plane;
   let wx = sx - f32(x0);
   let wy = sy - f32(y0);
   let wz = sz - f32(z0);
+{{ reflected_axis_taps() }}
+  let v000 = voxel({{ voxelArgs }}{{ "iz.x, iy.x, ix.x" if resolveReflectionTaps else "z0, y0, x0" }});
+  let v001 = voxel({{ voxelArgs }}{{ "iz.x, iy.x, ix.y" if resolveReflectionTaps else "z0, y0, x0 + 1" }});
+  let v010 = voxel({{ voxelArgs }}{{ "iz.x, iy.y, ix.x" if resolveReflectionTaps else "z0, y0 + 1, x0" }});
+  let v011 = voxel({{ voxelArgs }}{{ "iz.x, iy.y, ix.y" if resolveReflectionTaps else "z0, y0 + 1, x0 + 1" }});
+  let v100 = voxel({{ voxelArgs }}{{ "iz.y, iy.x, ix.x" if resolveReflectionTaps else "z0 + 1, y0, x0" }});
+  let v101 = voxel({{ voxelArgs }}{{ "iz.y, iy.x, ix.y" if resolveReflectionTaps else "z0 + 1, y0, x0 + 1" }});
+  let v110 = voxel({{ voxelArgs }}{{ "iz.y, iy.y, ix.x" if resolveReflectionTaps else "z0 + 1, y0 + 1, x0" }});
+  let v111 = voxel({{ voxelArgs }}{{ "iz.y, iy.y, ix.y" if resolveReflectionTaps else "z0 + 1, y0 + 1, x0 + 1" }});
   let front = (1.0 - wy) * ((1.0 - wx) * v000 + wx * v001) + wy * ((1.0 - wx) * v010 + wx * v011);
   let back = (1.0 - wy) * ((1.0 - wx) * v100 + wx * v101) + wy * ((1.0 - wx) * v110 + wx * v111);
   let result = (1.0 - wz) * front + wz * back;

build/webgpu/manifest.json CHANGED Viewed

@@ -22,8 +22,9 @@
     "align_corners": { "values": [0, 1] }
   },
   "typeConstraints": { "T": ["float32", "float16"] },
-  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
     "reportedNonWave32Adapter": "not wave32Adapter and (has(device.adapterInfo, \"subgroupMinSize\") or has(device.adapterInfo, \"subgroupMaxSize\"))",
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
@@ -31,11 +32,13 @@
     "channelWidth": 4,
     "scalar": "dtypes.T",
     "volumeWidth": "min(channelWidth, pow(2, log2ceil(max(1, dim(shapes.x, 1)))))",
-    "volumeWorkgroup": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)"
   },
   "bindings": {
     "params": {
-      "buffer": "uniform",
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
         { "name": "C", "type": "u32", "value": "dim(shapes.x, 1)" },
@@ -50,19 +53,14 @@
     {
       "id": "nchw_rank4_channel_vector",
       "priority": 5,
-      "when": ["rank4Ok", "not reportedNonWave32Adapter", "dim(shapes.x, 1) >= 2"],
       "passes": [
         {
           "id": "main",
           "name": "GridSample.ChannelX4",
           "shader": "grid-sample.wgsl.jinja",
-          "derive": {
-            "modeSpec": "attrs.mode",
-            "paddingMode": "attrs.padding_mode",
-            "alignCorners": "attrs.align_corners != 0",
-            "channelWidthSpec": "channelWidth",
-            "channelTail": "dim(shapes.x, 1) % channelWidth != 0"
-          },
           "bindings": ["x", "grid", "y", "params"],
           "dispatch": {
             "x": "min(ceilDiv((dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), channelWidth) * dim(shapes.y, 2) * dim(shapes.y, 3)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -80,11 +78,6 @@
           "id": "main",
           "name": "GridSample",
           "shader": "grid-sample.wgsl.jinja",
-          "derive": {
-            "modeSpec": "attrs.mode",
-            "paddingMode": "attrs.padding_mode",
-            "alignCorners": "attrs.align_corners != 0"
-          },
           "bindings": ["x", "grid", "y", "params"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -104,11 +97,9 @@
           "name": "GridSample.VolumetricChannels",
           "shader": "grid-sample3d.wgsl.jinja",
           "derive": {
-            "modeSpec": "attrs.mode",
-            "paddingMode": "attrs.padding_mode",
-            "alignCorners": "attrs.align_corners != 0",
             "channelWidthSpec": "volumeWidth",
-            "channelTail": "dim(shapes.x, 1) % volumeWidth != 0"
           },
           "bindings": [
             "x",
@@ -150,9 +141,7 @@
           "name": "GridSample.Volumetric",
           "shader": "grid-sample3d.wgsl.jinja",
           "derive": {
-            "modeSpec": "attrs.mode",
-            "paddingMode": "attrs.padding_mode",
-            "alignCorners": "attrs.align_corners != 0"
           },
           "bindings": [
             "x",

     "align_corners": { "values": [0, 1] }
   },
   "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 }, "RESOLVE_REFLECTION_TAPS": { "default": 1 } },
   "derive": {
+    "appleAdapter": "has(device.adapterInfo, \"vendor\") and device.adapterInfo.vendor == \"apple\"",
     "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
     "reportedNonWave32Adapter": "not wave32Adapter and (has(device.adapterInfo, \"subgroupMinSize\") or has(device.adapterInfo, \"subgroupMaxSize\"))",
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "channelWidth": 4,
     "scalar": "dtypes.T",
     "volumeWidth": "min(channelWidth, pow(2, log2ceil(max(1, dim(shapes.x, 1)))))",
+    "volumeWorkgroup": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "alignCorners": "attrs.align_corners != 0",
+    "modeSpec": "attrs.mode",
+    "paddingMode": "attrs.padding_mode"
   },
   "bindings": {
     "params": {
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
         { "name": "C", "type": "u32", "value": "dim(shapes.x, 1)" },
     {
       "id": "nchw_rank4_channel_vector",
       "priority": 5,
+      "when": ["rank4Ok", "dim(shapes.x, 1) >= 2"],
+      "demoteWhen": ["not appleAdapter"],
       "passes": [
         {
           "id": "main",
           "name": "GridSample.ChannelX4",
           "shader": "grid-sample.wgsl.jinja",
+          "derive": { "channelWidthSpec": "channelWidth", "channelTail": "dim(shapes.x, 1) % channelWidth != 0" },
           "bindings": ["x", "grid", "y", "params"],
           "dispatch": {
             "x": "min(ceilDiv((dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), channelWidth) * dim(shapes.y, 2) * dim(shapes.y, 3)), (tunables.WORKGROUP_SIZE)), 65535)",
           "id": "main",
           "name": "GridSample",
           "shader": "grid-sample.wgsl.jinja",
           "bindings": ["x", "grid", "y", "params"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "GridSample.VolumetricChannels",
           "shader": "grid-sample3d.wgsl.jinja",
           "derive": {
             "channelWidthSpec": "volumeWidth",
+            "channelTail": "dim(shapes.x, 1) % volumeWidth != 0",
+            "resolveReflectionTaps": "attrs.padding_mode == \"reflection\" and attrs.mode != \"nearest\" and tunables.RESOLVE_REFLECTION_TAPS != 0"
           },
           "bindings": [
             "x",
           "name": "GridSample.Volumetric",
           "shader": "grid-sample3d.wgsl.jinja",
           "derive": {
+            "resolveReflectionTaps": "attrs.padding_mode == \"reflection\" and attrs.mode != \"nearest\" and tunables.RESOLVE_REFLECTION_TAPS != 0"
           },
           "bindings": [
             "x",

build/webgpu/metadata.json CHANGED Viewed

@@ -1,22 +1,22 @@
 {
   "name": "ai.onnx.GridSample",
-  "id": "_ai_onnx_gridsample_webgpu_f340775",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "GKt3XLyZG7PU01DKekKBLzn6wJaj5/4VX+430PWZLas=",
-      "grid-sample.wgsl.jinja": "azjm8+LAvCxuA3CAJtB0ElRwBbDCOy05cOL8rJzeeK0=",
-      "grid-sample3d.wgsl.jinja": "zLErkeaeWt4gEn6iAkcEYGDEhauXCj+czW4rO0kbv2k=",
-      "manifest.json": "PcNyEEhLUbrU5SY4T6mBxI/MFk1fnH4BQWcpSnleDHU=",
-      "test.json": "mLgrc4Eep+6yiX6JmU0Hpt/dMAjXnL6kxOFsL+/hO38="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "nchw_rank4_channel_vector": ["grid-sample.wgsl.jinja"],
       "nchw_rank4": ["grid-sample.wgsl.jinja"],

 {
   "name": "ai.onnx.GridSample",
+  "id": "_ai_onnx_gridsample_webgpu_11c7e55",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "zRyVJ7WL0vHjv1U48z8EM9u2l5FufDQYPq5VfWYxVyw=",
+      "grid-sample.wgsl.jinja": "+smIQHwp1GEltoAVpbIKI0oJbsjAXGVGPy2ZhVHvF8Y=",
+      "grid-sample3d.wgsl.jinja": "e638b7ZNv0/EmEXA5+M4PlzQ81vHXpRlj3t5iVieohY=",
+      "manifest.json": "UyKUdbD0viJTcRHht1BWQpHNviVreaI5JZjLthQDKM0=",
+      "test.json": "5SMIjxyMjqQYJ+GGIjZ6md6MjqRGQwkKIbZn5NAIGRw="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "nchw_rank4_channel_vector": ["grid-sample.wgsl.jinja"],
       "nchw_rank4": ["grid-sample.wgsl.jinja"],

build/webgpu/test.json CHANGED Viewed

@@ -5,7 +5,8 @@
     "onnx_backend_far_coords_zeros_padding_input_grid": [-10, -10, -5, -5, -0.2, -0.2, 10, 10, 10, 10, -0.2, -0.2, 5, 5, 10, 10],
     "onnx_backend_gridsample_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15],
     "onnx_backend_gridsample_aligncorners_true_input_grid": [-1, -1, -0.5, -0.5, -0.20000000298023224, -0.20000000298023224, 0, 0, 0, 0, -0.20000000298023224, -0.20000000298023224, 0.5, 0.5, 1, 1],
-    "onnx_backend_gridsample_border_padding_input_grid": [-10, -10, -5, -5, -0.20000000298023224, -0.20000000298023224, 10, 10, 10, 10, -0.20000000298023224, -0.20000000298023224, 5, 5, 10, 10]
   },
   "cases": [
     {
@@ -227,7 +228,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/grid_sample_test_custom.cc",
         "test": "GridSampleCustomTest.test_grid_sample_20_4D_bilinear_reflection_extreme_coords",
-        "notes": "Uses smaller finite coordinates than ORT's extreme-coordinate regression while preserving the reflection boundary path."
       },
       "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
       "inputs": {
@@ -2739,6 +2740,312 @@
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     }
   ]
 }

     "onnx_backend_far_coords_zeros_padding_input_grid": [-10, -10, -5, -5, -0.2, -0.2, 10, 10, 10, 10, -0.2, -0.2, 5, 5, 10, 10],
     "onnx_backend_gridsample_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15],
     "onnx_backend_gridsample_aligncorners_true_input_grid": [-1, -1, -0.5, -0.5, -0.20000000298023224, -0.20000000298023224, 0, 0, 0, 0, -0.20000000298023224, -0.20000000298023224, 0.5, 0.5, 1, 1],
+    "onnx_backend_gridsample_border_padding_input_grid": [-10, -10, -5, -5, -0.20000000298023224, -0.20000000298023224, 10, 10, 10, 10, -0.20000000298023224, -0.20000000298023224, 5, 5, 10, 10],
+    "reflection_coordinate_f32_rounding_rank4_input_x": [1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1, 1, -1]
   },
   "cases": [
     {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/grid_sample_test_custom.cc",
         "test": "GridSampleCustomTest.test_grid_sample_20_4D_bilinear_reflection_extreme_coords",
+        "notes": "Grid coordinates of +/-5, well outside the normalized [-1,1] range but smaller than ONNX Runtime's extreme-coordinate test values, check reflection padding at finite far-out-of-bounds coordinates with bilinear sampling."
       },
       "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
       "inputs": {
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_ac0_float32",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3, 4, 7], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_ac1_float32",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3, 4, 7], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_tail_linear_reflection_float32_per_voxel",
+      "provenance": {
+        "notes": "Explicit per-voxel comparison path retains the reflection behavior on incomplete channel vectors."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 7, 9, 11, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 7, 9, 11], "tolerance": 0.0001, "relTolerance": 0.0001 } },
+      "tunables": { "RESOLVE_REFLECTION_TAPS": 0 }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_ac0_float16",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 3, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3, 4, 7], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "volume_scalar_linear_reflection_ac1_float16",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3, 4, 7], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "volume_tail_linear_reflection_float16_per_voxel",
+      "provenance": {
+        "notes": "Explicit per-voxel comparison path retains the reflection behavior on incomplete channel vectors."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 5, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 7, 9, 11, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 7, 9, 11], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "tunables": { "RESOLVE_REFLECTION_TAPS": 0 }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_ac0_float32",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3, 4, 7], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_ac1_float32",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3, 4, 7], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_tail_cubic_reflection_float32_per_voxel",
+      "provenance": {
+        "notes": "Explicit per-voxel comparison path retains the reflection behavior on incomplete channel vectors."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 7, 9, 11, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 7, 9, 11], "tolerance": 0.0001, "relTolerance": 0.0001 } },
+      "tunables": { "RESOLVE_REFLECTION_TAPS": 0 }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_ac0_float16",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 3, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3, 4, 7], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "volume_scalar_cubic_reflection_ac1_float16",
+      "provenance": {
+        "notes": "Scalar volumetric reflection with out-of-range taps; aligned cases include a dimension-one input to exercise the zero-width reflection interval."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 7, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3, 4, 7], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "volume_tail_cubic_reflection_float16_per_voxel",
+      "provenance": {
+        "notes": "Explicit per-voxel comparison path retains the reflection behavior on incomplete channel vectors."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 5, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.17, "cosStep": 0.09 }
+        },
+        "grid": {
+          "dtype": "float16",
+          "shape": [2, 7, 9, 11, 3],
+          "data": { "kind": "linspace", "start": -3.5, "end": 3.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 7, 9, 11], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "tunables": { "RESOLVE_REFLECTION_TAPS": 0 }
+    },
+    {
+      "name": "reflection_coordinate_f32_rounding_rank4",
+      "provenance": {
+        "notes": "Float32 coordinate normalization outside the image must round before interpolation. Alternating pixels amplify the f32 versus f64 coordinate difference; expected value independently checked with ONNX Runtime's CPU provider."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 256, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/reflection_coordinate_f32_rounding_rank4_input_x" }
+          }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 2],
+          "data": { "kind": "values", "values": [0.0, 3.7643392086029053] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1],
+          "data": { "kind": "values", "values": [-0.0933837890625] },
+          "tolerance": 0.0001,
+          "relTolerance": 0.0001
+        }
+      }
+    },
+    {
+      "name": "reflection_coordinate_f32_rounding_rank5",
+      "provenance": {
+        "notes": "Float32 coordinate normalization outside the image must round before interpolation. Alternating pixels amplify the f32 versus f64 coordinate difference; expected value independently checked with ONNX Runtime's CPU provider."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 256, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/reflection_coordinate_f32_rounding_rank4_input_x" }
+          }
+        },
+        "grid": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 3],
+          "data": { "kind": "values", "values": [0.0, 3.7643392086029053, 0.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [-0.0933837890625] },
+          "tolerance": 0.0001,
+          "relTolerance": 0.0001
+        }
+      }
     }
   ]
 }