sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +21 -10
build/webgpu/bench.json +286 -1
build/webgpu/grid-sample.wgsl.jinja +54 -57
build/webgpu/grid-sample3d.wgsl.jinja +79 -52
build/webgpu/manifest.json +122 -122
build/webgpu/metadata.json +16 -8
build/webgpu/test.json +787 -8

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `GridSample` spec](https://onnx.ai/onnx/operators/onnx__GridSample
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | Input tensor of shape `(N, C, D1, ..., Dr)` whose values are sampled. | required |
-| `grid` | `grid` | `T` | — | — | Flow-field of shape `(N, D1_out, ..., Dr_out, r)` with normalized sampling coordinates in `[-1, 1]`. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | same as `grid` | derived; see description | Output tensor of shape `(N, C, D1_out, ..., Dr_out)` containing the interpolated samples. | required |
 ## Attributes
@@ -35,9 +35,9 @@ Default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
-| `mode` | `"linear"` | Interpolation method: `linear` (bilinear or trilinear, depending on rank), `nearest`, or `cubic`. Cubic interpolation is supported for rank-4 (2-D spatial) inputs. |
-| `padding_mode` | `"zeros"` | How out-of-bound grid positions are handled: `zeros` pads with 0, `border` clamps to the border value, or `reflection` reflects coordinates back into the valid range. |
 | `align_corners` | `0` | When 1, extrema values `-1` and `1` map to the center of the corner pixels; when 0 (default) they map to the outer edge of corner pixels, making sampling resolution-agnostic. |
 ## Type constraints
@@ -45,9 +45,16 @@ Default values (overridable per request):
 | --- | --- |
 | `T` | `float32`, `float16` |
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, 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
@@ -56,10 +63,14 @@ Default values (overridable per request):
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `x` | `X` | `T` | — | — | Input tensor of shape `(N, C, D1, ..., Dr)` whose values are sampled. | required |
+| `grid` | — | `T` | — | — | Flow-field of shape `(N, D1_out, ..., Dr_out, r)` with normalized sampling coordinates in `[-1, 1]`. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | same as `grid` | derived | Output tensor of shape `(N, C, D1_out, ..., Dr_out)` containing the interpolated samples. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `align_corners` | `0` | When 1, extrema values `-1` and `1` map to the center of the corner pixels; when 0 (default) they map to the outer edge of corner pixels, making sampling resolution-agnostic. |
+| `mode` | `"linear"` | Interpolation method: `linear` (bilinear or trilinear, depending on rank), `nearest`, or `cubic` (bicubic for rank-4 inputs and tricubic for rank-5 inputs). |
+| `padding_mode` | `"zeros"` | How out-of-bound grid positions are handled: `zeros` pads with 0, `border` clamps to the border value, or `reflection` reflects coordinates back into the valid range. |
 ## Type constraints
 | --- | --- |
 | `T` | `float32`, `float16` |
+## Implementation variants
+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.
 ## Files
+- [`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
 ## 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.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.GridSample",
   "cases": [
     {
       "name": "1x3x256x256_to_256_linear",
@@ -127,6 +126,292 @@
       "bench": {
         "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
       }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "1x3x256x256_to_256_linear",
       "bench": {
         "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
       }
+    },
+    {
+      "name": "volume_channels_linear_zeros_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16300 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16400 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_zeros_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16301 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16401 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_border_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "border", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16302 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16402 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_border_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "border", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16303 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16403 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 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, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 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, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_zeros_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16306 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16406 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_zeros_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16307 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16407 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_border_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "border", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16308 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16408 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_border_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "border", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16309 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16409 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_reflection_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16310 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16410 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_nearest_reflection_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "nearest", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16311 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16411 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_zeros_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16312 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16412 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_zeros_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16313 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 0.9, "seed": 16413 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_border_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16314 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16414 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_border_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16315 },
+        "grid": { "dtype": "float16", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16415 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_reflection_float32_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8, 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, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_cubic_reflection_float16_c8",
+      "preset": "stress",
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 8, 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, 8, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 2" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float32_c2",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 2, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16318 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16418 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float32_c3",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16319 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16419 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float32_c5",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 5, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16320 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16420 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 5, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
+    },
+    {
+      "name": "volume_channels_linear_reflection_float32_c16",
+      "preset": "stress",
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 16, 12, 16, 16], "dist": "normal", "scale": 0.3, "seed": 16321 },
+        "grid": { "dtype": "float32", "shape": [1, 24, 32, 32, 3], "dist": "normal", "scale": 1.5, "seed": 16421 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 16, 24, 32, 32] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.grid) + numel(shapes.y)) * 4" }]
+      }
     }
   ]
 }

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

@@ -1,43 +1,40 @@
 {% 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
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
 {% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; 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
-  // maxComputeWorkgroupsPerDimension limit (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
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "dispatch" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit.
 {% endif %}
 {% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {%- endif %}
 {% endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
-  {% if source.alignCorners %}
   return (n + 1.0) * 0.5 * f32(length - 1u);
   {% else %}
   return ((n + 1.0) * f32(length) - 1.0) * 0.5;
@@ -46,7 +43,7 @@ fn denormalize(n: f32, length: u32) -> f32 {
 fn sanitize_coord(v: f32) -> f32 {
   if (!(v >= -2147483000.0 && v <= 2147483000.0)) {
-    {% if source.alignCorners %}
     return 0.0;
     {% else %}
     return -0.5;
@@ -54,7 +51,7 @@ fn sanitize_coord(v: f32) -> f32 {
   }
   return v;
 }
-{% if source.paddingMode == "reflection" %}
 fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   let range = hi - lo;
@@ -77,17 +74,17 @@ fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
 }
 {% endif %}
-{% if source.mode != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
@@ -98,14 +95,14 @@ fn sample_floor(v: f32) -> i32 {
 {%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
@@ -116,7 +113,7 @@ fn sample_round(v: f32) -> i32 {
 {%- endif %}
-{% if source.mode == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
@@ -133,17 +130,17 @@ fn cubic_one(x0: f32, a: f32) -> f32 {
 }
 {% endif -%}
-{% if source.channelWidth is not defined %}
 fn pixel(base: u32, h: i32, w: i32) -> f32 {
-  {% if source.paddingMode == "zeros" %}
   if (h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) { return 0.0; }
   let hh = u32(h);
   let ww = u32(w);
-  {% elif source.paddingMode == "border" %}
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
-  {% if source.alignCorners %}
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
   {% else %}
@@ -158,22 +155,22 @@ fn pixel(base: u32, h: i32, w: i32) -> f32 {
 {% endif %}
-{% if source.channelWidth is defined %}
-{% set channelVec = "vec" ~ source.channelWidth ~ "<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 source.paddingMode != "reflection" or source.mode == "nearest" %}
 fn pixel_channels(n: u32, c0: u32, h: i32, w: i32) -> {{ channelVec }} {
-  {% if source.paddingMode == "zeros" %}
   if (h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) { return {{ channelVec }}(0.0); }
   let hh = u32(h);
   let ww = u32(w);
-  {% elif source.paddingMode == "border" %}
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
-  {% if source.alignCorners %}
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
   {% else %}
@@ -187,23 +184,23 @@ fn pixel_channels(n: u32, c0: u32, h: i32, w: i32) -> {{ channelVec }} {
   let offset = hh * params.inW + ww;
   let nb = n * params.C;
   return {{ channelVec }}(
-{% for lane in range(source.channelWidth) %}
-    f32(x[(nb + {% if source.channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
 {% endif %}
-{% if source.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 source.mode != "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;
   let nb = n * params.C;
   return {{ channelVec }}(
-{% for lane in range(source.channelWidth) %}
-    f32(x[(nb + {% if source.channelTail %}min(c0 + {{ lane }}u, params.C - 1u){% else %}c0 + {{ lane }}u{% endif %}) * plane + offset]){% if not loop.last %},{% else %});{% endif %}
 {% endfor %}
 }
@@ -213,34 +210,34 @@ fn pixel_channels_resolved(n: u32, c0: u32, h: u32, w: u32) -> {{ channelVec }}
 {% endif %}
 {% endif -%}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d("i", "") }}
-{% if source.channelWidth is defined %}
   let out_plane = params.outH * params.outW;
-  let channel_blocks = (params.C + {{ source.channelWidth - 1 }}u) / {{ source.channelWidth }}u;
   let block_count = (params.count / params.C) * channel_blocks;
   if (i >= block_count) { return; }
   let spatial = i % out_plane;
   let cb = (i / out_plane) % channel_blocks;
   let n = i / (out_plane * channel_blocks);
-  let c = cb * {{ source.channelWidth }}u;
   let grid_base = (n * out_plane + spatial) * 2u;
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
   let sy = sanitize_coord(denormalize(f32(grid[grid_base + 1u]), params.inH));
-  {% if source.mode == "nearest" %}
   let result = pixel_channels(n, c, sample_round(sy), sample_round(sx));
-  {% elif source.mode == "cubic" %}
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let dx = sx - f32(x0 + 1);
   let dy = sy - f32(y0 + 1);
   let cx = cubic_coeffs(dx);
   let cy = cubic_coeffs(dy);
-{% if source.paddingMode == "reflection" %}
   var reflected_y: array<u32, 4>;
   var reflected_x: array<u32, 4>;
   for (var k = 0u; k < 4u; k = k + 1u) {
-    {% if source.alignCorners %}
     reflected_y[k] = u32(clamp(i32(reflect_coord(f32(y0 + i32(k)), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
     reflected_x[k] = u32(clamp(i32(reflect_coord(f32(x0 + i32(k)), 0.0, f32(params.inW) - 1.0)), 0, i32(params.inW) - 1));
     {% else %}
@@ -251,7 +248,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 {% endif %}
   var result = {{ channelVec }}(0.0);
   for (var r = 0i; r < 4i; r = r + 1i) {
-{% if source.paddingMode == "reflection" %}
     let row = cx.x * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[0])
       + cx.y * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[1])
       + cx.z * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[2])
@@ -269,8 +266,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let y0 = sample_floor(sy);
   let wx = sx - f32(x0);
   let wy = sy - f32(y0);
-{% if source.paddingMode == "reflection" %}
-  {% if source.alignCorners %}
   let reflected_y0 = u32(clamp(i32(reflect_coord(f32(y0), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
   let reflected_y1 = u32(clamp(i32(reflect_coord(f32(y0 + 1), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
   let reflected_x0 = u32(clamp(i32(reflect_coord(f32(x0), 0.0, f32(params.inW) - 1.0)), 0, i32(params.inW) - 1));
@@ -295,8 +292,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
     + wy * ((1.0 - wx) * v10 + wx * v11);
   {% endif %}
   let out_base = (n * params.C + c) * out_plane + spatial;
-{% for lane in range(source.channelWidth) %}
-{% if source.channelTail and lane > 0 %}
   if (c + {{ lane }}u < params.C) { y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result.{{ components[lane] }}); }
 {% else %}
   y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result.{{ components[lane] }});
@@ -312,9 +309,9 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
   let sy = sanitize_coord(denormalize(f32(grid[grid_base + 1u]), params.inH));
   let img_base = (n * params.C + c) * params.inH * params.inW;
-  {% if source.mode == "nearest" %}
   let result = pixel(img_base, sample_round(sy), sample_round(sx));
-  {% elif source.mode == "cubic" %}
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let dx = sx - f32(x0 + 1);

 {% 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 {
+  {% if alignCorners %}
   return (n + 1.0) * 0.5 * f32(length - 1u);
   {% else %}
   return ((n + 1.0) * f32(length) - 1.0) * 0.5;
 fn sanitize_coord(v: f32) -> f32 {
   if (!(v >= -2147483000.0 && v <= 2147483000.0)) {
+    {% if alignCorners %}
     return 0.0;
     {% else %}
     return -0.5;
   }
   return v;
 }
+{% if paddingMode == "reflection" %}
 fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   let range = hi - lo;
 }
 {% endif %}
+{% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
 {%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
 {%- endif %}
+{% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.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; }
   let hh = u32(h);
   let ww = u32(w);
+  {% elif paddingMode == "border" %}
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
+  {% if alignCorners %}
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
   {% else %}
 {% 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" %}
 fn pixel_channels(n: u32, c0: u32, h: i32, w: i32) -> {{ channelVec }} {
+  {% if paddingMode == "zeros" %}
   if (h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) { return {{ channelVec }}(0.0); }
   let hh = u32(h);
   let ww = u32(w);
+  {% elif paddingMode == "border" %}
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
+  {% if alignCorners %}
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
   {% else %}
   let offset = hh * params.inW + ww;
   let nb = n * params.C;
   return {{ channelVec }}(
+{% for lane in range(channelWidthSpec) %}
+    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;
   let nb = n * params.C;
   return {{ channelVec }}(
+{% for lane in range(channelWidthSpec) %}
+    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("i", "") }}
+{% if channelWidthSpec is defined %}
   let out_plane = params.outH * params.outW;
+  let channel_blocks = (params.C + {{ channelWidthSpec - 1 }}u) / {{ channelWidthSpec }}u;
   let block_count = (params.count / params.C) * channel_blocks;
   if (i >= block_count) { return; }
   let spatial = i % out_plane;
   let cb = (i / out_plane) % channel_blocks;
   let n = i / (out_plane * channel_blocks);
+  let c = cb * {{ channelWidthSpec }}u;
   let grid_base = (n * out_plane + spatial) * 2u;
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
   let sy = sanitize_coord(denormalize(f32(grid[grid_base + 1u]), params.inH));
+  {% if modeSpec == "nearest" %}
   let result = pixel_channels(n, c, sample_round(sy), sample_round(sx));
+  {% elif modeSpec == "cubic" %}
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let dx = sx - f32(x0 + 1);
   let dy = sy - f32(y0 + 1);
   let cx = cubic_coeffs(dx);
   let cy = cubic_coeffs(dy);
+{% if paddingMode == "reflection" %}
   var reflected_y: array<u32, 4>;
   var reflected_x: array<u32, 4>;
   for (var k = 0u; k < 4u; k = k + 1u) {
+    {% if alignCorners %}
     reflected_y[k] = u32(clamp(i32(reflect_coord(f32(y0 + i32(k)), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
     reflected_x[k] = u32(clamp(i32(reflect_coord(f32(x0 + i32(k)), 0.0, f32(params.inW) - 1.0)), 0, i32(params.inW) - 1));
     {% else %}
 {% endif %}
   var result = {{ channelVec }}(0.0);
   for (var r = 0i; r < 4i; r = r + 1i) {
+{% if paddingMode == "reflection" %}
     let row = cx.x * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[0])
       + cx.y * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[1])
       + cx.z * pixel_channels_resolved(n, c, reflected_y[u32(r)], reflected_x[2])
   let y0 = sample_floor(sy);
   let wx = sx - f32(x0);
   let wy = sy - f32(y0);
+{% if paddingMode == "reflection" %}
+  {% if alignCorners %}
   let reflected_y0 = u32(clamp(i32(reflect_coord(f32(y0), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
   let reflected_y1 = u32(clamp(i32(reflect_coord(f32(y0 + 1), 0.0, f32(params.inH) - 1.0)), 0, i32(params.inH) - 1));
   let reflected_x0 = u32(clamp(i32(reflect_coord(f32(x0), 0.0, f32(params.inW) - 1.0)), 0, i32(params.inW) - 1));
     + wy * ((1.0 - wx) * v10 + wx * v11);
   {% endif %}
   let out_base = (n * params.C + c) * out_plane + spatial;
+{% for lane in range(channelWidthSpec) %}
+{% if channelTail and lane > 0 %}
   if (c + {{ lane }}u < params.C) { y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result.{{ components[lane] }}); }
 {% else %}
   y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result.{{ components[lane] }});
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
   let sy = sanitize_coord(denormalize(f32(grid[grid_base + 1u]), params.inH));
   let img_base = (n * params.C + c) * params.inH * params.inW;
+  {% if modeSpec == "nearest" %}
   let result = pixel(img_base, sample_round(sy), sample_round(sx));
+  {% elif modeSpec == "cubic" %}
   let x0 = sample_floor(sx) - 1;
   let y0 = sample_floor(sy) - 1;
   let dx = sx - f32(x0 + 1);

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

@@ -1,43 +1,40 @@
 {% 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
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
 {% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; 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
-  // maxComputeWorkgroupsPerDimension limit (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
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "dispatch" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit.
 {% endif %}
 {% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {%- endif %}
 {% endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 fn denormalize(n: f32, length: u32) -> f32 {
-  {% if source.alignCorners %}
   return (n + 1.0) * 0.5 * f32(length - 1u);
   {% else %}
   return ((n + 1.0) * f32(length) - 1.0) * 0.5;
@@ -46,7 +43,7 @@ fn denormalize(n: f32, length: u32) -> f32 {
 fn sanitize_coord(v: f32) -> f32 {
   if (!(v >= -2147483000.0 && v <= 2147483000.0)) {
-    {% if source.alignCorners %}
     return 0.0;
     {% else %}
     return -0.5;
@@ -54,7 +51,7 @@ fn sanitize_coord(v: f32) -> f32 {
   }
   return v;
 }
-{% if source.paddingMode == "reflection" %}
 fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   let range = hi - lo;
@@ -77,17 +74,17 @@ fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
 }
 {% endif %}
-{% if source.mode != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
@@ -98,14 +95,14 @@ fn sample_floor(v: f32) -> i32 {
 {%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
-    {% if source.paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
@@ -116,7 +113,7 @@ fn sample_round(v: f32) -> i32 {
 {%- endif %}
-{% if source.mode == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.0);
@@ -133,20 +130,23 @@ fn cubic_one(x0: f32, a: f32) -> f32 {
 }
 {% endif -%}
-fn voxel(base: u32, d: i32, h: i32, w: i32) -> f32 {
-  {% if source.paddingMode == "zeros" %}
   if (d < 0 || d >= i32(params.inD) || h < 0 || h >= i32(params.inH) || w < 0 || w >= i32(params.inW)) {
-    return 0.0;
   }
   let dd = u32(d);
   let hh = u32(h);
   let ww = u32(w);
-  {% elif source.paddingMode == "border" %}
   let dd = u32(clamp(d, 0, i32(params.inD) - 1));
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
-  {% if source.alignCorners %}
   let rd = i32(reflect_coord(f32(d), 0.0, f32(params.inD) - 1.0));
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
@@ -159,11 +159,20 @@ fn voxel(base: u32, d: i32, h: i32, w: i32) -> f32 {
   let hh = u32(clamp(rh, 0, i32(params.inH) - 1));
   let ww = u32(clamp(rw, 0, i32(params.inW) - 1));
   {% endif %}
   return f32(x[((base + dd) * params.inH + hh) * params.inW + ww]);
 }
-@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d(guardInline=true) }}
   let ow = i % params.outW;
@@ -172,8 +181,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   t = t / params.outH;
   let od = t % params.outD;
   t = t / params.outD;
   let c = t % params.C;
   let n = t / params.C;
   let grid_base = (((n * params.outD + od) * params.outH + oh) * params.outW + ow) * 3u;
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
@@ -181,11 +196,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let sz = sanitize_coord(denormalize(f32(grid[grid_base + 2u]), params.inD));
   let img_base = (n * params.C + c) * params.inD;
-  {% if source.mode == "nearest" %}
-  let result = voxel(img_base, sample_round(sz), sample_round(sy), sample_round(sx));
-  {% elif source.mode == "cubic" %}
-  // Tricubic is the separable product of the same 4-tap Keys kernel the 2D path
-  // uses, one axis at a time; 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;
@@ -193,14 +208,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   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 = 0.0;
   for (var kz = 0i; kz < 4i; kz = kz + 1i) {
-    var plane = 0.0;
     for (var ky = 0i; ky < 4i; ky = ky + 1i) {
-      let row = cx.x * voxel(img_base, z0 + kz, y0 + ky, x0)
-        + cx.y * voxel(img_base, z0 + kz, y0 + ky, x0 + 1)
-        + cx.z * voxel(img_base, z0 + kz, y0 + ky, x0 + 2)
-        + cx.w * voxel(img_base, z0 + kz, y0 + ky, x0 + 3);
       plane = plane + cy[u32(ky)] * row;
     }
     result = result + cz[u32(kz)] * plane;
@@ -212,17 +227,29 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let wx = sx - f32(x0);
   let wy = sy - f32(y0);
   let wz = sz - f32(z0);
-  let v000 = voxel(img_base, z0, y0, x0);
-  let v001 = voxel(img_base, z0, y0, x0 + 1);
-  let v010 = voxel(img_base, z0, y0 + 1, x0);
-  let v011 = voxel(img_base, z0, y0 + 1, x0 + 1);
-  let v100 = voxel(img_base, z0 + 1, y0, x0);
-  let v101 = voxel(img_base, z0 + 1, y0, x0 + 1);
-  let v110 = voxel(img_base, z0 + 1, y0 + 1, x0);
-  let v111 = voxel(img_base, 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;
   {% endif %}
   y[i] = {{ scalar }}(result);
 }

 {% 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 {
+  {% if alignCorners %}
   return (n + 1.0) * 0.5 * f32(length - 1u);
   {% else %}
   return ((n + 1.0) * f32(length) - 1.0) * 0.5;
 fn sanitize_coord(v: f32) -> f32 {
   if (!(v >= -2147483000.0 && v <= 2147483000.0)) {
+    {% if alignCorners %}
     return 0.0;
     {% else %}
     return -0.5;
   }
   return v;
 }
+{% if paddingMode == "reflection" %}
 fn reflect_coord(v: f32, lo: f32, hi: f32) -> f32 {
   let range = hi - lo;
 }
 {% endif %}
+{% if modeSpec != "nearest" %}
 fn sample_floor(v: f32) -> i32 {
   if (v > 2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
 {%- else %}
 fn sample_round(v: f32) -> i32 {
   if (v > 2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return 2147483000;
     {% endif %}
   }
   if (v < -2147483000.0) {
+    {% if paddingMode == "reflection" %}
     return 0;
     {% else %}
     return -2147483000;
 {%- endif %}
+{% if modeSpec == "cubic" %}
 fn cubic_coeffs(t: f32) -> vec4<f32> {
   let a = -0.75;
   let x0 = abs(t + 1.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 %};
   }
   let dd = u32(d);
   let hh = u32(h);
   let ww = u32(w);
+  {% elif paddingMode == "border" %}
   let dd = u32(clamp(d, 0, i32(params.inD) - 1));
   let hh = u32(clamp(h, 0, i32(params.inH) - 1));
   let ww = u32(clamp(w, 0, i32(params.inW) - 1));
   {% else %}
+  {% if alignCorners %}
   let rd = i32(reflect_coord(f32(d), 0.0, f32(params.inD) - 1.0));
   let rh = i32(reflect_coord(f32(h), 0.0, f32(params.inH) - 1.0));
   let rw = i32(reflect_coord(f32(w), 0.0, f32(params.inW) - 1.0));
   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.
+  return {{ volumeT }}(
+{% for lane in range(channelWidthSpec) %}
+    f32(x[((base + dd{% if lane > 0 %} + {% if channelTail %}min({{ lane }}u, params.C - 1u - c0){% else %}{{ lane }}u{% endif %} * params.inD{% endif %}) * params.inH + hh) * params.inW + ww]){% if not loop.last %},{% else %});{% endif %}
+{% endfor %}
+{% else %}
   return f32(x[((base + dd) * params.inH + hh) * params.inW + ww]);
+{% endif %}
 }
+@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;
   t = t / params.outH;
   let od = t % params.outD;
   t = t / params.outD;
+{% if volumeVector %}
+  let channel_blocks = (params.C + {{ channelWidthSpec - 1 }}u) / {{ channelWidthSpec }}u;
+  let c = (t % channel_blocks) * {{ channelWidthSpec }}u;
+  let n = t / channel_blocks;
+{% else %}
   let c = t % params.C;
   let n = t / params.C;
+{% endif %}
   let grid_base = (((n * params.outD + od) * params.outH + oh) * params.outW + ow) * 3u;
   let sx = sanitize_coord(denormalize(f32(grid[grid_base]), params.inW));
   let sz = sanitize_coord(denormalize(f32(grid[grid_base + 2u]), params.inD));
   let img_base = (n * params.C + c) * params.inD;
+  {% if modeSpec == "nearest" %}
+  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 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;
   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;
   {% endif %}
+{% if volumeVector %}
+  let out_plane = params.outD * params.outH * params.outW;
+  let out_base = (n * params.C + c) * out_plane + (od * params.outH + oh) * params.outW + ow;
+{% for lane in range(channelWidthSpec) %}
+{% if channelTail and lane > 0 %}
+  if (c + {{ lane }}u < params.C) { y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result[{{ lane }}u]); }
+{% else %}
+  y[out_base + {{ lane }}u * out_plane] = {{ scalar }}(result[{{ lane }}u]);
+{% endif %}
+{% endfor %}
+{% else %}
   y[i] = {{ scalar }}(result);
+{% endif %}
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,33 +2,19 @@
   "domain": "ai.onnx",
   "name": "GridSample",
   "sinceVersion": 20,
-  "description": "Samples values from input tensor `X` at positions defined by a flow-field `grid`, producing output `Y` with spatial dimensions taken from `grid`. Grid coordinates are normalized to `[-1, 1]` over the input spatial extent; positions outside this range are handled according to `padding_mode`. Supports spatial `(rank-4, NCHW)` and volumetric `(rank-5, NCDHW)` inputs with `linear`, `nearest`, or `cubic` interpolation.",
-  "inputs": [
-    {
-      "role": "X",
-      "dtype": "T",
-      "description": "Input tensor of shape `(N, C, D1, ..., Dr)` whose values are sampled."
-    },
-    {
-      "role": "grid",
-      "dtype": "T",
-      "description": "Flow-field of shape `(N, D1_out, ..., Dr_out, r)` with normalized sampling coordinates in `[-1, 1]`."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
       "dtype": "T",
       "rank": "ranks.grid",
-      "description": "Output tensor of shape `(N, C, D1_out, ..., Dr_out)` containing the interpolated samples.",
-      "shape": "prefix(shapes.X, 2) + prefix(suffix(shapes.grid, 1), ranks.grid - 2)"
     }
-  ],
-  "attributes": { "mode": "linear", "padding_mode": "zeros", "align_corners": 0 },
-  "attributeDescriptions": {
-    "mode": "Interpolation method: `linear` (bilinear or trilinear, depending on rank), `nearest`, or `cubic`. Cubic interpolation is supported for rank-4 (2-D spatial) inputs.",
-    "padding_mode": "How out-of-bound grid positions are handled: `zeros` pads with 0, `border` clamps to the border value, or `reflection` reflects coordinates back into the valid range.",
-    "align_corners": "When 1, extrema values `-1` and `1` map to the center of the corner pixels; when 0 (default) they map to the outer edge of corner pixels, making sampling resolution-agnostic."
   },
   "attributeConstraints": {
     "mode": { "values": ["linear", "nearest", "cubic"] },
@@ -36,71 +22,52 @@
     "align_corners": { "values": [0, 1] }
   },
   "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "grid": { "kind": "tensor", "semantic": "grid", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 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\"))",
-    "rank4Ok": "ranks.X == 4 and ranks.grid == 4 and ranks.Y == 4 and dim(shapes.grid, 0) == dim(shapes.X, 0) and dim(shapes.grid, 3) == 2 and dim(shapes.Y, 0) == dim(shapes.X, 0) and dim(shapes.Y, 1) == dim(shapes.X, 1) and dim(shapes.Y, 2) == dim(shapes.grid, 1) and dim(shapes.Y, 3) == dim(shapes.grid, 2) and f16Ok(dtypes.T)",
-    "channelWidth": 4
   },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
-  "bindingSets": {
-    "rank4": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
-      {
-        "name": "grid",
-        "arg": "grid",
-        "semantic": "grid",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-            { "name": "C", "type": "u32", "value": "dim(shapes.X, 1)" },
-            { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
-            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
-            { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
-            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" }
-          ]
-        }
-      }
-    ]
   },
   "variants": [
     {
       "id": "nchw_rank4_channel_vector",
       "priority": 5,
-      "when": ["rank4Ok", "not reportedNonWave32Adapter", "dim(shapes.X, 1) >= 2"],
       "passes": [
         {
           "id": "main",
           "name": "GridSample.ChannelX4",
-          "source": {
-            "shader": "grid-sample.wgsl.jinja",
-            "inputs": {
-              "mode": "attrs.mode",
-              "paddingMode": "attrs.padding_mode",
-              "alignCorners": "attrs.align_corners != 0",
-              "channelWidth": "channelWidth",
-              "channelTail": "dim(shapes.X, 1) % channelWidth != 0"
-            }
           },
-          "bindings": "rank4",
           "dispatch": {
-            "threads": "dim(shapes.Y, 0) * ceilDiv(dim(shapes.Y, 1), channelWidth) * dim(shapes.Y, 2) * dim(shapes.Y, 3)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]
@@ -112,71 +79,104 @@
         {
           "id": "main",
           "name": "GridSample",
-          "source": {
-            "shader": "grid-sample.wgsl.jinja",
-            "inputs": {
-              "mode": "attrs.mode",
-              "paddingMode": "attrs.padding_mode",
-              "alignCorners": "attrs.align_corners != 0"
-            }
           },
-          "bindings": "rank4",
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "ncdhw_rank5",
       "priority": 10,
-      "when": ["ranks.X == 5", "ranks.grid == 5", "ranks.Y == 5", "dim(shapes.grid, 0) == dim(shapes.X, 0)", "dim(shapes.grid, 4) == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.Y, 2) == dim(shapes.grid, 1)", "dim(shapes.Y, 3) == dim(shapes.grid, 2)", "dim(shapes.Y, 4) == dim(shapes.grid, 3)", "(attrs.mode == \"linear\" or attrs.mode == \"nearest\" or attrs.mode == \"cubic\")", "(attrs.padding_mode == \"zeros\" or attrs.padding_mode == \"border\" or attrs.padding_mode == \"reflection\")", "f16Ok(dtypes.T)"],
       "passes": [
         {
           "id": "main",
           "name": "GridSample.Volumetric",
-          "source": {
-            "shader": "grid-sample3d.wgsl.jinja",
-            "inputs": {
-              "mode": "attrs.mode",
-              "paddingMode": "attrs.padding_mode",
-              "alignCorners": "attrs.align_corners != 0"
-            }
           },
           "bindings": [
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "grid",
-              "arg": "grid",
-              "semantic": "grid",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
             {
               "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [
-                  { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-                  { "name": "C", "type": "u32", "value": "dim(shapes.X, 1)" },
-                  { "name": "inD", "type": "u32", "value": "dim(shapes.X, 2)" },
-                  { "name": "inH", "type": "u32", "value": "dim(shapes.X, 3)" },
-                  { "name": "inW", "type": "u32", "value": "dim(shapes.X, 4)" },
-                  { "name": "outD", "type": "u32", "value": "dim(shapes.Y, 2)" },
-                  { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 3)" },
-                  { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 4)" }
-                ]
-              }
             }
           ],
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "GridSample",
   "sinceVersion": 20,
+  "inputs": { "x": { "onnx": "X", "dtype": "T" }, "grid": { "dtype": "T" } },
+  "outputs": {
+    "y": {
+      "onnx": "Y",
       "dtype": "T",
       "rank": "ranks.grid",
+      "shape": "prefix(shapes.x, 2) + prefix(suffix(shapes.grid, 1), ranks.grid - 2)"
     }
+  },
+  "attributes": {
+    "mode": { "default": "linear" },
+    "padding_mode": { "default": "zeros" },
+    "align_corners": { "default": 0 }
   },
   "attributeConstraints": {
     "mode": { "values": ["linear", "nearest", "cubic"] },
     "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)",
+    "rank4Ok": "ranks.x == 4 and ranks.grid == 4 and ranks.y == 4 and dim(shapes.grid, 0) == dim(shapes.x, 0) and dim(shapes.grid, 3) == 2 and dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.x, 1) and dim(shapes.y, 2) == dim(shapes.grid, 1) and dim(shapes.y, 3) == dim(shapes.grid, 2) and f16Ok(dtypes.T)",
+    "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)" },
+        { "name": "inH", "type": "u32", "value": "dim(shapes.x, 2)" },
+        { "name": "inW", "type": "u32", "value": "dim(shapes.x, 3)" },
+        { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
+        { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" }
+      ]
+    }
   },
   "variants": [
     {
       "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)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), channelWidth) * dim(shapes.y, 2) * dim(shapes.y, 3)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
         {
           "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)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw_rank5_channel_vector",
+      "priority": 15,
+      "when": ["ranks.x == 5", "ranks.grid == 5", "ranks.y == 5", "dim(shapes.grid, 0) == dim(shapes.x, 0)", "dim(shapes.grid, 4) == 3", "dim(shapes.y, 0) == dim(shapes.x, 0)", "dim(shapes.y, 1) == dim(shapes.x, 1)", "dim(shapes.y, 2) == dim(shapes.grid, 1)", "dim(shapes.y, 3) == dim(shapes.grid, 2)", "dim(shapes.y, 4) == dim(shapes.grid, 3)", "(attrs.mode == \"linear\" or attrs.mode == \"nearest\" or attrs.mode == \"cubic\")", "(attrs.padding_mode == \"zeros\" or attrs.padding_mode == \"border\" or attrs.padding_mode == \"reflection\")", "f16Ok(dtypes.T)", "not reportedNonWave32Adapter", "dim(shapes.x, 1) >= 2", "attrs.mode != \"cubic\" or attrs.padding_mode != \"border\""],
+      "passes": [
+        {
+          "id": "main",
+          "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",
+            "grid",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                {
+                  "name": "count",
+                  "type": "u32",
+                  "value": "dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), volumeWidth) * dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)"
+                },
+                { "name": "C", "type": "u32", "value": "dim(shapes.x, 1)" },
+                { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
+                { "name": "inH", "type": "u32", "value": "dim(shapes.x, 3)" },
+                { "name": "inW", "type": "u32", "value": "dim(shapes.x, 4)" },
+                { "name": "outD", "type": "u32", "value": "dim(shapes.y, 2)" },
+                { "name": "outH", "type": "u32", "value": "dim(shapes.y, 3)" },
+                { "name": "outW", "type": "u32", "value": "dim(shapes.y, 4)" }
+              ]
+            }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), volumeWidth) * dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)), (volumeWorkgroup)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.y, 0) * ceilDiv(dim(shapes.y, 1), volumeWidth) * dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)), (volumeWorkgroup)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "ncdhw_rank5",
       "priority": 10,
+      "when": ["ranks.x == 5", "ranks.grid == 5", "ranks.y == 5", "dim(shapes.grid, 0) == dim(shapes.x, 0)", "dim(shapes.grid, 4) == 3", "dim(shapes.y, 0) == dim(shapes.x, 0)", "dim(shapes.y, 1) == dim(shapes.x, 1)", "dim(shapes.y, 2) == dim(shapes.grid, 1)", "dim(shapes.y, 3) == dim(shapes.grid, 2)", "dim(shapes.y, 4) == dim(shapes.grid, 3)", "(attrs.mode == \"linear\" or attrs.mode == \"nearest\" or attrs.mode == \"cubic\")", "(attrs.padding_mode == \"zeros\" or attrs.padding_mode == \"border\" or attrs.padding_mode == \"reflection\")", "f16Ok(dtypes.T)"],
       "passes": [
         {
           "id": "main",
           "name": "GridSample.Volumetric",
+          "shader": "grid-sample3d.wgsl.jinja",
+          "derive": {
+            "modeSpec": "attrs.mode",
+            "paddingMode": "attrs.padding_mode",
+            "alignCorners": "attrs.align_corners != 0"
           },
           "bindings": [
+            "x",
+            "grid",
+            "y",
             {
               "name": "params",
+              "struct": [
+                { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+                { "name": "C", "type": "u32", "value": "dim(shapes.x, 1)" },
+                { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
+                { "name": "inH", "type": "u32", "value": "dim(shapes.x, 3)" },
+                { "name": "inW", "type": "u32", "value": "dim(shapes.x, 4)" },
+                { "name": "outD", "type": "u32", "value": "dim(shapes.y, 2)" },
+                { "name": "outH", "type": "u32", "value": "dim(shapes.y, 3)" },
+                { "name": "outW", "type": "u32", "value": "dim(shapes.y, 4)" }
+              ]
             }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,27 @@
 {
   "name": "ai.onnx.GridSample",
-  "id": "_ai_onnx_gridsample_webgpu_3892192",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "G6tTb9EsyvpKUZLsd4qAik9iHOhPJGsUBR3ah5o7L+w=",
-      "grid-sample.wgsl.jinja": "vGFCLAzhkxtkeC071yc3MrcZa3C5mqPT4+ZbjxJn5Zo=",
-      "grid-sample3d.wgsl.jinja": "cSEW1lJIrYaiamCfxoP1KBWbzqcRfe+00AtqpeXRUNA=",
-      "manifest.json": "iZOT6Iz9u/5yFWpFbpO+1EgSPRoW9YPluK9w2unD2wk=",
-      "test.json": "iD0xQONckmqQHJylOBRTC7z57/IvRn28fEdFgYX1TNI="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.GridSample" }
 }

 {
   "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"],
+      "ncdhw_rank5_channel_vector": ["grid-sample3d.wgsl.jinja"],
+      "ncdhw_rank5": ["grid-sample3d.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.GridSample",
   "fixtureArrays": {
     "onnx_backend_gridsample_volumetric_input_grid": [-1, -1, -1, -1, -0.5, 0.30000001192092896, -0.5, -0.5, -0.5, 1, -0.6000000238418579, -1, -0.20000000298023224, -0.20000000298023224, -0.20000000298023224, 0.4000000059604645, 0.20000000298023224, 0.6000000238418579, 0, 0, 0, -1, 0, 0, 0, 0, 0, -1, 1, 0, -0.20000000298023224, -0.20000000298023224, -0.20000000298023224, 1, 0.4000000059604645, -0.20000000298023224, 0.5, 0.5, 0.5, -1, -0.800000011920929, 0.800000011920929, 1, 1, 1, 0.4000000059604645, 0.6000000238418579, -0.30000001192092896],
     "grid_sample_grid_16": [-1, -0.800000011920929, -0.6000000238418579, -0.5, -0.10000000149011612, -0.20000000298023224, 0.699999988079071, 0, 0, 0.4000000059604645, 0.20000000298023224, -0.20000000298023224, -0.30000001192092896, 0.5, -1, 1],
@@ -94,7 +93,7 @@
         "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormals to zero in the ALU; bilinear/trilinear interpolation arithmetic flushes the subnormal pixel/voxel on GPU. Permanent FTZ limitation."
       },
       "provenance": {
-        "notes": "Rank-5 companion: exact corner sampling should copy the positive subnormal source voxel through the trilinear path."
       },
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
       "inputs": {
@@ -1729,11 +1728,7 @@
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
       "inputs": {
         "x": { "dtype": "float32", "shape": [1, 1, 0, 2], "data": { "kind": "values", "values": [] } },
-        "grid": {
-          "dtype": "float32",
-          "shape": [1, 2, 2, 2],
-          "data": { "kind": "values", "values": [-1.0, -1.0, 1.0, -1.0, -1.0, 1.0, 1.0, 1.0] }
-        }
       },
       "outputs": {
         "y": {
@@ -1890,7 +1885,7 @@
       "name": "rank5_cubic_border_pinned_taps",
       "provenance": {
         "source": "ONNX GridSample cubic definition (Keys kernel, cubic_coeff_a = -0.75), evaluated independently",
-        "notes": "Ground truth computed independently from the separable 4-tap definition, so neither the kernel nor the reference is its own oracle. Sample points are dyadic and the field is in quarters, so the values are exact in binary; two of the four points push taps past the volume, which border clamps. Trilinear on the same points gives different answers, so a cubic path that silently ran as linear would fail here."
       },
       "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 0 },
       "inputs": {
@@ -1960,6 +1955,790 @@
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2, 2], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     }
   ]
 }

 {
   "fixtureArrays": {
     "onnx_backend_gridsample_volumetric_input_grid": [-1, -1, -1, -1, -0.5, 0.30000001192092896, -0.5, -0.5, -0.5, 1, -0.6000000238418579, -1, -0.20000000298023224, -0.20000000298023224, -0.20000000298023224, 0.4000000059604645, 0.20000000298023224, 0.6000000238418579, 0, 0, 0, -1, 0, 0, 0, 0, 0, -1, 1, 0, -0.20000000298023224, -0.20000000298023224, -0.20000000298023224, 1, 0.4000000059604645, -0.20000000298023224, 0.5, 0.5, 0.5, -1, -0.800000011920929, 0.800000011920929, 1, 1, 1, 0.4000000059604645, 0.6000000238418579, -0.30000001192092896],
     "grid_sample_grid_16": [-1, -0.800000011920929, -0.6000000238418579, -0.5, -0.10000000149011612, -0.20000000298023224, 0.699999988079071, 0, 0, 0.4000000059604645, 0.20000000298023224, -0.20000000298023224, -0.30000001192092896, 0.5, -1, 1],
         "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormals to zero in the ALU; bilinear/trilinear interpolation arithmetic flushes the subnormal pixel/voxel on GPU. Permanent FTZ limitation."
       },
       "provenance": {
+        "notes": "Exact corner sampling on a rank-5 input must copy the positive subnormal source voxel through the trilinear path."
       },
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
       "inputs": {
       "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
       "inputs": {
         "x": { "dtype": "float32", "shape": [1, 1, 0, 2], "data": { "kind": "values", "values": [] } },
+        "grid": { "dtype": "float32", "shape": [1, 2, 2, 2], "data": { "kind": "constant", "value": -1.0 } }
       },
       "outputs": {
         "y": {
       "name": "rank5_cubic_border_pinned_taps",
       "provenance": {
         "source": "ONNX GridSample cubic definition (Keys kernel, cubic_coeff_a = -0.75), evaluated independently",
+        "notes": "Expected values are derived from the separable four-tap cubic definition. Dyadic sample points and quarter-valued voxels make the arithmetic exact; two points require border clamping, and trilinear interpolation gives different results."
       },
       "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 0 },
       "inputs": {
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2, 2], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_zeros_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_zeros_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_zeros_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_zeros_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_border_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_border_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_border_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_border_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_reflection_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_reflection_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_reflection_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_linear_reflection_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_zeros_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_zeros_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_zeros_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_zeros_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_border_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_border_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_border_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_border_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_reflection_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_reflection_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_reflection_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_nearest_reflection_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "nearest", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_zeros_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_zeros_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_zeros_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_zeros_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "zeros", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_border_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_border_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "border", "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_border_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_border_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "border", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_reflection_ac0_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_reflection_ac0_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_reflection_ac1_float32",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_tail_cubic_reflection_ac1_float16",
+      "provenance": {
+        "notes": "Volumetric channel sharing across batches with a three-channel tail, all interpolation/padding modes and both coordinate conventions."
+      },
+      "attrs": { "mode": "cubic", "padding_mode": "reflection", "align_corners": 1 },
+      "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 } }
+    },
+    {
+      "name": "volume_channels_linear_reflection_c1",
+      "provenance": {
+        "notes": "Checks scalar, two-lane, complete four-lane and masked four-lane volumetric channel selection."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 1, 3, 5, 7], "data": { "kind": "fillFloat32", "scale": 0.3 } },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 5, 7, 9, 3],
+          "data": { "kind": "linspace", "start": -2.5, "end": 2.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_channels_linear_reflection_c2",
+      "provenance": {
+        "notes": "Checks scalar, two-lane, complete four-lane and masked four-lane volumetric channel selection."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 2, 3, 5, 7], "data": { "kind": "fillFloat32", "scale": 0.3 } },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 5, 7, 9, 3],
+          "data": { "kind": "linspace", "start": -2.5, "end": 2.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_channels_linear_reflection_c4",
+      "provenance": {
+        "notes": "Checks scalar, two-lane, complete four-lane and masked four-lane volumetric channel selection."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 4, 3, 5, 7], "data": { "kind": "fillFloat32", "scale": 0.3 } },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 5, 7, 9, 3],
+          "data": { "kind": "linspace", "start": -2.5, "end": 2.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "volume_channels_linear_reflection_c5",
+      "provenance": {
+        "notes": "Checks scalar, two-lane, complete four-lane and masked four-lane volumetric channel selection."
+      },
+      "attrs": { "mode": "linear", "padding_mode": "reflection", "align_corners": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 5, 3, 5, 7], "data": { "kind": "fillFloat32", "scale": 0.3 } },
+        "grid": {
+          "dtype": "float32",
+          "shape": [2, 5, 7, 9, 3],
+          "data": { "kind": "linspace", "start": -2.5, "end": 2.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 5, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     }
   ]
 }