sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +15 -12
build/webgpu/bench.json +0 -1
build/webgpu/manifest.json +57 -138
build/webgpu/metadata.json +15 -8
build/webgpu/pad-quad-vec4.wgsl.jinja +12 -15
build/webgpu/pad.wgsl.jinja +15 -17
build/webgpu/test.json +27 -34

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `Pad` spec](https://onnx.ai/onnx/operators/onnx__Pad.html) for the
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `data` | `data` | `T` | — | — | Input tensor to be padded. | required |
-| `constant_value` | `constant_value` | `T` | `0` | — | Optional scalar value used as the fill constant when `mode` is `constant`; defaults to 0. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `output` | `output` | `T` | same as `data` | — | Tensor after padding, with each axis enlarged by the corresponding begin and end pad counts. | required |
 ## Attributes
@@ -46,7 +46,7 @@ Attributes and default values (overridable per request):
 ## 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
@@ -55,15 +55,18 @@ Attributes and default values (overridable per request):
 ## Use with `@huggingface/kernels`
-The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
-The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
-- `output`
-Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
 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 | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `data` | `T` | — | — | Input tensor to be padded. | required |
+| `constant_value` | `T` | `0` | — | Optional scalar value used as the fill constant when `mode` is `constant`; defaults to 0. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `output` | `T` | same as `data` | — | Tensor after padding, with each axis enlarged by the corresponding begin and end pad counts. | required |
 ## Attributes
 ## 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
+```
+Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
+This example supplies explicit metadata for:
+- `output`
 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.Pad",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

build/webgpu/manifest.json CHANGED Viewed

@@ -2,139 +2,54 @@
   "domain": "ai.onnx",
   "name": "Pad",
   "sinceVersion": 19,
-  "description": "Pads a tensor along each axis using one of four modes: `constant` (fill with a scalar value), `reflect` (mirror edge values), `edge` (replicate boundary values), or `wrap` (torus-like wrap-around). Supply the required ONNX `pads` tensor, and optional `axes` tensor when present, through the `pads` request attribute as a full-rank list. Negative counts crop the corresponding edge.",
-  "inputs": [
-    { "role": "data", "dtype": "T", "description": "Input tensor to be padded." },
-    {
-      "role": "constant_value",
-      "dtype": "T",
-      "rank": 0,
-      "optional": true,
-      "description": "Optional scalar value used as the fill constant when `mode` is `constant`; defaults to 0."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "output",
-      "dtype": "T",
-      "description": "Tensor after padding, with each axis enlarged by the corresponding begin and end pad counts.",
-      "rank": "ranks.data"
-    }
-  ],
-  "attributes": { "mode": "constant" },
-  "attributeDescriptions": {
-    "pads": "Values of the required `pads` tensor, with an optional `axes` tensor expanded to full rank, supplied as `[x1_begin, ..., xN_begin, x1_end, ..., xN_end]`.",
-    "mode": "Padding mode: `constant` (default), `reflect`, `edge`, or `wrap`."
-  },
   "attributeConstraints": {
     "pads": { "required": true },
     "mode": { "values": ["constant", "reflect", "edge", "wrap"] }
   },
   "typeConstraints": { "T": ["float32", "float16", "uint32", "int32", "uint8", "int8", "bool"] },
-  "args": {
-    "data": { "kind": "tensor", "semantic": "data", "role": "input" },
-    "constant_value": { "kind": "tensor", "semantic": "constant_value", "role": "constant_value", "required": false },
-    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "derive": {
     "shapeContract": "ranks.data == ranks.output and ranks.output >= 1",
     "defaultValueContract": "not present.constant_value and shapeContract and f16Ok(dtypes.T)",
-    "vec4Contract": "defaultValueContract and numel(shapes.output) > 0 and dim(shapes.output, ranks.output - 1) % 4 == 0"
   },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "mode": "attrs.mode" },
-  "bindingSets": {
-    "valueInput": [
-      {
-        "name": "data",
-        "arg": "data",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "padValue",
-        "arg": "constant_value",
-        "semantic": "constant_value",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar",
-        "length": 1
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
-      }
-    ],
-    "copy": [
-      {
-        "name": "data",
-        "arg": "data",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      }
-    ],
-    "scalarPad": [
-      {
-        "name": "data",
-        "arg": "data",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
-      }
-    ]
   },
   "variants": [
     {
       "id": "constant_value_input",
       "priority": 30,
       "when": ["present.constant_value", "numel(shapes.constant_value) == 1", "shapeContract", "f16Ok(dtypes.T)"],
-      "constants": { "valueFromInput": true },
       "passes": [
         {
           "id": "main",
           "name": "Pad",
-          "source": {
-            "shader": "pad.wgsl.jinja",
-            "inputs": {
-              "dataShape": "shapes.data",
-              "outputShape": "shapes.output",
-              "rank": "ranks.output",
-              "pads": "attrs.pads",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "valueInput",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -142,47 +57,51 @@
       "id": "quad_vec4",
       "priority": 11,
       "when": ["vec4Contract"],
-      "constants": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Pad.quadVec4",
-          "source": {
-            "shader": "pad-quad-vec4.wgsl.jinja",
-            "inputs": {
-              "dataShape": "shapes.data",
-              "outputShape": "shapes.output",
-              "rank": "ranks.output",
-              "pads": "attrs.pads",
-              "value": 0
-            }
           },
-          "bindings": "copy",
-          "dispatch": { "gridStride": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "constant_rank1_to_4",
       "when": ["defaultValueContract"],
-      "constants": { "valueFromInput": false },
       "passes": [
         {
           "id": "main",
           "name": "Pad",
-          "source": {
-            "shader": "pad.wgsl.jinja",
-            "inputs": {
-              "dataShape": "shapes.data",
-              "outputShape": "shapes.output",
-              "rank": "ranks.output",
-              "pads": "attrs.pads",
-              "value": 0,
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "scalarPad",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Pad",
   "sinceVersion": 19,
+  "inputs": { "data": { "dtype": "T" }, "constant_value": { "dtype": "T", "rank": 0, "optional": true } },
+  "outputs": { "output": { "dtype": "T", "rank": "ranks.data" } },
+  "attributes": { "mode": { "default": "constant" }, "pads": {} },
   "attributeConstraints": {
     "pads": { "required": true },
     "mode": { "values": ["constant", "reflect", "edge", "wrap"] }
   },
   "typeConstraints": { "T": ["float32", "float16", "uint32", "int32", "uint8", "int8", "bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "shapeContract": "ranks.data == ranks.output and ranks.output >= 1",
     "defaultValueContract": "not present.constant_value and shapeContract and f16Ok(dtypes.T)",
+    "vec4Contract": "defaultValueContract and numel(shapes.output) > 0 and dim(shapes.output, ranks.output - 1) % 4 == 0",
+    "scalar": "dtypes.T",
+    "mode": "attrs.mode"
   },
+  "bindings": {
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
   },
   "variants": [
     {
       "id": "constant_value_input",
       "priority": 30,
       "when": ["present.constant_value", "numel(shapes.constant_value) == 1", "shapeContract", "f16Ok(dtypes.T)"],
+      "derive": { "valueFromInput": true },
       "passes": [
         {
           "id": "main",
           "name": "Pad",
+          "shader": "pad.wgsl.jinja",
+          "derive": {
+            "dataShape": "shapes.data",
+            "outputShape": "shapes.output",
+            "rank": "ranks.output",
+            "pads": "attrs.pads",
+            "itemsPerInvocation": 4
           },
+          "bindings": [
+            "data",
+            { "arg": "constant_value", "name": "padValue", "elementType": "$scalar", "length": 1 },
+            "output",
+            "params"
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "quad_vec4",
       "priority": 11,
       "when": ["vec4Contract"],
+      "derive": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Pad.quadVec4",
+          "shader": "pad-quad-vec4.wgsl.jinja",
+          "derive": {
+            "dataShape": "shapes.data",
+            "outputShape": "shapes.output",
+            "rank": "ranks.output",
+            "pads": "attrs.pads",
+            "value": 0
           },
+          "bindings": ["data", { "arg": "output", "elementType": "$outputElement" }],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": 1,
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "constant_rank1_to_4",
       "when": ["defaultValueContract"],
+      "derive": { "valueFromInput": false },
       "passes": [
         {
           "id": "main",
           "name": "Pad",
+          "shader": "pad.wgsl.jinja",
+          "derive": {
+            "dataShape": "shapes.data",
+            "outputShape": "shapes.output",
+            "rank": "ranks.output",
+            "pads": "attrs.pads",
+            "value": 0,
+            "itemsPerInvocation": 4
           },
+          "bindings": ["data", "output", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,26 @@
 {
   "name": "ai.onnx.Pad",
-  "id": "_ai_onnx_pad_webgpu_df32065",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "c5OFqXWDpHEMBxvQhbwziuREKzcvtHq7AD3mAWc/O/w=",
-      "manifest.json": "uxY+ml18prHKbui31BK4eUR7F4cesNrM9HuEeRpdSKc=",
-      "pad-quad-vec4.wgsl.jinja": "1CGMCSmy+Suuii9slZko0MXqHiIf0+iAb8UjG/mKLkU=",
-      "pad.wgsl.jinja": "J1ZYf3L07tN15khGSOR1teZ15QsxAFv3MjiBrAu4ykk=",
-      "test.json": "9NikPawZjt0OZY+069sKOCAJ/K5Xc7F6usOSJ6TrvpE="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Pad" }
 }

 {
   "name": "ai.onnx.Pad",
+  "id": "_ai_onnx_pad_webgpu_30935c6",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "8tPycaU74JVcyHJOuvQfZrz8BEwiW5yt868REE1nvEI=",
+      "manifest.json": "izmrfOCxDl4LceDqRx5gtmM5Q1dfUM4oJcnITWPO8mI=",
+      "pad-quad-vec4.wgsl.jinja": "nKXmrJKUYEPgQR1ugQHSwvGjzW1gaerShR3b8fHLnoA=",
+      "pad.wgsl.jinja": "uAqBSKweEQls2Ml+WxW1nkrsPhVMvLOYPNTtGcTwSgM=",
+      "test.json": "pw+IPrkqhtk9kFg2+5puugkg5AE10N7CwLH0HWkNtks="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "constant_value_input": ["pad.wgsl.jinja"],
+      "quad_vec4": ["pad-quad-vec4.wgsl.jinja"],
+      "constant_rank1_to_4": ["pad.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/pad-quad-vec4.wgsl.jinja CHANGED Viewed

@@ -6,12 +6,9 @@
 // map_coord returns -1 outside bounds in constant mode and for empty extents.
 // Its reflect period and crop_start/sliced_dim arithmetic define the exact
 // coordinate mapping.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% set rank = source.rank %}
 fn map_coord(coord: i32, dim: i32) -> i32 {
   if (coord >= 0i && coord < dim) {
     return coord;
@@ -43,19 +40,19 @@ fn map_coord(coord: i32, dim: i32) -> i32 {
 {% set countVecs = namespace(value=1) %}
 {% for j in range(rank) %}
-{% set countVecs.value = countVecs.value * source.outputShape[j] %}
 {% endfor %}
 {% set countVecs.value = ((countVecs.value / 4) | int) %}
-{% set lastPre = source.pads[rank - 1] %}
-{% set lastPost = source.pads[2 * rank - 1] %}
 {% set lastCrop = (0 - lastPre) if lastPre < 0 else 0 %}
 {% set lastTrim = (0 - lastPost) if lastPost < 0 else 0 %}
-{% set lastSlicedRaw = source.dataShape[rank - 1] - lastCrop - lastTrim %}
 {% set lastSliced = lastSlicedRaw if lastSlicedRaw > 0 else 0 %}
 {% set lastPad = lastPre if lastPre > 0 else 0 %}
 const COUNT: u32 = {{ countVecs.value }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const OUT_W_VECS: u32 = {{ ((source.outputShape[rank - 1] / 4) | int) }}u;
 const PAD_W: i32 = {{ lastPad }}i;
 const CROP_W: i32 = {{ lastCrop }}i;
 const SLICED_W: i32 = {{ lastSliced }}i;
@@ -65,7 +62,7 @@ fn main(
   @builtin(global_invocation_id) gid: vec3<u32>,
   @builtin(num_workgroups) nwg: vec3<u32>
 ) {
-  let fill = {{ scalar }}({{ source.value }});
   let stride = nwg.x * WG;
   for (var v = gid.x; v < COUNT; v += stride) {
     let ow0 = i32((v % OUT_W_VECS) * 4u);
@@ -77,18 +74,18 @@ fn main(
 {% for axis in range(rank - 1) %}
 {% set ostride = namespace(value=1) %}
 {% for j in range(axis + 1, rank - 1) %}
-{% set ostride.value = ostride.value * source.outputShape[j] %}
 {% endfor %}
 {% set safeOStride = ostride.value %}
 {% set dstride = namespace(value=1) %}
 {% for j in range(axis + 1, rank) %}
-{% set dstride.value = dstride.value * source.dataShape[j] %}
 {% endfor %}
-{% set pre = source.pads[axis] %}
-{% set post = source.pads[axis + rank] %}
 {% set crop = (0 - pre) if pre < 0 else 0 %}
 {% set trim = (0 - post) if post < 0 else 0 %}
-{% set slicedRaw = source.dataShape[axis] - crop - trim %}
 {% set sliced = slicedRaw if slicedRaw > 0 else 0 %}
 {% set padBefore = pre if pre > 0 else 0 %}
     let out_coord{{ axis }} = i32(rem / {{ safeOStride }}u);

 // map_coord returns -1 outside bounds in constant mode and for empty extents.
 // Its reflect period and crop_start/sliced_dim arithmetic define the exact
 // coordinate mapping.
 {{ env.wgsl.resourceDeclarations }}
+{% set rank = rank %}
 fn map_coord(coord: i32, dim: i32) -> i32 {
   if (coord >= 0i && coord < dim) {
     return coord;
 {% set countVecs = namespace(value=1) %}
 {% for j in range(rank) %}
+{% set countVecs.value = countVecs.value * outputShape[j] %}
 {% endfor %}
 {% set countVecs.value = ((countVecs.value / 4) | int) %}
+{% set lastPre = pads[rank - 1] %}
+{% set lastPost = pads[2 * rank - 1] %}
 {% set lastCrop = (0 - lastPre) if lastPre < 0 else 0 %}
 {% set lastTrim = (0 - lastPost) if lastPost < 0 else 0 %}
+{% set lastSlicedRaw = dataShape[rank - 1] - lastCrop - lastTrim %}
 {% set lastSliced = lastSlicedRaw if lastSlicedRaw > 0 else 0 %}
 {% set lastPad = lastPre if lastPre > 0 else 0 %}
 const COUNT: u32 = {{ countVecs.value }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const OUT_W_VECS: u32 = {{ ((outputShape[rank - 1] / 4) | int) }}u;
 const PAD_W: i32 = {{ lastPad }}i;
 const CROP_W: i32 = {{ lastCrop }}i;
 const SLICED_W: i32 = {{ lastSliced }}i;
   @builtin(global_invocation_id) gid: vec3<u32>,
   @builtin(num_workgroups) nwg: vec3<u32>
 ) {
+  let fill = {{ scalar }}({{ value }});
   let stride = nwg.x * WG;
   for (var v = gid.x; v < COUNT; v += stride) {
     let ow0 = i32((v % OUT_W_VECS) * 4u);
 {% for axis in range(rank - 1) %}
 {% set ostride = namespace(value=1) %}
 {% for j in range(axis + 1, rank - 1) %}
+{% set ostride.value = ostride.value * outputShape[j] %}
 {% endfor %}
 {% set safeOStride = ostride.value %}
 {% set dstride = namespace(value=1) %}
 {% for j in range(axis + 1, rank) %}
+{% set dstride.value = dstride.value * dataShape[j] %}
 {% endfor %}
+{% set pre = pads[axis] %}
+{% set post = pads[axis + rank] %}
 {% set crop = (0 - pre) if pre < 0 else 0 %}
 {% set trim = (0 - post) if post < 0 else 0 %}
+{% set slicedRaw = dataShape[axis] - crop - trim %}
 {% set sliced = slicedRaw if slicedRaw > 0 else 0 %}
 {% set padBefore = pre if pre > 0 else 0 %}
     let out_coord{{ axis }} = i32(rem / {{ safeOStride }}u);

build/webgpu/pad.wgsl.jinja CHANGED Viewed

@@ -1,22 +1,20 @@
 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
-  let begin = invocation * {{ source.itemsPerInvocation }}u;
-  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 fn map_coord(coord: i32, dim: i32) -> i32 {
@@ -51,18 +49,18 @@ fn map_coord(coord: i32, dim: i32) -> i32 {
 fn data_offset_or_invalid(out_index: u32) -> i32 {
   var rem = out_index;
   var offset = 0i;
-{% for axis in range(source.rank) %}
 {% set out_stride = namespace(value=1) %}
-{% for j in range(axis + 1, source.rank) %}
-{% set out_stride.value = out_stride.value * source.outputShape[j] %}
 {% endfor %}
 {% set out_stride_divisor = out_stride.value if out_stride.value != 0 else 1 %}
 {% set data_stride = namespace(value=1) %}
-{% for j in range(axis + 1, source.rank) %}
-{% set data_stride.value = data_stride.value * source.dataShape[j] %}
 {% endfor %}
-{% set pre = source.pads[axis] %}
-{% set post = source.pads[axis + source.rank] %}
 {% if pre < 0 %}
 {% set crop_start = 0 - pre %}
 {% else %}
@@ -73,7 +71,7 @@ fn data_offset_or_invalid(out_index: u32) -> i32 {
 {% else %}
 {% set crop_end_trim = 0 %}
 {% endif %}
-{% set crop_end = source.dataShape[axis] - crop_end_trim %}
 {% set sliced_dim = crop_end - crop_start %}
 {% if sliced_dim < 0 %}
 {% set sliced_dim = 0 %}
@@ -101,7 +99,7 @@ fn data_offset_or_invalid(out_index: u32) -> i32 {
 {% if valueFromInput %}
     output[i] = padValue[0];
 {% else %}
-    output[i] = {{ scalar }}({{ source.value }});
 {% endif %}
   } else {
     output[i] = data[u32(src)];

 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
+{% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn map_coord(coord: i32, dim: i32) -> i32 {
 fn data_offset_or_invalid(out_index: u32) -> i32 {
   var rem = out_index;
   var offset = 0i;
+{% for axis in range(rank) %}
 {% set out_stride = namespace(value=1) %}
+{% for j in range(axis + 1, rank) %}
+{% set out_stride.value = out_stride.value * outputShape[j] %}
 {% endfor %}
 {% set out_stride_divisor = out_stride.value if out_stride.value != 0 else 1 %}
 {% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, rank) %}
+{% set data_stride.value = data_stride.value * dataShape[j] %}
 {% endfor %}
+{% set pre = pads[axis] %}
+{% set post = pads[axis + rank] %}
 {% if pre < 0 %}
 {% set crop_start = 0 - pre %}
 {% else %}
 {% else %}
 {% set crop_end_trim = 0 %}
 {% endif %}
+{% set crop_end = dataShape[axis] - crop_end_trim %}
 {% set sliced_dim = crop_end - crop_start %}
 {% if sliced_dim < 0 %}
 {% set sliced_dim = 0 %}
 {% if valueFromInput %}
     output[i] = padValue[0];
 {% else %}
+    output[i] = {{ scalar }}({{ value }});
 {% endif %}
   } else {
     output[i] = data[u32(src)];

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Pad",
   "fixtureArrays": {
     "onnx_backend_constant_pad_input_data": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
@@ -29,7 +28,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
-        "notes": "Uses Pad's optional `constant_value` input. Pads remain the synthesized attrs.pads initializer; the constant_value_input route consumes the `constant_value` tensor directly."
       },
       "attrs": { "mode": "constant", "pads": [1, 2] },
       "inputs": {
@@ -107,7 +106,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
-        "notes": "Typed ORT int32 case using Pad's optional `constant_value` input. Pads remain attrs.pads; the constant_value_input route consumes the `constant_value` tensor directly."
       },
       "attrs": { "mode": "constant", "pads": [1, 2] },
       "inputs": {
@@ -138,7 +137,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesOutOfOrder",
-        "notes": "ORT supplies axes=[3,2] with pads=[1,0,1,0]; this framework represents the equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
@@ -165,7 +164,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxes",
-        "notes": "ORT supplies axes=[1,3] with pads=[0,1,0,1]; this framework represents the equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
@@ -182,7 +181,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesTest2",
-        "notes": "ORT supplies axes=[2,3] with pads=[1,1,1,1]; this framework represents the equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
       "inputs": {
@@ -199,7 +198,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesTest4",
-        "notes": "ORT ConstantPadNegativeAxes, ConstantPadAxesTest1/Test3, and ConstantPadAxesWithOneDimensionSpecified project to this same full-rank pads request."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
@@ -343,7 +342,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Wrap_WebGpu_PadGreaterThanInputDimension",
-        "notes": "CPU-valid upper-padding projection of ORT's WebGPU regression."
       },
       "attrs": { "pads": [0, 5], "mode": "wrap" },
       "inputs": {
@@ -363,7 +362,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Wrap_WebGpu_PadGreaterThanInputDimension",
-        "notes": "Lower-padding projection of the same ORT WebGPU regression: the negative in-coordinate path of the wrap modulo."
       },
       "attrs": { "pads": [5, 0], "mode": "wrap" },
       "inputs": {
@@ -488,16 +487,10 @@
     {
       "name": "constant_vec4_last_axis_crop_removes_every_column",
       "provenance": {
-        "notes": "A positive pad on the innermost axis and a negative one that crops away every source column: the output row is entirely pad value, so the sliced width the vec4 kernel carries is zero. That width has a floor at zero precisely for this shape and no case had ever reached it. The uniform output is the assertion -- no source value may survive."
       },
       "attrs": { "pads": [0, 8, 0, -4] },
-      "inputs": {
-        "data": {
-          "dtype": "float32",
-          "shape": [2, 4],
-          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.31, "scale": 2.0 }
-        }
-      },
       "outputs": { "output": { "dtype": "float32", "shape": [2, 8], "tolerance": 0 } }
     },
     {
@@ -876,7 +869,7 @@
       "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 7, 12] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constant_pad",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX Pad pads/value inputs represented as framework attributes."
       }
     },
     {
@@ -897,7 +890,7 @@
       "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 4, 12] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constant_pad_axes",
-        "notes": "ONNX int64 metadata uses supported framework storage where representable, and Pad inputs are projected to the full-rank pads attribute. The official test_constant_pad_negative_axes fixture projects to this same request."
       }
     },
     {
@@ -916,7 +909,7 @@
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_edge_pad",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX Pad pads/value inputs represented as framework attributes."
       }
     },
     {
@@ -935,7 +928,7 @@
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reflect_pad",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX Pad pads/value inputs represented as framework attributes."
       }
     },
     {
@@ -954,7 +947,7 @@
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_wrap_pad",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX Pad pads/value inputs represented as framework attributes."
       }
     },
     {
@@ -962,7 +955,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, 2, 2, 1, 2, 2] },
       "inputs": {
@@ -979,7 +972,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_3D_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "constant", "pads": [1, 1, 0, 1, 1, 0] },
       "inputs": {
@@ -993,7 +986,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, 1, 0, 1, 1, 0] },
       "inputs": {
@@ -1006,7 +999,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Last_Pad_Slice_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, -1, 0, 1, 1, 0] },
       "inputs": {
@@ -1019,7 +1012,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Last_Slice_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, -1, 0, 1, 0, 0] },
       "inputs": {
@@ -1032,7 +1025,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Reflect_3D_Last_Pad_Slice_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "reflect", "pads": [1, -1, 0, 1, 1, 0] },
       "inputs": {
@@ -1045,7 +1038,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Reflect_3D_Last_Slice_Inner_No_Padding",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "reflect", "pads": [1, -1, 0, 1, 0, 0] },
       "inputs": {
@@ -1058,7 +1051,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_wrap_3D_Inner_No_Padding2",
-        "notes": "ONNX Pad pads/value inputs represented as framework attributes."
       },
       "attrs": { "mode": "wrap", "pads": [1, 2, 0, 1, 2, 0] },
       "inputs": {
@@ -1099,7 +1092,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
-        "notes": "Signed companion for exact integer padding values: -16777217 is representable as int32 but rounds to -16777216 if routed through f32."
       },
       "attrs": { "mode": "constant", "pads": [1, 1] },
       "inputs": {
@@ -1218,7 +1211,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
-        "notes": "Scalar-path companion to the aligned int32/uint32 large-attribute-value cases. The output's last dimension is 3, which excludes the vec4 path and verifies that exact int32 attribute conversion is preserved by the scalar fallback."
       },
       "attrs": { "mode": "constant", "pads": [1, 1] },
       "inputs": {
@@ -1266,7 +1259,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_2D",
-        "notes": "uint8 with a non-constant (edge) mode was untested; edge/reflect/wrap were only covered for f32/int32. Pure value-copy, exact."
       },
       "attrs": { "mode": "edge", "pads": [1, 2, 1, 2] },
       "inputs": {
@@ -1340,7 +1333,7 @@
     {
       "name": "rank3_value_input_scalar_path_perf_compact",
       "provenance": {
-        "notes": "Compact sibling for the rank-3 value-input Pad benchmark; optional `constant_value` input selects constant_value_input and the unaligned inner extent keeps scalar address handling active."
       },
       "attrs": { "mode": "constant", "pads": [1, 0, 2, 1, 0, 2] },
       "inputs": {

 {
   "fixtureArrays": {
     "onnx_backend_constant_pad_input_data": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "Uses Pad's optional `constant_value` input. Pads remain supplied by the `pads` attribute, while the fill value comes directly from the input tensor."
       },
       "attrs": { "mode": "constant", "pads": [1, 2] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "Typed ONNX Runtime int32 case using Pad's optional `constant_value` input. Pads remain supplied by the `pads` attribute, while the fill value comes directly from the input tensor."
       },
       "attrs": { "mode": "constant", "pads": [1, 2] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesOutOfOrder",
+        "notes": "The axes-specific request axes=[3,2], pads=[1,0,1,0] is represented by its equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxes",
+        "notes": "The axes-specific request axes=[1,3], pads=[0,1,0,1] is represented by its equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesTest2",
+        "notes": "The axes-specific request axes=[2,3], pads=[1,1,1,1] is represented by its equivalent full-rank pads attribute."
       },
       "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.ConstantPadAxesTest4",
+        "notes": "A full-rank `pads` request exercises constant padding without an `axes` input."
       },
       "attrs": { "pads": [0, 0, 0, 1, 0, 0, 0, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Wrap_WebGpu_PadGreaterThanInputDimension",
+        "notes": "Upper wrap padding larger than the input length exercises repeated positive modulo mapping."
       },
       "attrs": { "pads": [0, 5], "mode": "wrap" },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Wrap_WebGpu_PadGreaterThanInputDimension",
+        "notes": "Lower wrap padding larger than the input length exercises negative coordinates before modulo normalization."
       },
       "attrs": { "pads": [5, 0], "mode": "wrap" },
       "inputs": {
     {
       "name": "constant_vec4_last_axis_crop_removes_every_column",
       "provenance": {
+        "notes": "A positive innermost pad plus negative cropping removes every source column. The vectorized path carries a zero sliced width, and the uniform output verifies that no source value survives."
       },
       "attrs": { "pads": [0, 8, 0, -4] },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 4], "data": { "kind": "constant", "value": 2.0 } } },
       "outputs": { "output": { "dtype": "float32", "shape": [2, 8], "tolerance": 0 } }
     },
     {
       "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 7, 12] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constant_pad",
+        "notes": "This WebGPU package stores representable ONNX int64 metadata in int32 or uint32 slots. The fixture represents `pads` and `constant_value` inputs as attributes."
       }
     },
     {
       "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 4, 12] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constant_pad_axes",
+        "notes": "This WebGPU package stores representable ONNX int64 metadata in int32 slots. The negative-axis Pad request is expanded to a full-rank `pads` attribute."
       }
     },
     {
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_edge_pad",
+        "notes": "This WebGPU package stores representable ONNX int64 metadata in int32 or uint32 slots. The fixture represents `pads` and `constant_value` inputs as attributes."
       }
     },
     {
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reflect_pad",
+        "notes": "This WebGPU package stores representable ONNX int64 metadata in int32 or uint32 slots. The fixture represents `pads` and `constant_value` inputs as attributes."
       }
     },
     {
       "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 6, 7] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_wrap_pad",
+        "notes": "This WebGPU package stores representable ONNX int64 metadata in int32 or uint32 slots. The fixture represents `pads` and `constant_value` inputs as attributes."
       }
     },
     {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, 2, 2, 1, 2, 2] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_3D_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "constant", "pads": [1, 1, 0, 1, 1, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, 1, 0, 1, 1, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Last_Pad_Slice_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, -1, 0, 1, 1, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_3D_Last_Slice_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "edge", "pads": [1, -1, 0, 1, 0, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Reflect_3D_Last_Pad_Slice_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "reflect", "pads": [1, -1, 0, 1, 1, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Reflect_3D_Last_Slice_Inner_No_Padding",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "reflect", "pads": [1, -1, 0, 1, 0, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_wrap_3D_Inner_No_Padding2",
+        "notes": "The fixture represents the ONNX `pads` and optional `constant_value` inputs as attributes."
       },
       "attrs": { "mode": "wrap", "pads": [1, 2, 0, 1, 2, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "The padding value -16,777,217 is representable as int32 but rounds to -16,777,216 if routed through f32."
       },
       "attrs": { "mode": "constant", "pads": [1, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "An output width of 3 exercises scalar padding and verifies exact conversion of an int32 attribute value above 2^24."
       },
       "attrs": { "mode": "constant", "pads": [1, 1] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
         "test": "PadOpTest.Pad_Edge_2D",
+        "notes": "Uint8 edge padding is a bit-exact value-copy operation."
       },
       "attrs": { "mode": "edge", "pads": [1, 2, 1, 2] },
       "inputs": {
     {
       "name": "rank3_value_input_scalar_path_perf_compact",
       "provenance": {
+        "notes": "A rank-3 Pad with an optional `constant_value` input and unaligned inner extent exercises scalar address handling."
       },
       "attrs": { "mode": "constant", "pads": [1, 0, 2, 1, 0, 2] },
       "inputs": {